diff --git a/CHANGELOG.md b/CHANGELOG.md index dc0469ca1..aac9ad0ea 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -23640,6 +23640,14 @@ See ADR-0792. integration remains intact. +- **The HIP integer ADM kernels read their buffer description through a + pointer again (ADR-0759).** The change listed as "HIP ADM: AdmBufferHip + passed by pointer" had been undone by a later merge, so four kernels were + still copying the 328-byte struct into their arguments on every launch. It + is back: each launch now passes one device pointer. Scores are + byte-identical and 1080p throughput is unchanged within run-to-run noise. + + - **`float_ansnr_hip`: `hipMemcpy2DAsync` direction tagged `hipMemcpyDeviceToDevice` for host→device transfer**: `submit_fex_hip` in `core/src/feature/hip/float_ansnr_hip.c:324,330` copied diff --git a/changelog.d/fixed/hip-adm-buffer-by-pointer-reapplied.md b/changelog.d/fixed/hip-adm-buffer-by-pointer-reapplied.md new file mode 100644 index 000000000..47ca2291e --- /dev/null +++ b/changelog.d/fixed/hip-adm-buffer-by-pointer-reapplied.md @@ -0,0 +1,6 @@ +- **The HIP integer ADM kernels read their buffer description through a + pointer again (ADR-0759).** The change listed as "HIP ADM: AdmBufferHip + passed by pointer" had been undone by a later merge, so four kernels were + still copying the 328-byte struct into their arguments on every launch. It + is back: each launch now passes one device pointer. Scores are + byte-identical and 1080p throughput is unchanged within run-to-run noise. diff --git a/core/src/feature/hip/AGENTS.md b/core/src/feature/hip/AGENTS.md index 671f5d86f..cd4a62cba 100644 --- a/core/src/feature/hip/AGENTS.md +++ b/core/src/feature/hip/AGENTS.md @@ -252,27 +252,6 @@ The `VMAF_HSACO_WEAK_STUB` macro in `hip_hsaco_stubs.c` is retained as a documented pattern for in-progress ports of *new* extractors; it is currently used by zero extractors. -## AdmBufferHip struct-by-value kernel parameters — P1 known issue (Research-0755) - -`AdmBufferHip` (defined in `integer_adm_hip.h:70–96`) is a ~272-byte struct -containing 6 DWT band sub-structs (each 4 device pointers) plus 8 additional -device-pointer fields. It is currently passed by value in multiple `__global__` -kernel signatures in `integer_adm/adm_csf.hip` and `integer_adm/adm_cm.hip`. - -This mirrors the PR #93 F3 finding on the CUDA side. Consequences: - -- Every GPU thread's stack receives a full 272-byte copy via the kernel-argument - buffer path. On RDNA/GCN this adds measurable argument-passing overhead. -- Structs this large risk hitting the HIP/AMDDriver kernel-argument limit (varies - per target; typically 1024–4096 bytes total across all args). - -**Recommended fix**: replace `AdmBufferHip buf` parameters with -`const AdmBufferHip * __restrict__ buf` (pass a pointer to a device-side copy -of the struct). No correctness impact — only the passing convention changes. - -Until fixed: do NOT add new `__global__` parameters of type `AdmBufferHip` by -value. Any new ADM kernel should take a pointer. - ## extern "C" macro-instantiation pattern is correct (Research-0755) Several ADM kernel files (`adm_csf.hip`, `adm_csf_den.hip`, `adm_dwt2.hip`) @@ -286,37 +265,39 @@ The pattern is load-bearing. Do not "fix" it by adding an additional `extern "C"` declaration inside the macro body — that would create a nested `extern "C"` which is legal in C++ but redundant and confusing to reviewers. -## AdmBufferHip MUST be passed by pointer — invariant (ADR-0759) - -**Resolved**: The P1 known issue documented above (struct-by-value in ADM kernel -signatures) has been fixed by ADR-0759 (PR perf/hip-adm-buffer-by-pointer-20260529). - -**Invariant going forward**: Any new `__global__` kernel that needs `AdmBufferHip` -(or any other large parameter struct) MUST accept it as a pointer parameter, not by -value. The host launch site must: - -1. Hold a device-side copy of the struct allocated in `init_fex_hip` (or equivalent - init path) via `hipMalloc`. -2. Populate it via `hipMemcpy(hipMemcpyHostToDevice)` after all device pointers inside - the struct are set. -3. Pass `&dev_ptr_var` (address of the device pointer variable) as the kernel arg. - -Pattern: - -```c -/* host dispatch helper — correct */ -AdmBufferHip *buf_dev = s->buf_dev; /* device pointer, set in init */ -void *args[] = {&buf_dev, /* ... */}; -hipModuleLaunchKernel(fn, ..., args, NULL); -``` - -Rationale: `AdmBufferHip` is ~272 bytes. Passing by value marshals the full struct -through the per-launch argument buffer on every call. Pointer passing reduces this -to 8 bytes (one pointer) per launch. - -The same rule applies to `AdmFixedParametersHip` (~244 bytes) once that follow-up -is scoped; see ADR-0759 alternatives table. Do not add new by-value large struct -parameters to ADM kernels without an explicit ADR justification. +## AdmBufferHip is passed by pointer — invariant (ADR-0759, T-HIP-ADM-ADR0759-REVERTED-2026-09-18) + +- `adm_csf_kernel_1_4`, `i4_adm_csf_kernel_1_4` (`integer_adm/adm_csf.hip`), + `i4_adm_cm_line_kernel` and `adm_cm_line_kernel_8` (`integer_adm/adm_cm.hip`) + take `const AdmBufferHip *__restrict__ buf_ptr`. By value, the 328-byte + struct was copied into every launch's kernel arguments. +- Host: `AdmStateHip::buf_dev` is a device copy of `s->buf`. + `adm_hip_upload_buf()` (`hipMalloc` + `hipMemcpy` HtoD) runs at the end of + `adm_hip_init_device()`, after `adm_hip_slice_bands()` and + `adm_hip_slice_results()`. `adm_hip_free_buf_dev()` frees it in + `close_fex_hip()` and on both init failure paths. +- Launch argument = `(void *)&s->buf_dev`, the address of the variable that + holds the device pointer (ADR-0537 rule above). +- Precondition: nothing writes `s->buf` between init and close, and no launch + passes a modified copy. Code that changes `s->buf` after init (per-scale + band pointers, a resize) must upload it again before the next launch, or + the kernels read stale pointers. +- New ADM kernels that need `AdmBufferHip` take a pointer. + `AdmFixedParametersHip` (248 bytes) and `WarpShift` are still passed by + value; changing them needs its own measurement. +- History: #101 (`31a51afb2`) implemented this; #102 (`92ea978a4`, a CUDA + ciede change cut from an older base) reverted it in a merge without + mentioning it. On a conflict in these files keep the pointer form; + `grep -n 'AdmBufferHip buf' core/src/feature/hip/integer_adm/*.hip` must + print nothing. +- Measured on gfx1036: each kernel's argument segment is 320 bytes smaller; + per-thread scratch and VGPRs do not change because of the pointer. The + 936-byte scratch of `adm_cm_line_kernel_8` is VGPR spilling (239 spills at + the 128-register cap), not the struct. End-to-end fps is unchanged within + noise. +- Not a CUDA mirror: the CUDA twin passes `AdmBufferCuda` by value and always + has (the ADR-0756 audit lists those kernels). Research-0759's statement that + CUDA uses a pointer, and ADR-0759's "matches the CUDA pattern", are wrong. ## ms_ssim_vert_lcs kernel and host partials must both be `double` (ADR-1071) diff --git a/core/src/feature/hip/integer_adm/adm_cm.hip b/core/src/feature/hip/integer_adm/adm_cm.hip index 1898a0eda..45082702f 100644 --- a/core/src/feature/hip/integer_adm/adm_cm.hip +++ b/core/src/feature/hip/integer_adm/adm_cm.hip @@ -389,14 +389,14 @@ s0_cm_block_reduce(const int64_t (&accum_row)[rows_per_thread], int y, int end_r * -------------------------------------------------------------------- */ template __device__ __forceinline__ void -adm_cm_line_kernel_body(AdmBufferHip buf, int h, int w, int start_row, int end_row, int start_col, - int end_col, int src_stride, AdmFixedParametersHip params, - int64_t *accum_global, WarpShift ws, const uint32_t shift_inner_accum, - const uint32_t add_shift_inner_accum) +adm_cm_line_kernel_body(const AdmBufferHip *__restrict__ buf_ptr, int h, int w, int start_row, + int end_row, int start_col, int end_col, int src_stride, + AdmFixedParametersHip params, int64_t *accum_global, WarpShift ws, + const uint32_t shift_inner_accum, const uint32_t add_shift_inner_accum) { - const hip_adm_dwt_band_t *ref = &buf.ref_dwt2; - const hip_adm_dwt_band_t *dis = &buf.dis_dwt2; - const hip_adm_dwt_band_t *csf_f = &buf.csf_f; + const hip_adm_dwt_band_t *ref = &buf_ptr->ref_dwt2; + const hip_adm_dwt_band_t *dis = &buf_ptr->dis_dwt2; + const hip_adm_dwt_band_t *csf_f = &buf_ptr->csf_f; const int band = blockIdx.z + 1; int16_t *const *flt_angles = csf_f->bands + 1; @@ -496,16 +496,22 @@ extern "C" { * i4_adm_cm_line_kernel — scales 1-3 compute. Writes accum_per_thread * scratch consumed by adm_cm_reduce_line_kernel_4. Same launch shape as * the CUDA twin (host TU controls grid dims). + * + * This kernel and adm_cm_line_kernel_8 take the device-resident + * AdmBufferHip by pointer (ADR-0759): the host uploads it once in + * init_fex_hip() and never changes it, so every launch reads the bytes the + * by-value form used to copy into the kernel arguments. * -------------------------------------------------------------------- */ -__global__ void i4_adm_cm_line_kernel(AdmBufferHip buf, int h, int w, int top, int bottom, int left, - int right, int start_row, int end_row, int start_col, - int end_col, int src_stride, int /* csf_a_stride */, - int scale, int buffer_h, int buffer_stride, - int32_t *accum_per_thread, AdmFixedParametersHip params) +__global__ void i4_adm_cm_line_kernel(const AdmBufferHip *__restrict__ buf_ptr, int h, int w, + int top, int bottom, int left, int right, int start_row, + int end_row, int start_col, int end_col, int src_stride, + int /* csf_a_stride */, int scale, int buffer_h, + int buffer_stride, int32_t *accum_per_thread, + AdmFixedParametersHip params) { - const hip_i4_adm_dwt_band_t *ref = &buf.i4_ref_dwt2; - const hip_i4_adm_dwt_band_t *dis = &buf.i4_dis_dwt2; - const hip_i4_adm_dwt_band_t *csf_f = &buf.i4_csf_f; + const hip_i4_adm_dwt_band_t *ref = &buf_ptr->i4_ref_dwt2; + const hip_i4_adm_dwt_band_t *dis = &buf_ptr->i4_dis_dwt2; + const hip_i4_adm_dwt_band_t *csf_f = &buf_ptr->i4_csf_f; const int band = blockIdx.z + 1; int32_t *const *flt_angles = csf_f->bands + 1; @@ -532,17 +538,17 @@ __global__ void i4_adm_cm_line_kernel(AdmBufferHip buf, int h, int w, int top, i } } -__global__ void adm_cm_line_kernel_8(AdmBufferHip buf, int h, int w, int /* top */, - int /* bottom */, int /* left */, int /* right */, - int start_row, int end_row, int start_col, int end_col, - int src_stride, int /* csf_a_stride */, int /* buffer_h */, - int /* buffer_stride */, int32_t * /* accum_per_block */, - AdmFixedParametersHip params, int /* scale */, - int64_t *accum_global, WarpShift ws, +__global__ void adm_cm_line_kernel_8(const AdmBufferHip *__restrict__ buf_ptr, int h, int w, + int /* top */, int /* bottom */, int /* left */, + int /* right */, int start_row, int end_row, int start_col, + int end_col, int src_stride, int /* csf_a_stride */, + int /* buffer_h */, int /* buffer_stride */, + int32_t * /* accum_per_block */, AdmFixedParametersHip params, + int /* scale */, int64_t *accum_global, WarpShift ws, const uint32_t shift_inner_accum, const uint32_t add_shift_inner_accum) { - adm_cm_line_kernel_body<8>(buf, h, w, start_row, end_row, start_col, end_col, src_stride, + adm_cm_line_kernel_body<8>(buf_ptr, h, w, start_row, end_row, start_col, end_col, src_stride, params, accum_global, ws, shift_inner_accum, add_shift_inner_accum); } diff --git a/core/src/feature/hip/integer_adm/adm_csf.hip b/core/src/feature/hip/integer_adm/adm_csf.hip index 9174d20d7..2cb1d5106 100644 --- a/core/src/feature/hip/integer_adm/adm_csf.hip +++ b/core/src/feature/hip/integer_adm/adm_csf.hip @@ -26,13 +26,21 @@ #include "adm_decouple_inline.hip" -template -static __device__ __forceinline__ void copy_vec_4(const int32_t *__restrict__ in, - int32_t *__restrict__ out) +/* The two extern "C" kernels at the bottom are resolved by name from the + * embedded HSACO (hipModuleGetFunction); everything else is a device helper + * with internal linkage. + * + * Both kernels take the device-resident AdmBufferHip by pointer (ADR-0759). + * The host uploads it once in init_fex_hip(), after the band pointers are + * sliced, and never changes it afterwards, so every launch reads the same + * bytes the by-value form used to copy into the kernel arguments. */ +namespace { - //__builtin_assume_aligned(in, 16); - //__builtin_assume_aligned(out, 16); +template +__device__ __forceinline__ void copy_vec_4(const int32_t *__restrict__ in, + int32_t *__restrict__ out) +{ static_assert(cols_per_thread % 4 == 0, "implemented only for a multiple of 4"); #pragma unroll for (int col = 0; col < cols_per_thread; col += 4) { @@ -41,12 +49,9 @@ static __device__ __forceinline__ void copy_vec_4(const int32_t *__restrict__ in } template -static __device__ __forceinline__ void copy_vec_4(const int16_t *__restrict__ in, - int16_t *__restrict__ out) +__device__ __forceinline__ void copy_vec_4(const int16_t *__restrict__ in, + int16_t *__restrict__ out) { - // __builtin_assume_aligned(in, 8); - // __builtin_assume_aligned(out, 8); - static_assert(cols_per_thread % 4 == 0, "implemented only for a multiple of 4"); #pragma unroll for (int col = 0; col < cols_per_thread; col += 4) { @@ -54,20 +59,31 @@ static __device__ __forceinline__ void copy_vec_4(const int16_t *__restrict__ in } } +/* The dis-band sample of the kernel's band (1 = H, 2 = V, 3 = D). */ +template __device__ __forceinline__ T csf_band_sample(int band, T h, T v, T d) +{ + if (band == 1) { + return h; + } + if (band == 2) { + return v; + } + return d; +} + /* Scales 1-3 CSF with inline decouple — reads ref/dis DWT2, writes only csf_f */ template -__device__ __forceinline__ void i4_adm_csf_kernel(AdmBufferHip buf, int scale, int top, int bottom, - int left, int right, int stride, - AdmFixedParametersHip params) +__device__ __forceinline__ void +i4_adm_csf_kernel(const AdmBufferHip *__restrict__ buf_ptr, int scale, int top, int bottom, + int left, int right, int stride, AdmFixedParametersHip params) { - const int band = blockIdx.z + 1; - const hip_i4_adm_dwt_band_t *ref = &buf.i4_ref_dwt2; - const hip_i4_adm_dwt_band_t *dis = &buf.i4_dis_dwt2; - int32_t *flt_ptr = buf.i4_csf_f.bands[band]; + const hip_i4_adm_dwt_band_t *ref = &buf_ptr->i4_ref_dwt2; + const hip_i4_adm_dwt_band_t *dis = &buf_ptr->i4_dis_dwt2; + int32_t *flt_ptr = buf_ptr->i4_csf_f.bands[band]; - int y = top + (blockIdx.y * blockDim.y + threadIdx.y) * rows_per_thread; - int x = left + (blockIdx.x * blockDim.x + threadIdx.x) * cols_per_thread; + const int y = top + (blockIdx.y * blockDim.y + threadIdx.y) * rows_per_thread; + const int x = left + (blockIdx.x * blockDim.x + threadIdx.x) * cols_per_thread; const uint32_t i_rfactor = params.i_rfactor[scale * 3 + blockIdx.z]; const uint32_t FIX_ONE_BY_30 = 143165577; @@ -88,28 +104,21 @@ __device__ __forceinline__ void i4_adm_csf_kernel(AdmBufferHip buf, int scale, i for (int col = 0; col < cols_per_thread; ++col) { const int idx = row_off + col; - int32_t oh = ref->band_h[idx]; - int32_t ov = ref->band_v[idx]; - int32_t od = ref->band_d[idx]; - int32_t th = dis->band_h[idx]; - int32_t tv = dis->band_v[idx]; - int32_t td = dis->band_d[idx]; - - int angle_flag = decouple_angle_flag_s123(oh, ov, th, tv); - int32_t r_val = decouple_r_s123(oh, ov, od, th, tv, td, band - 1, angle_flag, - adm_enhn_gain_limit); - - int32_t t_val; - if (band == 1) - t_val = th; - else if (band == 2) - t_val = tv; - else - t_val = td; - - int32_t a_val = t_val - r_val; - - int32_t dst_val = + const int32_t oh = ref->band_h[idx]; + const int32_t ov = ref->band_v[idx]; + const int32_t od = ref->band_d[idx]; + const int32_t th = dis->band_h[idx]; + const int32_t tv = dis->band_v[idx]; + const int32_t td = dis->band_d[idx]; + + const int angle_flag = decouple_angle_flag_s123(oh, ov, th, tv); + const int32_t r_val = decouple_r_s123(oh, ov, od, th, tv, td, band - 1, angle_flag, + adm_enhn_gain_limit); + + const int32_t t_val = csf_band_sample(band, th, tv, td); + const int32_t a_val = t_val - r_val; + + const int32_t dst_val = (int32_t)(((i_rfactor * int64_t(a_val)) + add_bef_shift_dst) >> shift_dst); flt_vec[col] = (int32_t)((((int64_t)FIX_ONE_BY_30 * abs(dst_val)) + add_bef_shift_flt) >> @@ -125,16 +134,17 @@ __constant__ const uint16_t i_shiftsadd[4] = {0, 16384, 16384, 65535}; /* Scale-0 CSF with inline decouple — reads ref/dis DWT2, writes only csf_f */ template -__device__ __forceinline__ void adm_csf_kernel(AdmBufferHip buf, int top, int bottom, int left, - int right, int stride, AdmFixedParametersHip params) +__device__ __forceinline__ void adm_csf_kernel(const AdmBufferHip *__restrict__ buf_ptr, int top, + int bottom, int left, int right, int stride, + AdmFixedParametersHip params) { const int band = blockIdx.z + 1; - const hip_adm_dwt_band_t *ref = &buf.ref_dwt2; - const hip_adm_dwt_band_t *dis = &buf.dis_dwt2; - int16_t *flt_ptr = buf.csf_f.bands[band]; - int y = top + (blockIdx.y * blockDim.y + threadIdx.y) * rows_per_thread; - int x = left + (blockIdx.x * blockDim.x + threadIdx.x) * cols_per_thread; + const hip_adm_dwt_band_t *ref = &buf_ptr->ref_dwt2; + const hip_adm_dwt_band_t *dis = &buf_ptr->dis_dwt2; + int16_t *flt_ptr = buf_ptr->csf_f.bands[band]; + const int y = top + (blockIdx.y * blockDim.y + threadIdx.y) * rows_per_thread; + const int x = left + (blockIdx.x * blockDim.x + threadIdx.x) * cols_per_thread; const uint32_t i_rfactor = params.i_rfactor[blockIdx.z]; const uint16_t FIX_ONE_BY_30 = 4369; //(1/30)*2^17 @@ -151,29 +161,22 @@ __device__ __forceinline__ void adm_csf_kernel(AdmBufferHip buf, int top, int bo for (int col = 0; col < cols_per_thread; ++col) { const int idx = row_off + col; - int16_t oh = ref->band_h[idx]; - int16_t ov = ref->band_v[idx]; - int16_t od = ref->band_d[idx]; - int16_t th = dis->band_h[idx]; - int16_t tv = dis->band_v[idx]; - int16_t td = dis->band_d[idx]; - - int angle_flag = decouple_angle_flag_s0(oh, ov, th, tv); - int16_t r_val = decouple_r_s0(oh, ov, od, th, tv, td, band - 1, angle_flag, - adm_enhn_gain_limit); - - int16_t t_val; - if (band == 1) - t_val = th; - else if (band == 2) - t_val = tv; - else - t_val = td; - - int16_t a_val = t_val - r_val; - - int32_t dst_val = i_rfactor * (uint32_t)a_val; - int16_t i16_dst_val = (dst_val + i_shiftsadd[band]) >> i_shifts[band]; + const int16_t oh = ref->band_h[idx]; + const int16_t ov = ref->band_v[idx]; + const int16_t od = ref->band_d[idx]; + const int16_t th = dis->band_h[idx]; + const int16_t tv = dis->band_v[idx]; + const int16_t td = dis->band_d[idx]; + + const int angle_flag = decouple_angle_flag_s0(oh, ov, th, tv); + const int16_t r_val = decouple_r_s0(oh, ov, od, th, tv, td, band - 1, angle_flag, + adm_enhn_gain_limit); + + const int16_t t_val = csf_band_sample(band, th, tv, td); + const int16_t a_val = t_val - r_val; + + const int32_t dst_val = i_rfactor * (uint32_t)a_val; + const int16_t i16_dst_val = (dst_val + i_shiftsadd[band]) >> i_shifts[band]; flt_vec[col] = ((FIX_ONE_BY_30 * abs((int32_t)i16_dst_val)) + 2048) >> 12; } copy_vec_4(flt_vec, flt_ptr + row_off); @@ -181,22 +184,24 @@ __device__ __forceinline__ void adm_csf_kernel(AdmBufferHip buf, int top, int bo } } +} // namespace + #define ADM_CSF_KERNEL(rows_per_thread, cols_per_thread) \ __global__ void adm_csf_kernel_##rows_per_thread##_##cols_per_thread( \ - AdmBufferHip buf, int top, int bottom, int left, int right, int stride, \ - AdmFixedParametersHip params) \ + const AdmBufferHip *__restrict__ buf_ptr, int top, int bottom, int left, int right, \ + int stride, AdmFixedParametersHip params) \ { \ - adm_csf_kernel(buf, top, bottom, left, right, stride, \ - params); \ + adm_csf_kernel(buf_ptr, top, bottom, left, right, \ + stride, params); \ } #define I4_ADM_CSF_KERNEL(rows_per_thread, cols_per_thread) \ __global__ void i4_adm_csf_kernel_##rows_per_thread##_##cols_per_thread( \ - AdmBufferHip buf, int scale, int top, int bottom, int left, int right, int stride, \ - AdmFixedParametersHip params) \ + const AdmBufferHip *__restrict__ buf_ptr, int scale, int top, int bottom, int left, \ + int right, int stride, AdmFixedParametersHip params) \ { \ - i4_adm_csf_kernel(buf, scale, top, bottom, left, right, \ - stride, params); \ + i4_adm_csf_kernel(buf_ptr, scale, top, bottom, left, \ + right, stride, params); \ } extern "C" { diff --git a/core/src/feature/hip/integer_adm_hip.c b/core/src/feature/hip/integer_adm_hip.c index d50ab5aaf..c3ba83991 100644 --- a/core/src/feature/hip/integer_adm_hip.c +++ b/core/src/feature/hip/integer_adm_hip.c @@ -116,6 +116,13 @@ typedef struct AdmStateHip { hipFunction_t func_adm_cm_line_kernel_8; hipFunction_t func_i4_adm_cm_line_kernel; + /* ADR-0759: device copy of `buf`. The two CSF and the two CM compute + * kernels take `const AdmBufferHip *` and read their band pointers from + * here instead of receiving the whole struct by value on every launch. + * Uploaded once by adm_hip_upload_buf(); `buf` does not change after + * that, so the copy stays equal to it until close(). */ + AdmBufferHip *buf_dev; + /* ADR-1211: device staging for the scale-0 luma plane. * The HIP backend is host-pic (ADR-0530): `VmafPicture::data[]` points at * HOST memory. The DWT2 kernel is a device kernel, so the plane has to be @@ -759,7 +766,7 @@ static int adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, i const int BLOCKX = 32; const int BLOCKY = 4; - void *args[] = {buf, &top, &bottom, &left, &right, &stride, p}; + void *args[] = {(void *)&s->buf_dev, &top, &bottom, &left, &right, &stride, p}; hipError_t rc = hipModuleLaunchKernel( s->func_adm_csf_kernel_1_4, (uint32_t)DIV_ROUND_UP(right - left, BLOCKX * cols_per_thread), (uint32_t)DIV_ROUND_UP(bottom - top, BLOCKY * rows_per_thread), 3, (uint32_t)BLOCKX, @@ -794,7 +801,7 @@ static int i4_adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int scale, i const int BLOCKX = 32; const int BLOCKY = 4; - void *args[] = {buf, &scale, &top, &bottom, &left, &right, &stride, p}; + void *args[] = {(void *)&s->buf_dev, &scale, &top, &bottom, &left, &right, &stride, p}; hipError_t rc = hipModuleLaunchKernel(s->func_i4_adm_csf_kernel_1_4, (uint32_t)DIV_ROUND_UP(right - left, BLOCKX * cols_per_thread), @@ -888,7 +895,7 @@ static int i4_adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, /* inner CM kernel */ { const int BLOCKX = 128; - void *args[] = {buf, + void *args[] = {(void *)&s->buf_dev, &h, &w, &top, @@ -976,7 +983,7 @@ static int adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, in const int rows_per_thread = 8; const int BLOCKX = 32; const int BLOCKY = 4; - void *args[] = {buf, + void *args[] = {(void *)&s->buf_dev, &h, &w, &top, @@ -1476,6 +1483,35 @@ static void adm_hip_slice_results(AdmStateHip *s) } } +/* ADR-0759: upload `s->buf` to the device copy the CSF and CM kernels read. + * Call it after adm_hip_slice_bands() and adm_hip_slice_results(), once every + * pointer in the struct is final. Nothing writes `s->buf` between init and + * close, so one upload serves every launch; code that changes `s->buf` after + * init must upload it again before the next launch. On failure, frees what + * it allocated. */ +static int adm_hip_upload_buf(AdmStateHip *s) +{ + void *dev = NULL; + hipError_t hip_err = hipMalloc(&dev, sizeof(s->buf)); + if (hip_err != hipSuccess) + return hip_rc(hip_err); + hip_err = hipMemcpy(dev, &s->buf, sizeof(s->buf), hipMemcpyHostToDevice); + if (hip_err != hipSuccess) { + (void)hipFree(dev); + return hip_rc(hip_err); + } + s->buf_dev = dev; + return 0; +} + +static void adm_hip_free_buf_dev(AdmStateHip *s) +{ + if (s->buf_dev != NULL) { + (void)hipFree(s->buf_dev); + s->buf_dev = NULL; + } +} + /* Every device resource init_fex_hip() needs, in dependency order. On * failure, releases what it created. */ static int adm_hip_init_device(AdmStateHip *s, unsigned w, unsigned h, unsigned bpc) @@ -1506,7 +1542,14 @@ static int adm_hip_init_device(AdmStateHip *s, unsigned w, unsigned h, unsigned adm_hip_slice_bands(s, h); adm_hip_slice_results(s); - return 0; + err = adm_hip_upload_buf(s); + if (err) { + adm_hip_free_luma(s); + adm_hip_free_buffers(s); + adm_hip_unload_modules(s); + adm_hip_destroy_stream(s); + } + return err; } #endif /* HAVE_HIPCC */ @@ -1569,6 +1612,7 @@ static int init_fex_hip(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt, if (s->feature_name_dict == NULL) { /* The framework never calls close() after a failed init(), so every * device resource is released here. */ + adm_hip_free_buf_dev(s); adm_hip_free_luma(s); adm_hip_free_buffers(s); adm_hip_unload_modules(s); @@ -1648,6 +1692,7 @@ static int close_fex_hip(VmafFeatureExtractor *fex) #ifdef HAVE_HIPCC rc = adm_hip_close_stream(s); adm_hip_unload_modules(s); + adm_hip_free_buf_dev(s); adm_hip_free_luma(s); adm_hip_free_buffers(s); #endif /* HAVE_HIPCC */ diff --git a/docs/rebase-notes.md b/docs/rebase-notes.md index 1cb185585..31cc76159 100644 --- a/docs/rebase-notes.md +++ b/docs/rebase-notes.md @@ -1,6 +1,32 @@ # Rebase notes +## perf/hip-adm-buffer-by-pointer — HIP integer ADM buffer by pointer, re-applied (2026-09-18) + +ADR-0759 (the four HIP ADM kernels that read `AdmBufferHip` take it by +pointer) landed in #101 and was silently reverted by the next merge, #102, +whose branch predated it (T-HIP-ADM-ADR0759-REVERTED-2026-09-18). The HIP +twin is fork-only, so an upstream sync cannot conflict with it; the risk is +another fork branch cut before this change. When merging or rebasing anything +that touches these files, keep: + +- `core/src/feature/hip/integer_adm/adm_csf.hip` and `adm_cm.hip`: + `adm_csf_kernel_1_4`, `i4_adm_csf_kernel_1_4`, `i4_adm_cm_line_kernel` and + `adm_cm_line_kernel_8` take `const AdmBufferHip *__restrict__ buf_ptr`. + `adm_csf.hip` also carries its lint restructure (helpers in an anonymous + namespace, `csf_band_sample()`); kernel names and launch layouts are + unchanged. +- `core/src/feature/hip/integer_adm_hip.c`: `AdmStateHip::buf_dev`, uploaded + by `adm_hip_upload_buf()` at the end of `adm_hip_init_device()` and freed by + `adm_hip_free_buf_dev()` in `close_fex_hip()` and the init failure paths; + the four launches pass `(void *)&s->buf_dev`. + +Check after any merge that touches them: +`grep -n 'AdmBufferHip buf' core/src/feature/hip/integer_adm/*.hip` must print +nothing. Kernel and host must change together: a by-value kernel launched +with `&s->buf_dev` gets 328 bytes copied from that address as the struct and +dereferences whatever follows `buf_dev` in `AdmStateHip`. + ## fix/gpu-adm-dwt2-16bit-overflow — GPU integer ADM 16-bit vertical DWT sum (2026-09-18) Upstream Netflix/vmaf's CUDA integer ADM sums the scale-0 vertical DWT response diff --git a/docs/research/research-0759-hip-adm-buffer-by-pointer.md b/docs/research/research-0759-hip-adm-buffer-by-pointer.md index 4b4b4ed27..9c4be2bb4 100644 --- a/docs/research/research-0759-hip-adm-buffer-by-pointer.md +++ b/docs/research/research-0759-hip-adm-buffer-by-pointer.md @@ -81,3 +81,29 @@ the change is bit-exact: kernel bodies access the same device memory via `buf_pt instead of copying `buf.field` onto the device register file. All load addresses are identical. Runtime verification on AMD hardware is deferred but the change is safe to merge into a DRAFT PR. + +## Addendum 2026-09-18: reverted, re-applied, measured + +- **Reverted.** The implementation (#101, `31a51afb2`) was undone by the next + merge, #102 (`92ea978a4`), a CUDA ciede change whose branch predated it. The + kernels passed the struct by value from then until + `perf/hip-adm-buffer-by-pointer` re-applied the change on the current code + (T-HIP-ADM-ADR0759-REVERTED-2026-09-18 in `docs/state.md`). +- **Size.** `AdmBufferHip` is 328 bytes, not ~272: the gfx1036 kernel argument + layout puts the next argument at offset 328. `AdmFixedParametersHip` is 248 + bytes. +- **CUDA parity statement above is wrong.** The CUDA twin passes + `AdmBufferCuda` by value; the ADR-0756 audit lists those kernels, and none + was changed. +- **Runtime verification done** on a gfx1036 iGPU (ROCm 7.2): HIP output is + byte-identical at `%.17g` before and after on 8, 10, 12 and 16-bit inputs, + odd frame sizes and 1080p clips. +- **Measured effect.** Each of the four kernels' argument segment shrinks by + 320 bytes (856 to 536 on the CSF kernels, 904 to 584 and 968 to 648 on the + CM kernels). Per-thread scratch and VGPR counts do not change because of the + pointer. The 936-byte scratch on `adm_cm_line_kernel_8` is VGPR spilling + (239 spills at the 128-register cap), not a copy of the struct; a + `__launch_bounds__(128)` experiment raised the cap to 256 and still spilled + 112 registers. End-to-end HIP ADM throughput on a 60-frame 1080p 10-bit clip + is unchanged within noise (median 29.5 fps before, 30.0 after, 10 + alternating runs, spread 27.4 to 31.1). diff --git a/docs/state.md b/docs/state.md index 13a2422e8..80da248b2 100644 --- a/docs/state.md +++ b/docs/state.md @@ -456,6 +456,7 @@ landed fix yet._ ## Recently closed +| **T-HIP-ADM-ADR0759-REVERTED-2026-09-18** | **A stale merge silently reverted ADR-0759, so the HIP integer ADM kernels passed `AdmBufferHip` by value again.** ADR-0759 (Accepted) moved `adm_csf_kernel_1_4`, `i4_adm_csf_kernel_1_4`, `i4_adm_cm_line_kernel` and `adm_cm_line_kernel_8` to `const AdmBufferHip *` with a device copy uploaded at init, in `31a51afb2` (#101). The next merge, `92ea978a4` (#102, a CUDA ciede change cut from an older base), restored the by-value signatures and dropped `buf_dev` without saying so. Since then every launch copied the 328-byte struct into its kernel arguments while the ADR, the HIP `AGENTS.md` and the changelog described the pointer form. The precondition still holds: `s->buf` is written only during init (alloc and band/result slicing) and on teardown, and every launch receives `&s->buf` unmodified. **FIXED** on today's structure rather than by reverting the revert: `adm_hip_upload_buf()` uploads `s->buf` at the end of `adm_hip_init_device()`, the four launches pass `&s->buf_dev`, and the copy is freed in `close()` and on both init failure paths. HIP output is byte-identical at `%.17g` (debug features on) before and after on 8/10/12/16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. On gfx1036 each kernel's argument segment shrinks by 320 bytes (e.g. 968 to 648 on `adm_cm_line_kernel_8`); scratch and VGPRs are unchanged by the pointer. The 936-byte scratch on `adm_cm_line_kernel_8` is VGPR spilling (239 spills at the 128-register cap), not the struct. 1080p 10-bit end-to-end fps is unchanged within noise (median 29.5 before, 30.0 after over 10 alternating runs, spread 27.4 to 31.1). `adm_csf.hip` goes from 35 clang-tidy findings to 0; that restructure lowers the two CSF kernels from 44/53 to 38/45 VGPRs. The CUDA twin passes `AdmBufferCuda` by value, contrary to ADR-0759 and Research-0759. | [ADR-0759](adr/0759-hip-adm-buffer-by-pointer.md), [Research-0759](research/research-0759-hip-adm-buffer-by-pointer.md) | `perf/hip-adm-buffer-by-pointer` | 2026-09-18 | closed | | **T-GPU-ADM-DWT2-16BIT-INT32-OVERFLOW-2026-09-18** | **The CUDA, HIP and Metal integer ADM twins summed the 16-bit vertical DWT response in int32 (undefined behaviour).** The GPU half of `T-ADM-DWT2-16BIT-INT32-OVERFLOW-2026-09-18`: the scale-0 vertical pass forms `sum(filter[k] * s[k])` before subtracting `46342 * 2^(bpc - 1)`, and with 16-bit samples the int32 partial sum passes `INT32_MAX` once three samples reach 42456. Scores agreed with the CPU only because the devices wrap. The SYCL twin already formed the sum in int64. **FIXED**: the CUDA and HIP fused kernel (`adm_dwt2.cu`, `adm_dwt2.hip`) accumulates in `DwtVertAccum::type`, int64 for `uint16_t` input and int32 for 8-bit, and the Metal raw vertical kernel sums in `long`. CUDA and HIP output is byte-identical at `%.17g` before and after on 8, 10, 12 and 16 bpc inputs, and CUDA 1080p 10-bit throughput is unchanged (median 443.8 fps after, 432.8 before, run-to-run spread about 25). Metal is changed by construction and unverified on Apple silicon. The same PR takes both kernel files from 14 clang-tidy findings to 0. New `test_gpu_adm_bright_16bit_parity` scores bright 16-bit noise on each twin against the scalar CPU; it passes on CUDA, HIP and SYCL. | [Research-2063](research/2063-upstream-sync-2026-09-adm-vif-simd.md) | `fix/gpu-adm-dwt2-16bit-overflow` | 2026-09-18 | closed | | **T-UPSTREAM-03B5562C5-ADM-DECOUPLE-AVX2-OVERSHOOT-2026-09-18** | `adm_decouple_avx2` bounded its 8-wide loop with `right - (right % 8)` although the loop starts at `left`, so the last store ran up to seven int16 columns past `right` (Netflix/vmaf `03b5562c5`). Measured with the fork's band stride: 372 of 992 band widths (9..1000) wrote past `right`, and band widths 32 and 40 (frame widths 63-64, 79-80) wrote one or two samples past the row into the next row or, on the last row, into the next slab band. No consumer reads those samples, so no score moved: region outputs are identical and 3,072 end-to-end AVX2 scores (w 17..400, h 36, noise and smooth) are bit-identical before and after. **FIXED** with upstream's `right - ((right - left) % 8)` (`core/src/feature/x86/adm_avx2.c:805`, the form `adm_decouple_s123_avx2` and `adm_decouple_avx512` already used). `test_integer_adm_simd` now sweeps band widths 8..80 over five heights with guard-filled outputs; with the old bound it reports 48 samples written outside the region. | [Research-2063](research/2063-upstream-sync-2026-09-adm-vif-simd.md) | `port/upstream-2026-09` | 2026-09-18 | closed | | **T-UPSTREAM-EA012E387-ADM-DWT2-NEON-OOB-WRITE-2026-09-18** | `adm_dwt2_8_neon`'s horizontal 8-wide loop ran to `half_w` with no tail (Netflix/vmaf `ea012e387`), so on every dispatched width it stored one sample at column `stride`. That went into the next row, which that row then rewrote, and on the last row into element `[0][0]` of the next band in the ADM slab, which was already written. The spilled value came from lanes over-read out of `tmp_ref`. On Linux AArch64 this changed `integer_adm_scale0` for frame widths 24, 32 and 40 at heights below 50 (e.g. 24x20 scalar 0.47136 vs NEON 0.50793), and repeated NEON runs on 32x48 disagreed. **FIXED**: the loop stops at the fork's guarded DWT2 bound `half_w - 1 - ((half_w - 2) % 8)`, and an `ind_x`-driven scalar tail covers the rest, including the mirrored last column that the old redo block recomputed. The kernel was split into row helpers for readability-function-size. NEON is bit-exact with scalar across widths 16..128 (step 8) at six heights plus 576x32. 576x324 and widths 48..200 at h 72 score identically before and after under QEMU. `test_adm_dwt2_neon` now uses the real band stride and a guard band; with the old bound it reports samples written outside the band. | [ADR-1257](adr/1257-retire-darwin-adm-dwt2-legacy-dispatch.md), [Research-2063](research/2063-upstream-sync-2026-09-adm-vif-simd.md) | `port/upstream-2026-09` | 2026-09-18 | closed | diff --git a/scripts/ci/tidy-baseline-hip.json b/scripts/ci/tidy-baseline-hip.json index d063bf03a..0ee40c205 100644 --- a/scripts/ci/tidy-baseline-hip.json +++ b/scripts/ci/tidy-baseline-hip.json @@ -4,7 +4,7 @@ "generator": "scripts/ci/tidy-ratchet.py", "clang_tidy_version": "22.1.8", "tus": 50, - "total_warnings": 1005, + "total_warnings": 970, "total_nolint_uncited": 16, "warnings": { "core/src/feature/hip/ciede_hip.c": 13, @@ -21,7 +21,6 @@ "core/src/feature/hip/float_vif/float_vif_score.hip": 21, "core/src/feature/hip/float_vif_hip.c": 29, "core/src/feature/hip/hip_hsaco_stubs.c": 1, - "core/src/feature/hip/integer_adm/adm_csf.hip": 35, "core/src/feature/hip/integer_adm/adm_csf_den.hip": 15, "core/src/feature/hip/integer_adm/adm_decouple.hip": 49, "core/src/feature/hip/integer_adm/adm_decouple_inline.hip": 31, @@ -104,6 +103,21 @@ } }, "previous_baseline_sha256": "990b0aad4fedf3e3adc1d834e7d80dad56bb06294cb0c12cffa2a65b2b6ae5ea" + }, + { + "sources": [ + "core/src/feature/hip/integer_adm/adm_csf.hip" + ], + "clang_tidy_version": "22.1.8", + "changes": { + "core/src/feature/hip/integer_adm/adm_csf.hip": { + "warnings": [ + 35, + 0 + ] + } + }, + "previous_baseline_sha256": "c90b1a4ce1da9eda866956484bbf885395ad5dc3f1b7448f29f10bc8b410321c" } ] }