Porting CUDA Kernels to HIP for Peak AMD Performance

• [How CUDA Patterns Map to HIP: Common Language and API Differences]
• [Avoiding Memory-Access Pitfalls: Memory Model, Synchronization, and Thread Mapping]
• [Squeezing RDNA/GCN: Performance-Tuning Techniques for AMD GPUs]
• [Practical Toolchain: hipify, rocprof, and Debugging Workflows]
• [Validation and Benchmarks: Platform-Specific Pitfalls and What to Watch]
• [Practical Porting Checklist — Step-by-Step Protocol]

Porting CUDA kernels to HIP is usually quick at the surface level, but the real work starts when you re-optimize for AMD silicon: wavefront width, register pressure, and the memory hierarchy determine whether a port will merely run or actually perform. Treat the port as a hardware-aware re-architecture rather than a pure mechanical translation.

Your build completes, tests pass, and yet your kernels throughput lags the reference — low GPU utilization, long stall times in the memory unit, and kernel runtimes that don’t improve despite obvious CPU-side tweaks. That is the symptom set this guide addresses: the port is functionally correct but misaligned with AMD execution and memory primitives, which means profiling, targeted rewrites, and platform-aware compile options are the only path to peak performance.

How CUDA Patterns Map to HIP: Common Language and API Differences

Keep the first rule simple: hip is a portability layer and a language dialect — it maps a large portion of CUDA’s runtime and kernel syntax, but small differences matter for correctness and for performance.

• Use hipify-clang/hipify-perl to translate code as a first pass. hipify-clang parses CUDA into an AST and does the safest translation for complex code; hipify-perl is faster for trivial replacements but less robust for templates and macros. Use the clangen-based tool as your baseline for non-trivial code.

• Kernel-launch mapping:

  • HIP supports the <<<>>> syntax and hipLaunchKernelGGL. When HIP uses hipLaunchKernelGGL, the macro requires the first five launcher parameters: kernelName, gridDim, blockDim, dynamicShared, stream. That difference matters when you rely on optional <<<...>>> args in CUDA. HIP_KERNEL_NAME wrappers may be injected by hipify for templated kernels.

Example — minimal CUDA → HIP translation (before / after):

// CUDA
__global__ void saxpy(float a, const float *x, float *y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}
cudaMalloc(&d_x, n*sizeof(float));
cudaMemcpy(d_x, h_x, n*sizeof(float), cudaMemcpyHostToDevice);
saxpy<<<(n+255)/256, 256>>>(a, d_x, d_y, n);
cudaDeviceSynchronize();

// HIP
#include <hip/hip_runtime.h>
__global__ void saxpy(float a, const float *x, float *y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}
hipMalloc(&d_x, n*sizeof(float));
hipMemcpy(d_x, h_x, n*sizeof(float), hipMemcpyHostToDevice);
hipLaunchKernelGGL(saxpy, dim3((n+255)/256), dim3(256), 0, 0, a, d_x, d_y, n);
hipDeviceSynchronize();

API mapping cheat-sheet (common items):

CUDA	HIP	Notes
cudaMalloc	hipMalloc	Same semantics; check return hipError_t
cudaFree	hipFree	—
cudaMemcpy	hipMemcpy	Same direction enums map (hipMemcpyHostToDevice)
cudaMemcpyAsync	hipMemcpyAsync	Same stream semantics
cudaStream_t	hipStream_t	Replace directly
cudaGetLastError()	hipGetLastError()	HIP semantics differ — check immediately after launch.
cuBLAS	rocBLAS/hipBLAS	Library mappings exist; see porting guide.

Practical notes:

• Dynamic parallelism (device-launched kernels) is not supported in HIP on many targets — plan to flatten control where present.
• Avoid assuming CUDA's behavior for cudaGetLastError — hipGetLastError may only reflect the immediately preceding runtime call; therefore call and check it right after launches during debug.

Avoiding Memory-Access Pitfalls: Memory Model, Synchronization, and Thread Mapping

Memory-bound kernels fail on AMD for different reasons than they fail on NVIDIA. Pay attention to access patterns, on-chip scratch (LDS), and wavefront behavior.

