hip-kernel-optimization

HIP Kernel Optimization

Purpose

Provide ready patterns for efficient HIP kernels and guide diagnosis of memory throughput, occupancy, and synchronization bottlenecks.

When to Use

• Implementing or reviewing HIP kernels for AMD MI/CDNA architectures or CUDA-portable code
• Porting CUDA code to HIP while retaining performance
• Preparing profiling runs with rocprof

Optimization Priority

Phase 1: Low-hanging fruit (try first, low risk)

1. #pragma unroll on hot loops with small, fixed trip counts
2. Enable -ffast-math compiler flag for floating-point kernels
3. Use 32B vectorized loads/stores instead of 16B
4. Add __launch_bounds__(maxThreads, minBlocks) to guarantee occupancy
5. Add const qualifiers on read-only pointers
6. Verify memory coalescing (consecutive threads → consecutive addresses)

Phase 2: Targeted improvements (profile first) 7. Profile with rocprof to confirm bottleneck 8. If memory-bound: CK-Tile buffer views with vectorization 9. If compute-bound: Shared memory tiling 10. Dynamically calculate block size based on problem dimensions 11. Replace large 2D shared arrays with atomicAdd for sparse patterns 12. Provide multiple block size configurations to avoid register spill 13. Add explicit rounding mode control for numerical correctness 14. Pre-compute workspace size to avoid dynamic allocation 15. Implement CSV-based tuning cache for repeated GEMM shapes

Phase 3: Complex transformations (high effort) 16. Algorithm changes (e.g., Top-K-only softmax) 17. gfx950: Use 16x16x32 MFMA instead of 2x 16x16x16 18. Kernel fusion (multi-op in single kernel) 19. Persistent kernels for repeatedly executed operations 20. Shape-based heuristic dispatching

Anti-patterns:

• Optimizing everything at once
• Manual loop unrolling (use #pragma unroll instead)
• Over-unrolling (factor > 8)
• Premature vectorization without alignment check
• Unnecessary buffer coherence flags (e.g., glc)

Core Optimization Patterns

1. Memory Access

• Coalescing: Map consecutive threads to consecutive addresses; prefer SoA over AoS
• Vectorization: Use CK-Tile buffer views for efficient I/O; prefer 32B loads over 16B
• Boundary handling: Separate fast vectorized path from slow boundary path

  if(idx + VEC_SIZE <= d) {
      vec_o out_vec;
      #pragma unroll
      for(size_t j = 0; j < VEC_SIZE; j++) {
          out_vec[j] = compute(x[j], y[j]);
      }
      buffer_out.template set(idx, 0, true, out_vec);  // Fast path
  } else {
      for(size_t j = 0; j < VEC_SIZE; j++) {          // Boundary path
          if(idx + j < d) ptr_out[idx + j] = compute(...);
      }
  }