Port Cuda Kernel To Hip

Skill: Port a CUDA Kernel to HIP

Purpose

Guide the agent through translating a CUDA kernel to AMD HIP for ROCm-compatible hardware (MI250, MI300, RDNA), preserving correctness and performance intent while adapting to the HIP execution model, memory model, and AMD-specific toolchain.

Use this when

• An existing, correct CUDA kernel needs to run on AMD GPUs (MI250, MI300X, RDNA3) via the ROCm stack.
• Building a codebase that needs to support both NVIDIA and AMD hardware from a single source using HIP's compatibility layer.
• Using hipify-perl or hipify-clang automated conversion and needing to audit and fix what the tool missed or got wrong.
• The target is MI300X and performance parity with the CUDA original is required (not just functional correctness).

Do not use this when

• The kernel uses CUDA features with no HIP equivalent and cannot be rewritten (e.g., NVLink-specific topology, NVIDIA Tensor Core PTX instructions for Hopper). In this case, write a backend-agnostic plan first.
• The kernel is trivially expressed as a higher-level library call (rocBLAS, MIOpen, rocFFT) that already targets AMD platforms. Use the library rather than porting a custom kernel.
• The goal is Triton instead of HIP — Triton's AMD backend (ROCm) is often a better path for compute-intensive kernels than manual HIP porting.

Inputs the agent should gather first

• AMD hardware target: CDNA (MI250X, MI300X) for datacenter inference/training, or RDNA (RX 7900, etc.) for consumer. Wavefront size differs: CDNA and RDNA3 are 64-wide wavefronts by default, not 32-wide warps like CUDA.
• ROCm version: ROCm 5.x vs 6.x — HIP API surface and library availability differs across versions.
• CUDA features used: identify which features require attention — warp primitives (__shfl_*, __ballot_sync, __any_sync, __all_sync), Tensor Core intrinsics (WMMA/WGMMA), cooperative groups, dynamic parallelism, unified memory, NVTX, NCCL.
• Performance requirements: is functional correctness sufficient (pass CI), or must performance match or exceed the CUDA original on comparable hardware?
• Single-source or fork: will the code maintain one source base with #ifdef __HIP_PLATFORM_AMD__ guards, or will there be a separate HIP file?

Required reasoning process

1. Run hipify as a first pass, then review every change. Use hipify-perl or hipify-clang on the source file to perform mechanical name substitutions (cudaMalloc → hipMalloc, cudaDeviceSynchronize → hipDeviceSynchronize, etc.). Do not trust the automated output blindly — review every changed line. The tool handles API names but not semantic differences.

2. Audit wavefront width differences. This is the most common source of correctness bugs in CUDA-to-HIP ports:

   • CUDA warps are 32 threads. AMD wavefronts are 64 threads on CDNA/GCN and 32 threads on RDNA3+ (via wavefront64 mode, but defaults differ).
   • Any code that hardcodes 32 for warp size must be replaced with warpSize (runtime constant) or the architecture-specific constant. Do not hardcode 32 or 64.
   • Warp reduction masks: CUDA uses 32-bit masks (unsigned int) for __ballot_sync, __activemask, etc. HIP uses 64-bit masks (unsigned long long) on 64-wide wavefronts. Mask type and shift values change accordingly.
   • Warp-level primitive sequences: for (int offset = warpSize/2; offset > 0; offset >>= 1) works correctly if warpSize is used.

3. Map CUDA warp primitives to HIP equivalents.

   CUDA	HIP
   __shfl_sync(mask, val, lane)	__shfl(val, lane) (HIP does not take a mask on older ROCm; __shfl_sync exists on newer ROCm ≥ 5.0)
   __ballot_sync(mask, pred)	__ballot(pred) (returns uint64_t on 64-wide wavefront)
   __any_sync(mask, pred)	__any(pred)
   __all_sync(mask, pred)	__all(pred)
   __activemask()	__activemask() (available ROCm ≥ 5.2, returns uint64_t)

   • On ROCm ≥ 5.0, __shfl_sync, __ballot_sync, etc. are available with mask parameters. Prefer these for portability. On older ROCm, use the maskless versions and handle the reduced API.

