ROCm - OpenCL alternatives: Question bank — Hardware deep dive

Definition The vocabulary these traps use — read first

Every acronym below is expanded here once, so no reveal surprises you.

kernel: the small function that runs on many GPU lanes at once.
work-item: one lane running that kernel.
work-group: a batch of work-items that run together on one hardware unit and can share fast on-chip memory.
ISA (Instruction Set Architecture): the raw machine language a specific GPU actually executes.
warp: NVIDIA's fixed batch of 32 work-items ("lanes") that execute in lockstep on one hardware unit.
wavefront: AMD's equivalent fixed batch, 64 work-items wide.
SM (Streaming Multiprocessor): one of NVIDIA's physical execution units; it owns a slice of fast shared memory.
CU (Compute Unit): AMD's equivalent physical execution unit; it owns the LDS (Local Data Share) fast memory.
HIP: AMD's CUDA-lookalike API — same built-in names, hip-prefixed functions.
JIT (Just-In-Time compilation): compiling the kernel at run time, on the machine that will run it, for the exact device present — instead of at build time.
PTX (Parallel Thread Execution): NVIDIA's virtual ISA, a portable intermediate step; NVPTX is the compiler backend that emits PTX for NVIDIA GPUs.
NVCC: NVIDIA's traditional CUDA compiler driver (the tool that turns CUDA .cu files into GPU binaries).
NDRange: OpenCL's name for the total grid of work-items, of size $(G_{x}, G_{y}, G_{z})$ .
$G_{x}$ : the global work size — total number of work-items along the x dimension.
$L_{x}$ : the local work size — number of work-items per work-group along x. The number of work-groups is $G_{x} / L_{x}$ .

True or false — justify

Cover the right side, argue your case, then reveal.

OpenCL code compiled once will run unchanged on any GPU brand.

False. OpenCL source is portable, but the kernel is compiled Just-In-Time (JIT) on each machine for that device's ISA — you ship source or intermediate representation, not a single universal binary.

HIP lets you run the same source on both NVIDIA and AMD GPUs.

True. HIP-Clang emits PTX (via the NVPTX backend) for NVIDIA and AMDGPU ISA for AMD from identical source; only the compile target changes, not the code you wrote.

