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);