CUDA & HIP

checking your setup

nvidia-smi| https://developer.nvidia.com/system-management-interface | shows driver version, CUDA version, and live GPU utilization/memory on Nvidia hardware |'cuh_chk1'

rocminfo && rocm-smi| https://rocm.docs.amd.com/ | rocminfo lists agents/ISA; rocm-smi shows utilization and memory, the AMD equivalent of nvidia-smi |'cuh_chk2'

nvcc --version| | check the CUDA toolkit compiler version installed, separate from the driver's reported CUDA version |'cuh_chk3'

kernel launch syntax

my_kernel<<<numBlocks, threadsPerBlock>>>(args);| https://docs.nvidia.com/cuda/cuda-c-programming-guide/ | CUDA's triple-chevron launch syntax — grid of blocks, each block a group of threads |'cuh_launch1'

hipLaunchKernelGGL(my_kernel, numBlocks, threadsPerBlock, 0, 0, args);| https://rocm.docs.amd.com/projects/HIP/ | HIP's portable launch macro; modern HIP also allows the same <<<>>> syntax as CUDA when compiled with hipcc |'cuh_launch2'

thread hierarchy (identical names in CUDA and HIP)

int idx = blockIdx.x * blockDim.x + threadIdx.x;| https://docs.nvidia.com/cuda/cuda-c-programming-guide/#thread-hierarchy | the standard 1D global-index formula every kernel starts with — block position * block size + thread position within the block |'cuh_thread1'

HIP deliberately mirrors CUDA's threadIdx/blockIdx/blockDim/gridDim naming so kernel bodies port with near-zero changes — the portability boundary is almost entirely in host-side API calls (memory allocation, launch, stream management), not in the kernel math itself.

compiling

nvcc -O3 -arch=sm_90 kernel.cu -o kernel| https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/ | compile for a specific Nvidia compute capability (sm_90 = Hopper); mismatched -arch silently falls back to slower generic code or fails at runtime |'cuh_comp1'

hipcc -O3 --offload-arch=gfx942 kernel.cpp -o kernel| https://rocm.docs.amd.com/projects/HIP/ | compile for a specific AMD GPU ISA target (gfx942 = MI300 series); find your target with rocminfo |'cuh_comp2'

porting CUDA to HIP

hipify-perl kernel.cu > kernel.cpp| https://rocm.docs.amd.com/projects/HIPIFY/ | automated source translation tool that rewrites CUDA API calls to their HIP equivalents as a starting point, not a guarantee of a finished port |'cuh_port1'

hipify handles the mechanical renames (cudaMalloc → hipMalloc, cudaMemcpy → hipMemcpy, etc). What it can't do automatically: library calls that don't have a 1:1 equivalent, warp-size-dependent code (32 vs 64 threads), and anything relying on Nvidia-specific intrinsics like warp shuffle variants that differ between wave32/wave64.

streams & async work

cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);| https://docs.nvidia.com/cuda/cuda-runtime-api/ | queues a copy on a non-default stream so it can overlap with kernel execution on another stream — the basis of compute/copy overlap |'cuh_stream1'

hipStreamCreate(&stream); hipMemcpyAsync(dst, src, size, hipMemcpyHostToDevice, stream);| https://rocm.docs.amd.com/projects/HIP/ | HIP's equivalent stream API, same overlap model as CUDA streams |'cuh_stream2'

reference

CUDA C++ Programming Guide
HIP Programming Guide
HIPIFY