• Architecture reality check: AMD hardware exposes differing wavefront sizes (the unit analogous to CUDA’s warp). Older GCN targets use wave64; RDNA and newer GPUs frequently use a native wave32 execution but many devices support 32 or 64; you cannot assume warpSize == 32. Test the device and write lanes generically. Hardware specs and per‑GPU wave sizes are documented in ROCm device tables.

• Unified/managed memory is supported on many AMD product lines (Vega and later), but behavior depends on kernel-mode driver and HMM/XNACK configuration. Use hipMallocManaged() only after checking hipDeviceAttributeManagedMemory, and set HSA_XNACK=1 for system-allocator-managed unified memory where required. Treat page-migration behavior as an explicit test case rather than a drop-in replacement.

Code snippet to detect managed-memory support:

int managed = 0;
hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, device_id);
if (managed) {
  hipMallocManaged(&ptr, N * sizeof(float));
}

• Synchronization and warp/wave intrinsics:

  • __syncthreads() exists and behaves as expected for block-level barriers.
  • Cross-lane intrinsics (shuffle, ballot, vote) exist in HIP, but __ballot returns a 64-bit mask on AMD; do not assume a 32-bit result. Prefer warpSize-aware code and test hasWarpShuffle/hasWarpBallot device properties during runtime guard.

• Fences and cache control:

  • __threadfence_system semantics differ and may not flush L2 the same way on all ROCm toolchains. The porting guide warns that threadfence_system functionality may be unavailable; workarounds (like HSA_DISABLE_CACHE=1) exist but carry costs. Profile before and after any such global cache-control changes.

Important: During port debugging call hipGetLastError() immediately after kernel launches; the semantics differ from cudaGetLastError() and failing to check it timely will hide launch-time errors.

Squeezing RDNA/GCN: Performance-Tuning Techniques for AMD GPUs

Getting the last 10–50% is where you earn your cred as a kernel engineer. AMD throughput depends on how you feed the vector ALUs across wavefronts and how you manage per-wave registers and LDS.

• Start from the hardware constraints:

  • Wavefront width (32/64) controls how many lanes must be busy to avoid serializing divergent work. Choose block sizes that are multiples of the native wave width when possible.
  • VGPR (vector GPR) and SGPR pressure limit concurrent waves per CU; excessive per-thread registers reduce occupancy. Use compiler feedback and rocprof to see active wave counts.

• Compiler flags that help tuning:

  • Use hipcc --offload-arch=gfx90a (or the target gfx value for your GPU family) to generate code for the right GPU, and iterate with -O2/-O3. hipcc is a wrapper around HIP-Clang/amdclang and accepts --offload-arch.
  • On RDNA you may toggle -mwavefrontsize64 / -mno-wavefrontsize64 to select wave64 vs wave32 for codegen experiments, and -mcumode to test CU vs WGP scheduling modes where available. Use these flags to experiment and re-profile.

• Practical tuning levers (ordered by expected impact):

  1. Memory-layout and alignment — convert AoS to SoA for vector math, pack loads into vector types (e.g., float4) where you can, and ensure contiguous accesses across lanes. Avoid strided per-lane access patterns that break cache-line locality.
  2. Stage data into LDS (HIP __shared__) for multi-lane reuse — tile-based GEMM and convolution benefit heavily from careful LDS tiling.
  3. Reduce register pressure — hoist temporaries to shared memory when that reduces per-thread VGPRs enough to increase active waves per CU.
  4. Prefer compute-friendly intrinsics — use __shfl*/__ballot-style operations for reductions and scans inside a wave to avoid global atomics.
  5. Micro-benchmark — single-kernel microbenchmarks help isolate memory vs ALU bottlenecks; use rocprof counters to measure MemUnitStalled and VALUInsts.

• Watch for platform-specific throughput quirks:

  • RDNA’s SIMD32 execution sometimes makes fewer registers per wave preferable compared with legacy wave64 code patterns; rebalancing work per thread (more work per thread, fewer threads per block) can help with fewer waves but higher per-thread throughput.

Practical Toolchain: hipify, rocprof, and Debugging Workflows

A pragmatic toolchain and repeatable profiling loop will save you weeks of guesswork.

