Portability

Note: AMD GPUs use 64-thread wavefronts (on CDNA) — code assuming warpSize == 32 must be generalized.

HIP (AMD ROCm) near-identical API to CUDA; hipify-perl/hipify-clang convert code; runs on AMD and NVIDIA GPUs SYCL / oneAPI DPC++

Khronos standard C++; SYCLomatic migrates CUDA code; Intel/AMD/NVIDIA backends older cross-vendor standard (C-based kernels) OpenACC / OpenMP offload directive-based (#pragma acc parallel loop, #pragma omp target teams distribute) C++ standard parallelism

std::transform(std::execution::par_unseq,...) compiled with nvc++ -stdpar=gpu Kokkos, RAJA, Alpaka performance-portability C++ libraries (DOE labs) Vulkan compute / WebGPU graphics-API compute shaders Triton, Mojo, Halide, TVM higher-level DSLs/compilers targeting multiple backends SCALE / ZLUDA projects aiming to run CUDA code on non-NVIDIA GPUs

Note: AMD GPUs use 64-thread wavefronts (on CDNA) — code assuming warpSize == 32 must be generalized.

// HIP looks almost identical to CUDA
hipMalloc(&d_a, bytes);
hipMemcpy(d_a, h_a, bytes, hipMemcpyHostToDevice);
hipLaunchKernelGGL(vecAdd, dim3(blocks), dim3(threads), 0, 0, d_a, d_b, d_c, n);   // or <<<>>>

🧪 Part III Exercises

Profile your tiled matmul with Nsight Compute; identify the top stall reason and plot the kernel on the roofline.

Fuse a three-kernel element-wise pipeline into one kernel; measure the speedup and memory traffic reduction.

Implement a reduction with cooperative groups (cg::reduce) and with CUB; compare with your hand-written version.

Use a grid-wide sync (cooperative launch) to implement an iterative Jacobi solver in a single kernel.

Capture 100 small kernel launches in a CUDA graph; compare total time with direct launches.

Replace cudaMalloc/cudaFree in a loop with cudaMallocAsync/cudaFreeAsync; measure the difference.

Implement an FP16 Tensor Core GEMM with WMMA; compare accuracy and speed with an FP32 kernel and cuBLAS.

Sort 100M key-value pairs with Thrust and with CUB DeviceRadixSort; report throughput.

Split a vector reduction across all available GPUs (or simulate with streams) and combine with NCCL AllReduce.

Run compute-sanitizer racecheck on a reduction with a missing __syncthreads(); fix the bug and confirm.