ROCm - OpenCL alternatives: Exercises — Hardware deep dive

Before we start, the one formula that half these problems lean on, restated in plain words:

The picture behind that formula, since three exercises depend on it:

Level 1 — Recognition

Exercise 1.1 (L1)

Match each OpenCL memory space to the AMD hardware it lands on: __global, __local, __private.

Recall Solution

__global → VRAM (the GPU's big off-chip memory, gigabytes, slow-ish).
__local → LDS (Local Data Share, on-chip scratchpad shared by one work-group).
__private → VGPRs (Vector General-Purpose Registers, one item's private scratch).

Why this ordering? It follows the distance-from-the-core rule: private is closest and fastest, local is shared-but-still-on-chip, global is farthest and largest. Same idea NVIDIA calls registers / shared memory / global memory.

Exercise 1.2 (L1)

Which platform is which? Pair the tool to its owner and its purpose: HIP, OpenCL, HSA.

Recall Solution

HIP — AMD's CUDA-like API; write CUDA-shaped code, compile to AMD or NVIDIA.
OpenCL — Khronos Group's cross-vendor standard; runs on GPU, CPU, FPGA, DSP.
HSA — Heterogeneous System Architecture; the low-level layer under ROCm giving direct memory access between host and device.

Level 2 — Application

Exercise 2.1 (L2)

You launch an OpenCL kernel over $G_{x} = 1, 048, 576$ elements with a local size $L_{x} = 64$ . How many work-groups are created along $x$ ?

Recall Solution

Number of groups $= G_{x} / L_{x} = 1, 048, 576/64 = 16, 384$ .

Why divide? Each group is a fixed-size batch of 64. Splitting a line of $1, 048, 576$ items into batches of 64 is exactly integer division — count how many full batches fit.

Exercise 2.2 (L2)

A work-item reports group_id = 65, local_id = 0, local_size = 64. What is its global id?

Recall Solution

$global_id = 65 \times 64 + 0 = 4160 .$ This is the worked example from the workhorse formula, run forward.

Exercise 2.3 (L2)

Reverse it. A work-item's global id is $4, 160$ , with local_size = 64. Find its group_id and local_id.

Recall Solution

To undo group_id × local_size + local_id, we split $4160$ into how-many-full-64s and the leftover:

group_id $= ⌊ 4160/64 ⌋ = 65$ (integer part of the division).
local_id $= 4160 mod 64 = 0$ (the remainder).

Why floor-and-remainder? Division-with-remainder is the inverse of "multiply then add a leftover smaller than the divisor". Since $0 \leq local_id < 64$ , this split is unique — there's only one way to write $4160$ as $(groups) \times 64 + (leftover < 64)$ .

Level 3 — Analysis

Exercise 3.1 (L3)

You process $N = 5000$ elements. Local size is fixed at $L_{x} = 256$ . (a) How many work-groups do you launch to cover all elements? (b) How many total work-items run? (c) How many of them do useless work?

Recall Solution

(a) Round up: $⌈ 5000/256 ⌉ = ⌈ 19.53 \dots ⌉ = 20$ groups. (b) Total items $= 20 \times 256 = 5120$ . (c) Wasted items $= 5120 - 5000 = 120$ . These have global_id from $5000$ to $5119$ ; the kernel guard if (i < 5000) makes them exit immediately.

Why we accept waste: the hardware executes in fixed-width batches (AMD wavefront = 64, and 256 = four wavefronts). You cannot launch a partial group, so the last group is always fully populated even if it overhangs the data. The overhang is the price of tidy, uniform batches.

Exercise 3.2 (L3)

An AMD GPU has a wavefront of 64 threads; an NVIDIA GPU has a warp of 32. You choose local_size = 100 for portability. Explain, per vendor, why this is a poor choice and give a better number.

Recall Solution

A wavefront/warp is the smallest indivisible batch the hardware physically issues. If local_size is not a whole-number multiple of it, the last batch runs partly empty — those lanes are switched off but still occupy a slot.

On AMD (64): $100 = 64 + 36$ . The second wavefront runs with only 36 of 64 lanes active → $28/64 \approx 44%$ of that wavefront wasted.
On NVIDIA (32): $100 = 3 \times 32 + 4 = 96 + 4$ . The fourth warp runs 4 of 32 lanes → $28/32 \approx 88%$ of that warp wasted.

Better: pick a multiple of both, i.e. a multiple of $lcm (32, 64) = 64$ . Choosing $128$ (or 256) gives full wavefronts on AMD ( $128 = 2 \times 64$ ) and full warps on NVIDIA ( $128 = 4 \times 32$ ) — zero idle lanes on either vendor. This is why portable code favours powers of two ≥ 64.

Level 4 — Synthesis

Exercise 4.1 (L4)

You have a CUDA saxpy kernel (computes $y [i] = a \cdot x [i] + y [i]$ ). Port it to run on both AMD and NVIDIA with a single source. List the concrete API substitutions and the compilation targets, and state why the kernel body stays byte-for-byte identical.

Recall Solution

Use HIP. The kernel body is unchanged because HIP defines the same built-ins:

__global__ void saxpy(float a, float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

Host-side API swaps (CUDA → HIP), a near 1:1 rename:

cudaMalloc → hipMalloc
cudaMemcpy → hipMemcpy
cudaMemcpyHostToDevice → hipMemcpyHostToDevice
cudaFree → hipFree
kernel<<<blocks,threads>>>(...) → hipLaunchKernelGGL(kernel, blocks, threads, 0, 0, ...)

Compilation targets (both via HIP-Clang, not NVCC):

AMD path: HIP source → HIP-Clang → AMDGPU ISA → ROCm binary
NVIDIA path: HIP source → HIP-Clang → NVPTX → NVIDIA binary

Why the body is identical: blockIdx, blockDim, threadIdx are just names that each backend maps to its own intrinsics (workgroup/workitem IDs on AMD, native CUDA registers on NVIDIA). The arithmetic — index math and the a*x+y — is plain C++ that compiles to whatever machine has. Only the vendor-specific plumbing (allocation, launch) needed renaming.

Exercise 4.2 (L4)

Take the same vector-add task and count the explicit setup calls OpenCL forces you to make that CUDA/HIP hide. Then explain the trade-off in one sentence.

Recall Solution

OpenCL exposes the whole pipeline. Explicit stages (from the parent's example):

clGetPlatformIDs — pick a platform.
clGetDeviceIDs — pick a device.
clCreateContext — a container for that device.
clCreateCommandQueue — a queue to submit work.
clCreateBuffer — device memory.
clCreateProgramWithSource + clBuildProgram — JIT-compile the kernel at runtime.
clCreateKernel — get a handle.
clSetKernelArg — bind each argument by hand.
clEnqueueWriteBuffer / clEnqueueNDRangeKernel — transfer and launch.
clReleaseMemObject / clReleaseKernel / clReleaseCommandQueue — tear down.

That's ≈10 stages CUDA compresses into cudaMalloc / cudaMemcpy / <<<>>> / cudaFree.

Trade-off (one sentence): OpenCL's verbosity buys you runtime portability across any vendor's device at the cost of boilerplate, whereas CUDA's brevity buys convenience at the cost of NVIDIA lock-in.

Level 5 — Mastery

Exercise 5.1 (L5)

Design decision. You maintain a 200,000-line CUDA physics codebase. Management wants it to also run on AMD MI300 GPUs, ships only on Linux, and cannot afford a full rewrite. Choose ROCm/HIP or OpenCL, and justify with three concrete facts from the platform models. Then name the single biggest technical risk of your choice.

Recall Solution

Choose ROCm / HIP. Justification:

Automatic translation: HIP-ify tooling converts CUDA to HIP with ~80–95% done automatically, so a 200k-line base needs edits mostly on the ~5–20% edge cases — not a rewrite. OpenCL would demand rewriting every kernel and all host boilerplate from scratch.
Syntax parity: HIP mirrors CUDA built-ins (blockIdx, threadIdx) and the runtime API (hipMalloc ↔ cudaMalloc), so developer knowledge transfers directly. OpenCL's platform/context/queue model is a different mental model.
Target match: ROCm is Linux-first and MI300 is a first-class ROCm target — exactly the deployment constraint given. And HIP still compiles back to NVIDIA via NVPTX, so you keep the existing NVIDIA fleet.

Biggest technical risk: the non-translatable slice. The ~5–20% that HIP-ify cannot auto-convert is usually the performance-critical, hand-tuned code — warp-level intrinsics, PTX inline assembly, wavefront-size assumptions (NVIDIA warp 32 vs AMD wavefront 64). Any kernel that hard-codes "32" for a shuffle or reduction will produce wrong results silently on a 64-wide wavefront until it's audited by hand.

Exercise 5.2 (L5)

Numeric capstone. Your MI300 kernel needs each work-item to hold a $16 \times 16$ tile of floats in local memory for a matmul. LDS per compute unit is 64 KB. (a) How many bytes of LDS does one work-group need if it has 256 items each keeping its own tile? (b) Does it fit? (c) If not, what is the largest local size that fits, rounded down to a multiple of the 64-wide wavefront?

Recall Solution

One tile $= 16 \times 16 = 256$ floats $\times 4$ bytes $= 1024$ bytes per item. (a) $256 items \times 1024 B = 262, 144 B = 256 KB$ . (b) LDS is only $64 KB = 65, 536$ B, so $256 KB > 64 KB$ → does not fit (needs 4× the budget). (c) Max items $= ⌊ 65, 536/1024 ⌋ = 64$ items. That is already exactly $1 \times 64$ , a whole wavefront, so no rounding needed → local_size = 64.

Why this is the real ceiling: occupancy is bounded by the scarcest on-chip resource. Here LDS, not thread count, is the wall — a beautiful power-of-two 256 group is impossible not because the scheduler dislikes it but because the physics of 64 KB scratchpad forbids it.

Recall Quick self-check clozes

HIP compiles through HIP-Clang (not NVCC), emitting NVPTX for NVIDIA and AMDGPU ISA for AMD. An AMD wavefront is 64 threads; an NVIDIA warp is 32 threads. The global index formula is global_id = group_id × local_size + local_id. What undoes that formula? ::: Integer division: group_id = floor(global_id / local_size), local_id = global_id mod local_size. Why round the group count UP? ::: A partial work-group cannot launch, so you cover the overhang and bounds-check with if (i < n).