1. hipify: automatic port
   • Use hipify-clang as the default porting tool; run it with a compile_commands.json so translation understands your build flags and include paths. Use --print-stats to see what translated cleanly and what needs manual attention.

Example:

   hipify-clang -p build/compile_commands.json src/module.cu -o src/module.hip.cpp --print-stats

1. Build with hipcc / amdclang:
   • For AMD targets prefer hipcc (wrapper) or invoke amdclang++ directly to get fine-grained flags. Always set an explicit target: --offload-arch=gfx90a (or gfx1030, gfx1100, …). Use -O3 for production runs and keep -g -O0 for debugs.

Example:

   hipcc -std=c++17 --offload-arch=gfx90a -O3 -o myapp module.hip.cpp

To test RDNA32 vs RDNA64 codegen:

   hipcc -O3 --offload-arch=gfx1030 -mno-wavefrontsize64 -o myapp32 module.hip.cpp
   hipcc -O3 --offload-arch=gfx1030 -mwavefrontsize64 -o myapp64 module.hip.cpp

1. Profile with rocprof:
   • Use rocprof --stats or --hip-trace to collect kernel timings and activity. For counter-based profiling use an input file describing the pmc counters to collect. Outputs include results.stats.csv and trace JSONs you can visualize.

Example:

   # input.txt: a small list of perf counters
   rocprof -i input.txt ./myapp
   rocprof --stats --hip-trace ./myapp     # quick overview traces and CSVs

rocprof outputs results.stats.csv (per-kernel durations and averages) and results.hip_stats.csv (HIP runtime API stats). Use those to find hot kernels and disproportionate memcpy time.

1. Debug with ROCgdb:
   • For source-level GPU stepping and register dumps use rocgdb. It mimics gdb and supports dumping wavefront registers (info registers) and stepping into device code on supported platforms. Run on a node with ROCm installed; ensure any SELinux/containers are configured so ROCgdb has device access.

Example:

   rocgdb ./myapp
   (gdb) break main
   (gdb) run
   (gdb) info registers   # dumps wavefront registers

1. Iterate: edit → build → profile → measure. Use the profiler CSVs as your source of truth and limit changes to one knob at a time.

Validation and Benchmarks: Platform-Specific Pitfalls and What to Watch

Validation and benchmarking are a discipline: functional correctness first, then microbench correctness, then performance budgets.

• Library mapping and numerical parity:

  • Replace CUDA libraries with their ROCm counterparts: cuBLAS → rocBLAS (or hipBLAS wrapper), cuFFT → rocFFT/hipFFT, cuDNN → MIOpen. HIPIFY automates many calls but validate math results and tolerances (FP32 reductions may differ slightly between implementations).

• Common pitfalls checklist (quick reference):

Symptom	Likely cause	Quick check / fix
Silent kernel failure	hipGetLastError() semantics; error swallowed	Insert if (hipGetLastError() != hipSuccess) { ... } immediately after kernel.
Slow first-run kernel	Managed-memory page faults / migration	Warm pages (prefetch) or use hipMemPrefetchAsync, or enable correct HMM/XNACK settings.
Low occupancy despite many threads	High VGPR/SGPR use or big shared usage	Review compiler feedback, reduce in-kernel temporaries, split kernels.
Inconsistent perf across machines	Offload-arch mismatch or wrong HIP_PLATFORM	Ensure --offload-arch matches device and HIP_PLATFORM=amd is set in CI where required.

• Benchmarking protocol:

  1. Build with -O3 and --offload-arch for the target GPU.
  2. Run microbenchmarks that isolate memory vs compute (e.g., simple vector add / memcpy / GEMM).
  3. Collect rocprof --stats and inspect results.stats.csv for per-kernel average durations and results.hip_stats.csv for host-side API overhead.
  4. Use derived metrics: achieved GB/s (bytes processed / kernel time) and GFLOPS (flops / kernel time) to compare against theoretical bandwidth/compute for the target GPU (found on ROCm spec pages).

• Platform-specific sandboxing:

  • ROCm tools require appropriate kernel modules, /dev/kfd device access, and matching ROCM_PATH/HIP_CLANG_PATH in the environment to produce reliable builds and profiling runs. hipcc and ROCgdb behavior depends on these paths.