4. Handle WMMA / Tensor Core differences. CUDA WMMA (nvcuda::wmma) has no direct HIP equivalent for NVIDIA-specific tensor core shapes. AMD equivalent:

   • ROCm exposes matrix fused multiply-add via rocwmma library (ROCm ≥ 5.4) with similar programming model.
   • Alternative: use rocBLAS or hipBLAS for GEMM operations that would otherwise require tensor core code.
   • For FP8 GEMM on MI300X: ROCm provides hipBLASLt with FP8 support. Do not attempt to write raw MFMA (AMD matrix fused multiply-add) intrinsics unless library support is insufficient.
   • If the CUDA kernel uses PTX for WGMMA (Hopper): this has no equivalent on AMD. The port must use rocwmma or library GEMM.

5. Adapt memory model and atomic operations. HIP's memory model is based on the HSA memory model, which differs from CUDA's weakly-ordered GPU memory model in subtle ways:

   • __threadfence() → __threadfence() (supported in HIP)
   • __threadfence_block() → __threadfence_block() (supported)
   • Atomic operations: atomicAdd, atomicCAS, etc. behave similarly. FP16 and BF16 atomics availability differs by GFX architecture — verify for the target.
   • For inter-block synchronization patterns that use global atomics: these are supported but performance characteristics differ. On AMD, global atomic contention can be more severe than on NVIDIA for certain access patterns.

6. Check launch configuration limits. AMD GPU limits:

   • Maximum threads per block: 1024 (same as CUDA).
   • Wavefront (warp) size: 64 on CDNA (MI-series), 32 on RDNA3. blockDim.x should be a multiple of the wavefront size. A blockDim.x = 128 that works well on NVIDIA (4 warps) behaves as 2 wavefronts on CDNA or 4 wavefronts on RDNA3.
   • LDS (shared memory) per block: 64 KB on most CDNA targets (same as CUDA Ampere). Verify for your target.
   • Maximum grid dimensions: generally the same as CUDA.

7. Update device queries and capability checks. Replace:

   • cudaGetDeviceProperties → hipGetDeviceProperties. The HIP_DEVICE_PROP struct fields differ in naming.
   • __CUDA_ARCH__ preprocessor macro → __HIP_DEVICE_COMPILE__ or __gfx90a__ (for specific AMD targets). For portable code: wrap architecture-specific code in #if defined(__HIP_PLATFORM_AMD__) && defined(__gfx90a__).
   • Compute capability checks (if (smVersion >= 80)) → GFX version checks (__HIP_ARCH_GFX90A__ or similar HIP architecture macros).

8. Replace CUDA-specific libraries. Common replacements:

   CUDA library	HIP/ROCm equivalent
   cuBLAS	rocBLAS / hipBLAS
   cuDNN	MIOpen
   NCCL	RCCL
   cuRAND	rocRAND
   cuFFT	rocFFT
   Thrust	rocThrust (hipThrust)
   NVTX	rocTX / roctx
   cuda-memcheck	ROCm debugger (rocgdb) or AddressSanitizer

9. Compile and test iteratively. Build with hipcc targeting the AMD device. For single-source CUDA/HIP code: build with nvcc on NVIDIA and hipcc on AMD. Key flags:

   • hipcc --offload-arch=gfx90a for MI300X.
   • Enable strict error checking: compile with -Wall and treat warnings as errors for the HIP-specific sections.
   • Validate correctness against the CUDA original using the same test inputs before measuring performance.

10. Evaluate performance and close gaps. After correctness is confirmed:

    • Profile with rocprof or Omniperf (AMD equivalent of Nsight Compute).
    • If the HIP kernel is significantly slower than the CUDA original on comparable hardware: investigate LDS (shared memory) bank conflicts (AMD LDS banking differs from CUDA), occupancy (wavefront count per CU vs warp count per SM), and memory access patterns for AMD's NUMA-style HBM topology on MI300X.
    • For MI300X specifically: the unified memory architecture means CPU and GPU share the same HBM pool. hipMalloc and normal host pointers can coexist, but prefetch behavior differs from cudaMallocManaged.

Kernel design rules

• Never hardcode warp size as 32. Always use warpSize or __AMDGCN_WAVEFRONT_SIZE for AMD-specific code.
• Always use the _sync variants of shuffle and ballot primitives when available (ROCm ≥ 5.0) for forward compatibility with AMD's evolving memory model.
• Use hipDeviceProp_t.warpSize at runtime to query the actual wavefront width if the code must work across RDNA (32) and CDNA (64) targets.
• For AMD, prefer block sizes that are multiples of 64 to ensure full wavefront utilization on CDNA targets.
• Use __HIP_PLATFORM_AMD__ and __HIP_PLATFORM_NVCC__ for platform-specific guard macros, not __CUDA__ or CUDA_VERSION.
• Do not assume hipLaunchKernelGGL is necessary — it is the legacy HIP launch syntax. Modern HIP supports the <<<>>> launch syntax.

