From a82f4f08139e013eb823af8a84a0ea23fcffc39b Mon Sep 17 00:00:00 2001 From: "OReilly, Ossian" Date: Mon, 18 May 2026 12:10:32 -0500 Subject: [PATCH 1/3] Add helper function to API to return allocation size for uncompressed buffer. Modify CPU sweep test to use fewer low thread count configurations (reduces unit test time). --- hip/hipCompress.cpp | 12 ++++++++++++ hip/hipCompress.h | 3 +++ tests/test_compress_api_hip.cpp | 2 +- 3 files changed, 16 insertions(+), 1 deletion(-) diff --git a/hip/hipCompress.cpp b/hip/hipCompress.cpp index 770be8a..1a10698 100644 --- a/hip/hipCompress.cpp +++ b/hip/hipCompress.cpp @@ -363,6 +363,18 @@ hipError_t hipDecompress( return launch_err; } +hipError_t hipCompressBufferSize(const hipCompressPlan* plan, size_t* size) +{ + if (!plan) return hipErrorInvalidValue; + if (!size) { + plan->last_error = HIP_COMPRESS_ERROR_NULL_OUTPUT; + return hipErrorInvalidValue; + } + plan->last_error = HIP_COMPRESS_SUCCESS; + *size = (size_t)plan->nx * plan->ny * plan->nz * sizeof(float); + return hipSuccess; +} + hipError_t hipCompressMaxOutputSize(const hipCompressPlan* plan, size_t* size) { if (!plan) return hipErrorInvalidValue; diff --git a/hip/hipCompress.h b/hip/hipCompress.h index 11af240..660c115 100644 --- a/hip/hipCompress.h +++ b/hip/hipCompress.h @@ -158,6 +158,9 @@ hipError_t hipDecompress( hipCompressPlan* plan, hipStream_t user_stream); +// Uncompressed wavelet-layout buffer size in bytes (for allocation). +hipError_t hipCompressBufferSize(const hipCompressPlan* plan, size_t* size); + // Upper bound on compressed output size in bytes (for allocation). hipError_t hipCompressMaxOutputSize(const hipCompressPlan* plan, size_t* size); diff --git a/tests/test_compress_api_hip.cpp b/tests/test_compress_api_hip.cpp index c51ff96..fc23939 100644 --- a/tests/test_compress_api_hip.cpp +++ b/tests/test_compress_api_hip.cpp @@ -1834,7 +1834,7 @@ static void bench_radial_sinc_sweep() int best_fwd_threads = 1, best_inv_threads = 1; std::vector thread_sweep; - for (int t = 1; t <= max_threads; t *= 2) thread_sweep.push_back(t); + for (int t = 8; t <= max_threads; t *= 2) thread_sweep.push_back(t); if (thread_sweep.back() != max_threads) thread_sweep.push_back(max_threads); for (int nt : thread_sweep) { From 1070d3cd8a4a82183d215136c47bad9220596a0e Mon Sep 17 00:00:00 2001 From: "OReilly, Ossian" Date: Mon, 18 May 2026 12:41:20 -0500 Subject: [PATCH 2/3] Add tests for hipCompressBufferSize API. --- tests/test_compress_api_hip.cpp | 31 +++++++++++++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/tests/test_compress_api_hip.cpp b/tests/test_compress_api_hip.cpp index fc23939..f87e4b2 100644 --- a/tests/test_compress_api_hip.cpp +++ b/tests/test_compress_api_hip.cpp @@ -106,6 +106,23 @@ static bool test_plan_lifecycle() size_t expected = (size_t)(8 + 8 * nb + 4) + (size_t)nb * 4 * 32768; if (max_sz != expected) { printf(" FAIL: MaxOutputSize=%zu expected=%zu\n", max_sz, expected); hipCompressDestroyPlan(plan); return false; } printf(" MaxOutputSize=%zu: PASS\n", max_sz); + + // BufferSize + size_t buf_sz = 0; + HIPCHECK(hipCompressBufferSize(plan, &buf_sz)); + size_t expected_buf = (size_t)64 * 64 * 64 * sizeof(float); + if (buf_sz != expected_buf) { printf(" FAIL: BufferSize=%zu expected=%zu\n", buf_sz, expected_buf); hipCompressDestroyPlan(plan); return false; } + printf(" BufferSize=%zu: PASS\n", buf_sz); + hipCompressDestroyPlan(plan); + + // BufferSize with non-cubic dimensions + err = hipCompressCreatePlan(&plan, 128, 64, 32, 0); + if (err != hipSuccess) { printf(" FAIL: create 128x64x32 plan\n"); return false; } + buf_sz = 0; + HIPCHECK(hipCompressBufferSize(plan, &buf_sz)); + expected_buf = (size_t)128 * 64 * 32 * sizeof(float); + if (buf_sz != expected_buf) { printf(" FAIL: BufferSize 128x64x32=%zu expected=%zu\n", buf_sz, expected_buf); hipCompressDestroyPlan(plan); return false; } + printf(" BufferSize 128x64x32=%zu: PASS\n", buf_sz); hipCompressDestroyPlan(plan); return true; @@ -2367,6 +2384,12 @@ static bool test_error_codes() if (err != hipErrorInvalidValue) { printf(" FAIL [MaxOutputSize null plan]\n"); pass = false; } else printf(" ok [MaxOutputSize null plan]\n"); } + { + size_t sz = 0; + hipError_t err = hipCompressBufferSize(nullptr, &sz); + if (err != hipErrorInvalidValue) { printf(" FAIL [BufferSize null plan]\n"); pass = false; } + else printf(" ok [BufferSize null plan]\n"); + } // Allocate a valid plan for the remaining tests hipCompressPlan* plan = nullptr; @@ -2571,6 +2594,14 @@ static bool test_error_codes() pass = false; } + // --- BufferSize errors --- + { + hipError_t err = hipCompressBufferSize(plan, nullptr); + if (!check_error("BufferSize null size", plan, err, + hipErrorInvalidValue, HIP_COMPRESS_ERROR_NULL_OUTPUT)) + pass = false; + } + // --- GetLastError / ErrorString --- { hipCompressError_t e = hipCompressGetLastError(nullptr); From 7f8b224f265f6e75ef7a26a6056ebec8d9a68ea0 Mon Sep 17 00:00:00 2001 From: "OReilly, Ossian" Date: Thu, 28 May 2026 20:05:39 -0500 Subject: [PATCH 3/3] add hipComputeRMS convenience wrapper --- HIP_API.md | 1 + hip/hipCompress.cpp | 20 ++++++++++++++++++++ hip/hipCompress.h | 13 +++++++++++++ tests/test_compress_api_hip.cpp | 28 ++++++++++++++++++++++++++++ 4 files changed, 62 insertions(+) diff --git a/HIP_API.md b/HIP_API.md index 6eb5f1f..4383a08 100644 --- a/HIP_API.md +++ b/HIP_API.md @@ -64,6 +64,7 @@ All functions are declared in [`hip/hipCompress.h`](hip/hipCompress.h). | Function | Description | |----------|-------------| | `hipCopyToWaveletLayout` | Copy from strided grid → contiguous wavelet buffer (zero-pad + optional RMS) | +| `hipComputeRMS` | Convenience wrapper: RMS over extraction window only (no copy) | | `hipCopyFromWaveletLayout` | Copy from wavelet buffer → strided grid (extraction window only) | ### Compression / Decompression diff --git a/hip/hipCompress.cpp b/hip/hipCompress.cpp index 1a10698..77da05b 100644 --- a/hip/hipCompress.cpp +++ b/hip/hipCompress.cpp @@ -292,6 +292,26 @@ hipError_t hipCopyToWaveletLayout( return launch_err; } +hipError_t hipComputeRMS( + const float* d_src, + int ldimx, int ldimxy, + int x0, int y0, int z0, + int ex, int ey, int ez, + double* d_rms_out, + hipCompressPlan* plan, + hipStream_t user_stream) +{ + if (!plan) return hipErrorInvalidValue; + plan->last_error = HIP_COMPRESS_SUCCESS; + if (!d_rms_out) + PLAN_ERROR(plan, HIP_COMPRESS_ERROR_NULL_OUTPUT, hipErrorInvalidValue); + return hipCopyToWaveletLayout( + d_src, ldimx, ldimxy, + x0, y0, z0, ex, ey, ez, + /*d_dst=*/nullptr, d_rms_out, + plan, user_stream); +} + hipError_t hipCopyFromWaveletLayout( const float* d_src, float* d_dst, diff --git a/hip/hipCompress.h b/hip/hipCompress.h index 660c115..7a32174 100644 --- a/hip/hipCompress.h +++ b/hip/hipCompress.h @@ -105,6 +105,19 @@ hipError_t hipCopyToWaveletLayout( hipCompressPlan* plan, hipStream_t user_stream); +// Convenience wrapper: RMS-only mode of hipCopyToWaveletLayout. +// Computes RMS over the ex*ey*ez extraction window; writes no wavelet buffer. +// Equivalent to hipCopyToWaveletLayout(..., d_dst=NULL, d_rms_out, ...). +// Returns HIP_COMPRESS_ERROR_NULL_OUTPUT if d_rms_out is NULL. +hipError_t hipComputeRMS( + const float* d_src, + int ldimx, int ldimxy, + int x0, int y0, int z0, + int ex, int ey, int ez, + double* d_rms_out, + hipCompressPlan* plan, + hipStream_t user_stream); + // Copy from a wavelet-layout buffer back to a strided destination volume. // Only the ex*ey*ez extraction samples are written; padding is skipped. // Wavelet dims are derived from the extraction window. diff --git a/tests/test_compress_api_hip.cpp b/tests/test_compress_api_hip.cpp index f87e4b2..7f7f1b8 100644 --- a/tests/test_compress_api_hip.cpp +++ b/tests/test_compress_api_hip.cpp @@ -1478,6 +1478,20 @@ static bool test_copy_modes() rms_match ? "ok" : "FAIL"); if (!rms_match) pass = false; + // hipComputeRMS wrapper: should match RMS-only path. + double h_rms3 = 0; + HIPCHECK(hipComputeRMS( + d_src, N, N * N, + 0, 0, 0, N, N, N, + plan->d_rms, plan, 0)); + HIPCHECK(hipDeviceSynchronize()); + HIPCHECK(hipMemcpy(&h_rms3, plan->d_rms, sizeof(double), hipMemcpyDeviceToHost)); + bool wrapper_match = (fabsf((float)h_rms3 - (float)h_rms1) < 1e-10f); + printf(" hipComputeRMS wrapper: rms=%.4e (match=%s): %s\n", + h_rms3, wrapper_match ? "yes" : "no", + wrapper_match ? "ok" : "FAIL"); + if (!wrapper_match) pass = false; + // Both NULL should error hipError_t err = hipCopyToWaveletLayout( d_src, N, N * N, @@ -1487,6 +1501,20 @@ static bool test_copy_modes() printf(" both-null: error=%s: %s\n", err_ok ? "yes" : "no", err_ok ? "ok" : "FAIL"); if (!err_ok) pass = false; + // hipComputeRMS with null d_rms_out -> NULL_OUTPUT + hipError_t err_rms = hipComputeRMS( + d_src, N, N * N, + 0, 0, 0, N, N, N, + nullptr, plan, 0); + hipCompressError_t lib_err = hipCompressGetLastError(plan); + bool err_rms_ok = (err_rms != hipSuccess) && + (lib_err == HIP_COMPRESS_ERROR_NULL_OUTPUT); + printf(" hipComputeRMS null-output: error=%s code=%s: %s\n", + err_rms != hipSuccess ? "yes" : "no", + hipCompressErrorString(lib_err), + err_rms_ok ? "ok" : "FAIL"); + if (!err_rms_ok) pass = false; + printf(" mode dispatch: %s\n", pass ? "PASS" : "FAIL"); hipFree(d_src); hipFree(d_dst); hipCompressDestroyPlan(plan);