Practical Porting Checklist — Step-by-Step Protocol

1. Inventory & baseline:

   • Run your CUDA test-suite and record golden outputs and runtimes on NVIDIA (if available).
   • Add compile_commands.json for your build (CMake: CMAKE_EXPORT_COMPILE_COMMANDS=ON).

2. Automated port:

   • Run hipify-clang with the compile DB and --print-stats. Inspect files for unsupported constructs and missing library mappings.

   hipify-clang -p build/compile_commands.json src/foo.cu -o src/foo.hip.cpp --print-stats

1. Manual fixes:

   • Replace driver-API-only uses with runtime equivalents or rework the logic.
   • Swap CUDA-specific libraries to ROCm libs or hip wrappers (check function availability).
   • Fix kernel launch argument order when hipify used hipLaunchKernelGGL incorrectly for templates.

2. Compile & smoke-test:

   • Build with hipcc targeting your GPU:

   hipcc -std=c++17 --offload-arch=gfx90a -O3 -o myapp src/foo.hip.cpp

• For debug builds use -g -O0 so ROCgdb can step into device code.

1. Sanity profiling:

   • Run rocprof --stats to get the first pass timings and CSVs. Identify the top 3 kernels by total time.

2. Micro-optimize kernels:

   • For each hot kernel: reduce register temporaries, stage reused data into __shared__, vectorize loads/stores, and align block/thread sizes to the device wavefront width. Rebuild with -mno-wavefrontsize64 vs -mwavefrontsize64 experiments on RDNA to decide best codegen.

3. Counter-based profiling:

   • Create a rocprof input file listing pmc counters (e.g., MemUnitStalled, VALUInsts) and run rocprof -i counters.txt ./myapp. Inspect input.csv and results.stats.csv to quantify memory stalls vs ALU utilization.

4. Regression & numeric validation:

   • Compare outputs against golden datasets with tolerances. When behavior differs between rocBLAS and cuBLAS, investigate algorithmic differences and test different solver/plan options.

5. CI & packaging:

   • Pin ROCM_PATH and add --offload-arch or GPU_TARGETS settings to your CMake files so build servers produce reproducible binaries. Note GPU_TARGETS is the current recommended CMake variable name for ROCm builds.

6. Finalize:

   • Sweep for error handling: ensure hipGetLastError() checks exist and convert cudaDeviceSynchronize() checks into hipDeviceSynchronize() while checking returned errors.

Sources

HIPIFY: Convert CUDA to Portable C++ Code - Official HIPIFY GitHub repository and documentation; used for guidance on hipify-clang vs hipify-perl and practical hipification workflow.

GPU hardware specifications — ROCm Documentation - Per-GPU tables listing wavefront size, LDS, and cache characteristics; used to pick wave sizes and hardware constraints.

Using rocprof — ROCProfiler Documentation - rocprof usage, trace modes, and output formats (results.stats.csv); used for profiling commands and interpreting CSV outputs.

Unified memory management — HIP Runtime API (HIP docs) - hipMallocManaged, __managed__, and HMM/XNACK behavior and requirements for managed memory on AMD GPUs.

ROCm compiler reference (rocmcc / hipcc) - hipcc/amdclang flags including --offload-arch, -mwavefrontsize64 / -mno-wavefrontsize64, -mcumode, and environment variables affecting compilation.

Using El Capitan Systems: Known Issues — LLNL HPC docs - Practical debugging note: call hipGetLastError() immediately after kernel launches because its semantics differ from cudaGetLastError().

Kernel Language Syntax — HIP Documentation - hipLaunchKernelGGL parameter ordering, kernel qualifiers, and language differences between CUDA and HIP.

Kernel Language Syntax — HIP (intrinsics notes) - Cross-lane intrinsics, __ballot return width, and warp/wave cautions; used for shuffle/ballot semantics.

ROCgdb quick start — ROCgdb Documentation - How to use ROCgdb for heterogeneous (CPU+GPU) debugging, including info registers on wavefronts.

HIP porting guide — HIP Documentation - Library mapping guidance (cuBLAS → rocBLAS/hipBLAS, cuDNN → MIOpen), feature coverage, and portability notes.