Correctness requirements

• The HIP kernel must produce output numerically equivalent to the CUDA kernel for the same inputs, within floating-point non-determinism bounds (< 1e-5 relative error for fp32 reductions).
• Warp-level operations with 64-wide wavefronts must produce the same logical result as the 32-wide CUDA original. Reductions that assume 32-thread warp structure MUST be verified.
• Mask types for ballot/activemask operations must be the correct width (64-bit on CDNA). Using 32-bit masks on a 64-wide wavefront silently drops the upper 32 lanes from the operation.
• Memory ordering: any kernel that relies on __threadfence() for inter-block communication must verify this pattern is correctly expressed in HIP.

Performance requirements

• A correctly ported HIP kernel may not match CUDA performance due to architectural differences (LDS banking, wavefront scheduling, cache topology). This is expected and not a bug.
• On MI300X vs H100, FP16 GEMM throughput is roughly comparable. FP8 on MI300X (hipBLASLt FP8) should be compared to H100 FP8 cuBLAS, not to the raw kernel implementation.
• Profile with rocprof --stats as a first pass to identify the dominant HBM bandwidth or compute bottleneck. LDS bandwidth issues on AMD are detected via rocprof's LDS bank conflict counters.
• If performance is significantly below expectations, check: (1) wavefront occupancy per CU, (2) LDS usage per wavefront, (3) vector register spilling, (4) uncoalesced global memory accesses.

Output format

The final response must include:

1. Assessment of CUDA features used: list each CUDA feature and its HIP equivalent or the required design change.
2. Wavefront width impact analysis: identify any code paths that assume a 32-thread warp and describe the required changes.
3. Ported HIP kernel: the complete translated kernel with all changes annotated.
4. Build instructions: hipcc command and offload-arch flag for the target.
5. Correctness validation plan: how to compare HIP and CUDA outputs for the same inputs.
6. Known performance differences: any architectural reasons the HIP kernel may behave differently than the CUDA original.

Common failure modes

• Hardcoded warp size of 32: the most common and impactful CUDA-to-HIP port bug. A warp reduction loop that iterates offset = 16, 8, 4, 2, 1 (for 32-wide) gives wrong results on a 64-wide wavefront because it only reduces half the wavefront.
• 32-bit mask used for 64-wide ballot: unsigned int mask = __ballot_sync(0xffffffff, pred) — on a 64-wide wavefront, the mask is 64 bits and the top 32 threads are dropped from the result.
• WMMA code not ported: nvcuda::wmma code is compiled out silently by hipify without a replacement, leaving the compute path empty or using a fallback that is orders of magnitude slower.
• __CUDA_ARCH__ guards blocking HIP compilation: code guarded with #ifdef __CUDA_ARCH__ is excluded from the HIP device compilation path. Replace with #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) for shared device code.
• Library ABI mismatch: linking against cuBLAS headers but linking the ROCm/HIP binary against a HIP runtime that also pulls in CUDA stubs. Ensure the build system links exactly one GPU backend.
• Cooperative groups not available: some HIP versions have incomplete cooperative groups support. Code that uses cg::grid_group for grid-level synchronization may require a workaround using ROCm's cooperative launch API directly.

Review checklist

• Has warpSize been used instead of the constant 32 everywhere warp width is relevant?
• Have all __ballot_sync return types been changed to uint64_t for 64-wide targets?
• Have all WMMA operations been replaced with rocwmma or library GEMM equivalents?
• Have all CUDA library calls been replaced with their ROCm equivalents?
• Have all __CUDA_ARCH__ guards been updated to include __HIP_DEVICE_COMPILE__?
• Has the kernel been compiled with hipcc --offload-arch=<target> without warnings?
• Has the kernel output been validated against the CUDA reference using the same inputs?
• Has the kernel been profiled with rocprof to identify any AMD-specific bottlenecks?
• For 64-wide wavefront targets: has the block size been verified to be a multiple of 64?
• Is the single-source guard strategy (__HIP_PLATFORM_AMD__ / __HIP_PLATFORM_NVCC__) consistent across all files?