HIP achieves NVIDIA support by calling NVCC (NVIDIA's CUDA compiler driver) under the hood.

False. HIP uses HIP-Clang, which itself generates PTX through its NVPTX backend for NVIDIA targets; NVCC is not required in the modern toolchain.

ROCm and OpenCL exist mainly to make GPUs faster than CUDA.

False. Their purpose is portability and open standards (avoiding vendor lock-in), not raw speed — a native CUDA path can still win on NVIDIA silicon.

A __local variable in OpenCL always maps to the exact same physical memory on every GPU.

False. It maps to whatever fast on-chip shared memory the vendor has — NVIDIA "shared memory" (48–128 KB per SM, the Streaming Multiprocessor) or AMD "LDS" (64 KB per CU, the Compute Unit) — sizes and behaviour differ, that's exactly what the abstraction hides.

If a kernel runs correctly on AMD, the same OpenCL kernel is guaranteed correct on NVIDIA.

False. Different local-memory sizes, wavefront/warp widths (64 vs 32), and rounding can expose bugs (races, over-large work-groups) that only surface on one vendor.

OpenCL can target CPUs and FPGAs, not just GPUs.

True. OpenCL's platform model treats any conforming device as a "compute device," so multicore CPUs, DSPs and FPGAs are valid targets.

HIP-ifying CUDA code is usually a 100% automatic, zero-effort translation.

False. The tooling automates roughly 80–95%; the remaining fraction (vendor-specific intrinsics, tuned assembly, unsupported libraries) needs manual porting.

ROCm has the same first-class Windows support as CUDA.

False. ROCm is Linux-first; Windows support is limited, which is a real deployment constraint versus CUDA.

Spot the error

Each line contains a mistaken statement; the reveal explains what's actually true.

"A work-group of 100 items maps perfectly onto an AMD wavefront."

An AMD wavefront (its fixed lockstep batch) is 64 lanes, so 100 items span two wavefronts with the second half-idle — pick a multiple of the wavefront/warp size for efficiency.

"get_global_id(0) returns the item's position inside its work-group."

No — that's get_local_id(0), which runs 0 to

L_{x} - 1

. get_global_id returns the item's index across the entire NDRange (0 to

G_{x} - 1

"To port CUDA to HIP you must rename threadIdx.x to hipThreadIdx_x."

In modern HIP the built-ins keep the same names (threadIdx.x, blockIdx.x, blockDim.x); it's the API calls that gain the hip prefix (cudaMalloc → hipMalloc).

"OpenCL is short because it hides platform, device, context and queue from you."

The opposite — OpenCL is verbose precisely because it exposes that whole pipeline explicitly, whereas CUDA/HIP hide most of it.

"The number of work-groups is the global size

G_{x}

times the local size

L_{x}

It's the quotient:

Num Groups_{x} = G_{x} / L_{x}

. Multiplying would explode the count instead of chunking the work.

"Since HIP mirrors CUDA, hipMemcpy uses a different argument order than cudaMemcpy."

HIP mirrors the CUDA runtime API nearly 1:1; the argument order matches, only the prefix and copy-direction enum names change (hipMemcpyHostToDevice).

"OpenCL compiles source straight to GPU machine code with no intermediate step."

There is an intermediate stage: source → LLVM IR → device-specific ISA. The IR is what lets one source retarget many devices.

Why questions

Predict the reason, then check.

Why do we bother splitting the NDRange into work-groups at all instead of one flat list of items?

Hardware physically executes items in fixed batches — warps of 32 on an NVIDIA SM, wavefronts of 64 on an AMD CU — and only items in the same group can share fast local memory and synchronize, so groups give the logical structure that matches that physical reality.

Why is JIT (Just-In-Time) compilation central to OpenCL's portability?

Because the driver compiles at run time for the actual hardware present, it can target the exact ISA (RDNA 3, Xe-HPG, etc.) rather than guessing at build time, making one source adapt to any device.

Why does HIP need two separate compilation paths if the source is identical?

Because NVIDIA and AMD have different instruction sets, so code generation must emit different targets (PTX via NVPTX vs AMDGPU) — the divergence is in the backend, not the source.

Why is "vendor lock-in" treated as a problem worth solving rather than just a business detail?

Locking thousands of lines of code to one vendor removes competitive pressure and hardware choice, which raises costs and slows innovation — portability restores the ability to switch.

Why does the same OpenCL kernel not need rewriting when LDS/shared-memory sizes differ across vendors?

The __local qualifier is an abstract request; the runtime maps it to each device's on-chip memory (LDS on a CU, shared memory on an SM), so your code names the space by role, not by a hardware-specific size or address.

Edge cases

The scenarios people forget to test.

What happens if the global work size

G_{x}

is not a whole multiple of the local size

L_{x}

, e.g. 1000 items with groups of 64?

1000/64 is not an integer, so either the launch is rejected or you must pad the NDRange up to a multiple (1024) and add an if (i < n) guard inside the kernel to ignore the extra lanes.

Why does the SAXPY/vector-add kernel include if (i < n) before writing y[i]?

Because rounding the launch up to full work-groups spawns items with indices ==beyond n==; without the guard those extra lanes write past the array, corrupting memory or crashing.

What is the local ID of global item 4160 when the local size

L_{x}

is 64?

Zero — since

4160 = 65 \times 64 + 0

, it sits at the start of group 65, so get_local_id(0) returns 0 while get_group_id(0) returns 65.

Zero-element edge case: what should a well-written OpenCL launch do when

G_{x} = 0

Launch nothing — a global work size of 0 means no work-items, so no groups are dispatched and the kernel body never runs; the host code should skip the enqueue rather than divide by a local size.

What breaks if you request a work-group larger than the device's maximum work-group size, and how do you find that limit?

The kernel launch fails — each device caps work-group size (limited by register/local-memory budget); query it beforehand with clGetKernelWorkGroupInfo(..., CL_KERNEL_WORK_GROUP_SIZE, ...) (or the device-level CL_DEVICE_MAX_WORK_GROUP_SIZE) so an oversized group is rejected at enqueue, not silently split.

Recall Fast self-check

Portable source but per-device binary — which platform? ::: Both OpenCL and HIP: you ship source/IR, the JIT (Just-In-Time) or per-target compile produces the device-specific binary. One-sentence reason HIP source is identical across vendors. ::: HIP defines the same built-ins and API names, and only the compiler backend (PTX via NVPTX vs AMDGPU) differs. Warp vs wavefront width? ::: Warp = 32 lanes (NVIDIA SM); wavefront = 64 lanes (AMD CU). OpenCL query for the largest work-group a kernel allows? ::: CL_KERNEL_WORK_GROUP_SIZE via clGetKernelWorkGroupInfo.