Porting CUDA to HIP with hipify

the short version

HIP was designed so that porting from CUDA is mostly a rename. The kernel language is nearly identical, the runtime API mirrors CUDA call for call, and a tool does the mechanical substitution for you. A straightforward codebase can be ported in an afternoon.

What the tool cannot do is notice that your kernel assumed 32 threads per group and is now running on hardware with 64. That assumption is usually implicit, rarely commented, and silently produces wrong answers.

what hipify does

It performs source-to-source translation of the CUDA API surface into HIP's equivalents.

CUDA	HIP
`cudaMalloc`	`hipMalloc`
`cudaMemcpy`	`hipMemcpy`
`cudaStreamCreate`	`hipStreamCreate`
`cudaDeviceSynchronize`	`hipDeviceSynchronize`
`cublasSgemm`	`hipblasSgemm`
`<cuda_runtime.h>`	`<hip/hip_runtime.h>`

The kernel body usually needs nothing at all. __global__, __device__, __shared__, threadIdx, blockIdx, blockDim and __syncthreads are spelled the same in both.

the two tools

hipify-perl kernel.cu -o kernel.cpp | https://rocm.docs.amd.com/projects/HIPIFY/en/latest/ | regex-based translation; no compiler needed, fast, and happy to run on code it cannot fully parse |'pt_perl'

hipify-clang kernel.cu -o kernel.cpp | https://rocm.docs.amd.com/projects/HIPIFY/en/latest/ | clang-based translation; understands the code properly and is more accurate, but needs the CUDA headers present to parse it |'pt_clang'

hipexamine-perl.sh . | https://rocm.docs.amd.com/projects/HIPIFY/en/latest/ | dry run over a tree: reports what would be converted and what it does not recognize, before changing anything |'pt_examine'

Start with the examine script. Knowing how much of your codebase the tool does not recognize is a better estimate of the work than counting lines.

the wavefront trap

This is the one that costs people days. CUDA code frequently assumes a warp is 32 threads, and that assumption is usually invisible: a hardcoded 32, a shuffle-based reduction that halves from 16, a mask of 0xffffffff, or a block size chosen because it was exactly one or two warps.

On AMD CDNA hardware a wavefront is 64. None of the code above is translated, because none of it mentions a CUDA API function. It compiles cleanly and produces wrong results.


// Assumes warpSize == 32. Silently wrong on wave64.
for (int offset = 16; offset > 0; offset /= 2)
    val += __shfl_down_sync(0xffffffff, val, offset);

// Portable: ask the hardware.
for (int offset = warpSize / 2; offset > 0; offset /= 2)
    val += __shfl_down(val, offset);

warpSize is a built-in in both dialects and reports the real group size, so writing against it rather than a literal is the portable habit. Grep a codebase you are about to port for the literal 32, for 0xffffffff, and for any shuffle or ballot intrinsic. That grep is the actual port.

the other three things it cannot do

Libraries without a one-to-one mapping. cuBLAS maps cleanly onto hipBLAS, and cuDNN largely onto MIOpen, but coverage is not total and some functions have no counterpart. Anything Nvidia-specific with no AMD analogue has to be rewritten or dropped.

Warp-level intrinsics with different semantics. Beyond the size issue, the synchronizing variants Nvidia added after Volta do not all have identical AMD equivalents, and ballot returns a 64-bit mask rather than 32-bit. Code that packs a ballot result into an unsigned int loses half of it.

Inline PTX. Any asm volatile block containing PTX is Nvidia machine-specific and must be rewritten, either as portable C++ or as the AMD equivalent.

a workflow that works

Run the examine script to size the job. Convert with hipify, preferring the clang variant if the CUDA headers are available. Build for a real gfx target and fix what the compiler rejects, which is the easy half. Then grep for the wave-size assumptions above and fix those, which is the half that would otherwise have shipped.

Finally, validate numerically against the CUDA build on the same inputs rather than eyeballing output. A port that compiles and runs is not a port that is correct, and reduction order alone will produce small floating point differences that you need to distinguish from real bugs.

keeping one source for both

Once ported, you generally do not want two copies. HIP code compiles for Nvidia as well, with hipcc using nvcc underneath, so a single HIP source tree can target both vendors. That makes HIP the pragmatic choice for anything intended to be portable, at the cost of being one abstraction step away from the newest CUDA-only features.

Where the vendors genuinely differ, isolate it behind a macro or a small platform header rather than scattering conditionals through kernels.

reference

AMD HIPIFY documentation
AMD HIP documentation
NVIDIA CUDA C++ Programming Guide