From 792d3cb394fe9e49543cbd71d02a12331f1bbade Mon Sep 17 00:00:00 2001 From: Skyler Ruiter Date: Mon, 12 Oct 2026 17:38:53 +0000 Subject: [PATCH] Fix 2-D/3-D compressed size missing the rate-array alignment padding The 2-D and 3-D compress entry points reported glob_sync + blockNum, but the kernels place block payloads after the rate array padded to a multiple of 4 bytes (rate_ofs = (blockNum + 3) / 4 * 4). Callers that store exactly cmpSize bytes lost the last 0-3 bytes, the top bit plane of the last block with a payload, so values there could miss the error bound by a power-of-two multiple of 2*eb. Report the padded length, matching rate_ofs and the 1-D entry points. --- src/cuSZp_entry_2D_f32.cu | 6 +++--- src/cuSZp_entry_2D_f64.cu | 6 +++--- src/cuSZp_entry_3D_f32.cu | 6 +++--- src/cuSZp_entry_3D_f64.cu | 6 +++--- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/src/cuSZp_entry_2D_f32.cu b/src/cuSZp_entry_2D_f32.cu index 025dd2b..42c34f5 100644 --- a/src/cuSZp_entry_2D_f32.cu +++ b/src/cuSZp_entry_2D_f32.cu @@ -44,7 +44,7 @@ void cuSZp_compress_2D_fixed_f32(float* d_oriData, unsigned char* d_cmpBytes, si // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -142,7 +142,7 @@ void cuSZp_compress_2D_plain_f32(float* d_oriData, unsigned char* d_cmpBytes, si // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -240,7 +240,7 @@ void cuSZp_compress_2D_outlier_f32(float* d_oriData, unsigned char* d_cmpBytes, // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); diff --git a/src/cuSZp_entry_2D_f64.cu b/src/cuSZp_entry_2D_f64.cu index 749fbb2..a6ead71 100644 --- a/src/cuSZp_entry_2D_f64.cu +++ b/src/cuSZp_entry_2D_f64.cu @@ -44,7 +44,7 @@ void cuSZp_compress_2D_fixed_f64(double* d_oriData, unsigned char* d_cmpBytes, s // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -142,7 +142,7 @@ void cuSZp_compress_2D_plain_f64(double* d_oriData, unsigned char* d_cmpBytes, s // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -240,7 +240,7 @@ void cuSZp_compress_2D_outlier_f64(double* d_oriData, unsigned char* d_cmpBytes, // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); diff --git a/src/cuSZp_entry_3D_f32.cu b/src/cuSZp_entry_3D_f32.cu index 70f99c0..ac996ff 100644 --- a/src/cuSZp_entry_3D_f32.cu +++ b/src/cuSZp_entry_3D_f32.cu @@ -45,7 +45,7 @@ void cuSZp_compress_3D_fixed_f32(float* d_oriData, unsigned char* d_cmpBytes, si // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -145,7 +145,7 @@ void cuSZp_compress_3D_plain_f32(float* d_oriData, unsigned char* d_cmpBytes, si // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -245,7 +245,7 @@ void cuSZp_compress_3D_outlier_f32(float* d_oriData, unsigned char* d_cmpBytes, // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); diff --git a/src/cuSZp_entry_3D_f64.cu b/src/cuSZp_entry_3D_f64.cu index 6ac5ffe..a8bd704 100644 --- a/src/cuSZp_entry_3D_f64.cu +++ b/src/cuSZp_entry_3D_f64.cu @@ -45,7 +45,7 @@ void cuSZp_compress_3D_fixed_f64(double* d_oriData, unsigned char* d_cmpBytes, s // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -145,7 +145,7 @@ void cuSZp_compress_3D_plain_f64(double* d_oriData, unsigned char* d_cmpBytes, s // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset); @@ -245,7 +245,7 @@ void cuSZp_compress_3D_outlier_f64(double* d_oriData, unsigned char* d_cmpBytes, // Obtain compression ratio and move data back to CPU. cudaMemcpy(&glob_sync, d_cmpOffset+cmpOffSize-1, sizeof(unsigned int), cudaMemcpyDeviceToHost); - *cmpSize = (size_t)glob_sync + blockNum; + *cmpSize = (size_t)glob_sync + (blockNum + 3) / 4 * 4; // include the rate array's 4-byte alignment padding (matches rate_ofs in the kernels) // Free memory that is used. cudaFree(d_cmpOffset);