Skip to content

perf(hip): pass AdmBufferHip by pointer in ADM kernels (F3 fix, ADR-0759) #101

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
lusoris merged 1 commit into from
May 29, 2026
Merged

perf(hip): pass AdmBufferHip by pointer in ADM kernels (F3 fix, ADR-0759) #101

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
14 changes: 14 additions & 0 deletions changelog.d/perf/hip-adm-buffer-by-pointer.md
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,14 @@
## HIP ADM: AdmBufferHip passed by pointer (ADR-0759)

`AdmBufferHip` (~272 bytes) was previously passed by value in four HIP `__global__`
kernel signatures in `adm_csf.hip` and `adm_cm.hip`. Each kernel launch marshalled
the full struct through the per-launch argument buffer.

The struct now passes as `const AdmBufferHip * __restrict__ buf_ptr`. A device-side
copy is allocated once at init time and stable for the extractor's lifetime.

Impact: reduced per-launch argument-buffer overhead on all four ADM kernels
(`adm_csf_kernel_1_4`, `i4_adm_csf_kernel_1_4`, `i4_adm_cm_line_kernel`,
`adm_cm_line_kernel_8`) on every AMD GPU target. Numerically transparent.

Mirrors the CUDA F3 fix (PR #93 / PR #96). ADR-0759.
30 changes: 30 additions & 0 deletions core/src/feature/hip/AGENTS.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -262,3 +262,33 @@ preprocessor expands the macro at the point of instantiation (inside
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.
20 changes: 10 additions & 10 deletions core/src/feature/hip/integer_adm/adm_cm.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -183,17 +183,17 @@ extern "C" {
*
* Pure __global__; no template helpers required.
* -------------------------------------------------------------------- */
__global__ void i4_adm_cm_line_kernel(AdmBufferHip buf, int h, int w, int top, int bottom,
__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)
{
(void)csf_a_stride;

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;

Expand Down Expand Up @@ -274,7 +274,7 @@ __global__ void i4_adm_cm_line_kernel(AdmBufferHip buf, int h, int w, int top, i

template <int rows_per_thread>
__device__ __forceinline__ static void
adm_cm_line_kernel_body(AdmBufferHip buf, int h, int w, int top, int bottom, int left, int right,
adm_cm_line_kernel_body(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,
Expand All @@ -287,9 +287,9 @@ adm_cm_line_kernel_body(AdmBufferHip buf, int h, int w, int top, int bottom, int
(void)accum_per_block;
(void)scale;

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;

Expand Down Expand Up @@ -379,7 +379,7 @@ adm_cm_line_kernel_body(AdmBufferHip buf, int h, int w, int top, int bottom, int
atomicAdd((uint64_cu *)&accum_global[band2], (uint64_cu)shifted);
}

extern "C" __global__ void adm_cm_line_kernel_8(AdmBufferHip buf, int h, int w, int top, int bottom, int left,
extern "C" __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,
Expand All @@ -388,7 +388,7 @@ extern "C" __global__ void adm_cm_line_kernel_8(AdmBufferHip buf, int h, int w,
const uint32_t shift_inner_accum,
const uint32_t add_shift_inner_accum)
{
adm_cm_line_kernel_body<8>(buf, h, w, top, bottom, left, right, start_row, end_row, start_col,
adm_cm_line_kernel_body<8>(buf_ptr, h, w, top, bottom, left, right, start_row, end_row, start_col,
end_col, src_stride, csf_a_stride, buffer_h, buffer_stride,
accum_per_block, params, scale, accum_global, ws, shift_inner_accum,
add_shift_inner_accum);
Expand Down
24 changes: 12 additions & 12 deletions core/src/feature/hip/integer_adm/adm_csf.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,15 +56,15 @@ static __device__ __forceinline__ void copy_vec_4(const int16_t *__restrict__ in

/* Scales 1-3 CSF with inline decouple — reads ref/dis DWT2, writes only csf_f */
template <int rows_per_thread, int cols_per_thread>
__device__ __forceinline__ void i4_adm_csf_kernel(AdmBufferHip buf, int scale, int top, int bottom,
__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;
Expand Down Expand Up @@ -125,14 +125,14 @@ __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 <int rows_per_thread, int cols_per_thread>
__device__ __forceinline__ void adm_csf_kernel(AdmBufferHip buf, int top, int bottom, int left,
__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];
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];
int y = top + (blockIdx.y * blockDim.y + threadIdx.y) * rows_per_thread;
int x = left + (blockIdx.x * blockDim.x + threadIdx.x) * cols_per_thread;

Expand Down Expand Up @@ -183,19 +183,19 @@ __device__ __forceinline__ void adm_csf_kernel(AdmBufferHip buf, int top, int bo

#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, \
const AdmBufferHip * __restrict__ buf_ptr, int top, int bottom, int left, int right, int stride, \
AdmFixedParametersHip params) \
{ \
adm_csf_kernel<rows_per_thread, cols_per_thread>(buf, top, bottom, left, right, stride, \
adm_csf_kernel<rows_per_thread, cols_per_thread>(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, \
const AdmBufferHip * __restrict__ buf_ptr, int scale, int top, int bottom, int left, int right, int stride, \
AdmFixedParametersHip params) \
{ \
i4_adm_csf_kernel<rows_per_thread, cols_per_thread>(buf, scale, top, bottom, left, right, \
i4_adm_csf_kernel<rows_per_thread, cols_per_thread>(buf_ptr, scale, top, bottom, left, right, \
stride, params); \
}

Expand Down
70 changes: 51 additions & 19 deletions core/src/feature/hip/integer_adm_hip.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,7 @@
typedef struct AdmStateHip {
size_t integer_stride;
AdmBufferHip buf;
AdmBufferHip *buf_dev; /* device-side copy of buf (ADR-0759: pointer-passing convention) */
bool debug;
double adm_enhn_gain_limit;
double adm_norm_view_dist;
Expand Down Expand Up @@ -557,8 +558,8 @@ static int adm_dwt2_s123_combined_device_hip(AdmStateHip *s, const int32_t *d_i4
return 0;
}

static int adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, int stride,
AdmFixedParametersHip *p, hipStream_t c_stream)
static int adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, AdmBufferHip *buf_dev, int w,
int h, int stride, AdmFixedParametersHip *p, hipStream_t c_stream)
{
for (int band = 0; band < 3; ++band)
assert(((size_t)(buf->csf_f.bands[band]) & 15) == 0);
Expand All @@ -583,16 +584,17 @@ static int adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, i
const int rows_per_thread = 1;
const int BLOCKX = 32, BLOCKY = 4;

void *args[] = {buf, &top, &bottom, &left, &right, &stride, p};
void *args[] = {&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,
(uint32_t)BLOCKY, 1, 0, c_stream, args, NULL);
return hip_rc(rc);
}

static int i4_adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int scale, int w, int h,
int stride, AdmFixedParametersHip *p, hipStream_t c_stream)
static int i4_adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, AdmBufferHip *buf_dev,
int scale, int w, int h, int stride, AdmFixedParametersHip *p,
hipStream_t c_stream)
{
for (int band = 0; band < 3; ++band)
assert(((size_t)(buf->i4_csf_f.bands[band]) & 15) == 0);
Expand All @@ -617,7 +619,7 @@ static int i4_adm_csf_device_hip(AdmStateHip *s, AdmBufferHip *buf, int scale, i
const int rows_per_thread = 1;
const int BLOCKX = 32, BLOCKY = 4;

void *args[] = {buf, &scale, &top, &bottom, &left, &right, &stride, p};
void *args[] = {&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),
Expand Down Expand Up @@ -691,9 +693,9 @@ typedef struct WarpShiftHip {
uint32_t add_shift_sq[3];
} WarpShiftHip;

static int i4_adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, int src_stride,
int csf_a_stride, int scale, AdmFixedParametersHip *p,
hipStream_t c_stream)
static int i4_adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, AdmBufferHip *buf_dev, int w,
int h, int src_stride, int csf_a_stride, int scale,
AdmFixedParametersHip *p, hipStream_t c_stream)
{
int left = (int)(w * (float)(ADM_BORDER_FACTOR)-0.5f);
int top = (int)(h * (float)(ADM_BORDER_FACTOR)-0.5f);
Expand All @@ -712,7 +714,7 @@ static int i4_adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h,
{
const int BLOCKX = 128;
void *args[] = {
buf, &h, &w, &top, &bottom, &left,
&buf_dev, &h, &w, &top, &bottom, &left,
&right, &start_row, &end_row, &start_col, &end_col, &src_stride,
&csf_a_stride, &scale, &buffer_h, &buffer_stride, &buf->tmp_accum, p};
hipError_t rc = hipModuleLaunchKernel(
Expand All @@ -739,8 +741,9 @@ static int i4_adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h,
return 0;
}

static int adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, int w, int h, int src_stride,
int csf_a_stride, AdmFixedParametersHip *p, hipStream_t c_stream)
static int adm_cm_device_hip(AdmStateHip *s, AdmBufferHip *buf, AdmBufferHip *buf_dev, int w, int h,
int src_stride, int csf_a_stride, AdmFixedParametersHip *p,
hipStream_t c_stream)
{
int scale = 0;
int left = (int)(w * (float)(ADM_BORDER_FACTOR)-0.5f);
Expand Down Expand Up @@ -776,7 +779,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, BLOCKY = 4;
void *args[] = {buf,
void *args[] = {&buf_dev,
&h,
&w,
&top,
Expand Down Expand Up @@ -938,11 +941,12 @@ static int integer_compute_adm_hip(AdmStateHip *s, VmafPicture *ref_pic, VmafPic
if (err)
return err;

err = adm_csf_device_hip(s, buf, w, h, (int)buf_stride, &p, s->str);
err = adm_csf_device_hip(s, buf, s->buf_dev, w, h, (int)buf_stride, &p, s->str);
if (err)
return err;

err = adm_cm_device_hip(s, buf, w, h, (int)buf_stride, (int)buf_stride, &p, s->str);
err = adm_cm_device_hip(s, buf, s->buf_dev, w, h, (int)buf_stride, (int)buf_stride, &p,
s->str);
if (err)
return err;
} else {
Expand All @@ -964,12 +968,13 @@ static int integer_compute_adm_hip(AdmStateHip *s, VmafPicture *ref_pic, VmafPic
if (err)
return err;

err = i4_adm_csf_device_hip(s, buf, (int)scale, w, h, (int)buf_stride, &p, s->str);
err = i4_adm_csf_device_hip(s, buf, s->buf_dev, (int)scale, w, h, (int)buf_stride, &p,
s->str);
if (err)
return err;

err = i4_adm_cm_device_hip(s, buf, w, h, (int)buf_stride, (int)buf_stride, (int)scale,
&p, s->str);
err = i4_adm_cm_device_hip(s, buf, s->buf_dev, w, h, (int)buf_stride, (int)buf_stride,
(int)scale, &p, s->str);
if (err)
return err;
}
Expand Down Expand Up @@ -1169,13 +1174,36 @@ static int init_fex_hip(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt,
}
}

/* ADR-0759: allocate and populate device-side copy of AdmBufferHip.
* The CSF and CM kernels take const AdmBufferHip * __restrict__ buf_ptr instead
* of AdmBufferHip by value (~272 bytes), eliminating per-launch argument-buffer
* overhead on every kernel call. The struct contains only device pointers set up
* above; they are stable for the extractor's lifetime, so a single copy suffices. */
hip_err = hipMalloc(&s->buf_dev, sizeof(AdmBufferHip));
if (hip_err != hipSuccess)
goto fail_host; /* buf_dev not yet allocated; jump directly to results_host teardown */
hip_err = hipMemcpy(s->buf_dev, &s->buf, sizeof(AdmBufferHip), hipMemcpyHostToDevice);
if (hip_err != hipSuccess)
goto fail_buf_dev;

s->feature_name_dict =
vmaf_feature_name_dict_from_provided_features(fex->provided_features, fex->options, s);
if (s->feature_name_dict == NULL)
goto fail_host;
goto fail_feature_dict;

return 0;

fail_feature_dict:
/* feature_name_dict failed: buf_dev was already allocated, free it */
(void)hipFree(s->buf_dev);
s->buf_dev = NULL;
/* fall through */
fail_buf_dev:
/* hipMemcpy of buf_dev failed: buf_dev was allocated, free it */
if (s->buf_dev != NULL) {
(void)hipFree(s->buf_dev);
s->buf_dev = NULL;
}
fail_host:
(void)hipHostFree(s->buf.results_host);
s->buf.results_host = NULL;
Expand Down Expand Up @@ -1347,6 +1375,10 @@ static int close_fex_hip(VmafFeatureExtractor *fex)
(void)hipFree(s->buf.data_buf);
s->buf.data_buf = NULL;
}
if (s->buf_dev != NULL) {
(void)hipFree(s->buf_dev);
s->buf_dev = NULL;
}
#endif /* HAVE_HIPCC */

if (s->feature_name_dict != NULL) {
Expand Down
Loading
Loading