Skip to content

Fix 2-D/3-D compressed size missing the rate-array alignment padding #14

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
hyfshishen merged 1 commit into from
Oct 13, 2026
Merged

Fix 2-D/3-D compressed size missing the rate-array alignment padding #14

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions src/cuSZp_entry_2D_f32.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down
6 changes: 3 additions & 3 deletions src/cuSZp_entry_2D_f64.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down
6 changes: 3 additions & 3 deletions src/cuSZp_entry_3D_f32.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down
6 changes: 3 additions & 3 deletions src/cuSZp_entry_3D_f64.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down
Loading