opus-kernel-best-practice

OPUS Kernel Compile-Time Best Practices

Techniques for reducing HIP/C++ kernel compile time when using opus.hpp. These patterns were developed while optimizing a GQA flash attention kernel from 4.8s to 1.5s (70% reduction) in device-only compilation.

Required headers and include paths

For kernel development with OPUS, use these headers from csrc/include/:

• opus/opus.hpp — the OPUS template library + device intrinsic wrappers. This is the only include needed for device code. Provides opus::thread_id_x(), opus::block_id_x(), opus::sync_threads(), opus::warp_all(), etc.
• opus/hip_minimal.hpp — minimal HIP host-side only declarations (dim3, hipMalloc, hipLaunchKernelGGL, etc.). Use on the host pass instead of <hip/hip_runtime.h>.

hipcc my_kernel.cu -I<aiter_root>/csrc/include -D__HIPCC_RTC__ -std=c++20 -O3 --offload-arch=gfx950

HIP runtime	opus:: wrapper	LLVM builtin
threadIdx.x	opus::thread_id_x()	__builtin_amdgcn_workitem_id_x()
blockIdx.x	opus::block_id_x()	__builtin_amdgcn_workgroup_id_x()
blockDim.x	opus::block_size_x()	__builtin_amdgcn_workgroup_size_x()
gridDim.x * blockDim.x	opus::grid_size_x()	__builtin_amdgcn_grid_size_x()
__syncthreads()	opus::sync_threads()	__builtin_amdgcn_s_barrier()
__all(pred)	opus::warp_all(pred)	—

If anything is missing, contact the maintainer ([email protected]) for adding support.

0. Always Separate Device and Host Code (Most Important)

This is the single most impactful technique. hipcc always performs two compilation passes on every .hip/.cu file — one for the host (x86_64) and one for the device (AMDGPU). The heavy opus.hpp template library is only needed on the device side, but without a guard, hipcc parses it on BOTH passes, doubling the frontend cost.

Always structure your kernel files like this:

// my_kernel.cu
#ifdef __HIP_DEVICE_COMPILE__
// ── Device pass: include opus.hpp and define kernels ──
#include "opus/opus.hpp"

__global__ __launch_bounds__(256, 2)
void my_kernel(const float* src, float* dst, int n) {
    // ... opus layout, load, store, MMA, etc.
}

#else
// ── Host pass: minimal declarations + launcher only ──
#include "opus/hip_minimal.hpp"

__global__ void my_kernel(const float* src, float* dst, int n);  // declaration only

extern "C" void run_my_kernel(const void* d_src, void* d_dst, int n) {
    dim3 grid((n + 255) / 256), block(256);
    hipLaunchKernelGGL(my_kernel, grid, block, 0, 0,
                       (const float*)d_src, (float*)d_dst, n);
    hipDeviceSynchronize();
}
#endif