ROCm - OpenCL alternatives: Exercises — Hinglish deep dive

6.2.15 · D4 · Hardware › GPU Architecture › ROCm - OpenCL alternatives

Shuru karne se pehle, woh ek formula jo in mein se aadhe problems mein kaam aata hai, plain words mein:

Us formula ke peeche ki picture, kyunki teen exercises iske upar depend karti hain:

Level 1 — Recognition

Exercise 1.1 (L1)

Har OpenCL memory space ko us AMD hardware se match karo jahan woh land hoti hai: __global, __local, __private.

Recall Solution

__global → VRAM (GPU ki badi off-chip memory, gigabytes, thodi slow).
__local → LDS (Local Data Share, on-chip scratchpad jo ek work-group share karta hai).
__private → VGPRs (Vector General-Purpose Registers, ek item ki private scratch).

Yeh ordering kyun? Yeh distance-from-the-core rule follow karta hai: private sabse close aur fast hai, local shared-but-still-on-chip hai, global sabse door aur bada hai. NVIDIA issi idea ko registers / shared memory / global memory kehta hai.

Exercise 1.2 (L1)

Kaun sa platform kaun sa hai? Tool ko uske owner aur purpose se pair karo: HIP, OpenCL, HSA.

Recall Solution

HIP — AMD ka CUDA-jaisa API; CUDA-shaped code likho, AMD ya NVIDIA ke liye compile karo.
OpenCL — Khronos Group ka cross-vendor standard; GPU, CPU, FPGA, DSP sab par chalta hai.
HSA — Heterogeneous System Architecture; ROCm ke neeche low-level layer jo host aur device ke beech direct memory access deta hai.

Level 2 — Application

Exercise 2.1 (L2)

Aap ek OpenCL kernel $G_{x} = 1, 048, 576$ elements par launch karte ho local size $L_{x} = 64$ ke saath. $x$ direction mein kitne work-groups create honge?

Recall Solution

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

Divide kyun karte hain? Har group 64 ka fixed-size batch hai. $1, 048, 576$ items ki line ko 64 ke batches mein todna exactly integer division hai — count karo kitne full batches fit hote hain.

Exercise 2.2 (L2)

Ek work-item report karta hai group_id = 65, local_id = 0, local_size = 64. Uska global id kya hai?

Recall Solution

$global_id = 65 \times 64 + 0 = 4160 .$ Yeh workhorse formula ka worked example hai, forward run kiya gaya.

Exercise 2.3 (L2)

Isko ulta karo. Ek work-item ka global id $4, 160$ hai, local_size = 64 ke saath. Uska group_id aur local_id nikalo.

Recall Solution

group_id × local_size + local_id ko undo karne ke liye, $4160$ ko kitne-full-64s aur bacha-khucha mein todna hoga:

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

Floor-aur-remainder kyun? Division-with-remainder "multiply then add a leftover smaller than the divisor" ka inverse hai. Kyunki $0 \leq local_id < 64$ , yeh split unique hai — $4160$ ko $(groups) \times 64 + (leftover < 64)$ likhne ka sirf ek hi tarika hai.

Level 3 — Analysis

Exercise 3.1 (L3)

Aap $N = 5000$ elements process kar rahe ho. Local size $L_{x} = 256$ fix hai. (a) Sab elements cover karne ke liye kitne work-groups launch karoge? (b) Total kitne work-items run karte hain? (c) Unme se kitne useless kaam karte hain?

Recall Solution

(a) Round up karo: $⌈ 5000/256 ⌉ = ⌈ 19.53 \dots ⌉ = 20$ groups. (b) Total items $= 20 \times 256 = 5120$ . (c) Wasted items $= 5120 - 5000 = 120$ . Inka global_id $5000$ se $5119$ tak hai; kernel guard if (i < 5000) inhe immediately exit karwa deta hai.

Waste kyun accept karte hain: hardware fixed-width batches mein execute karta hai (AMD wavefront = 64, aur 256 = char wavefronts). Aap partial group launch nahi kar sakte, isliye aakhri group hamesha fully populated hota hai chahe woh data se bahar jaaye. Yeh overhang tidy, uniform batches ki keemat hai.

Exercise 3.2 (L3)

Ek AMD GPU ka wavefront 64 threads ka hai; ek NVIDIA GPU ka warp 32 threads ka hai. Aap portability ke liye local_size = 100 choose karte ho. Explain karo, har vendor ke liye, yeh poor choice kyun hai aur ek better number do.

Recall Solution

Wavefront/warp woh sabse chhota indivisible batch hai jo hardware physically issue karta hai. Agar local_size uska whole-number multiple nahi hai, to aakhri batch partly empty run karta hai — woh lanes off hote hain lekin phir bhi ek slot occupy karte hain.

AMD (64) par: $100 = 64 + 36$ . Doosra wavefront sirf 36/64 lanes active ke saath run karta hai → $28/64 \approx 44%$ us wavefront ka waste.
NVIDIA (32) par: $100 = 3 \times 32 + 4 = 96 + 4$ . Chautha warp 4/32 lanes run karta hai → $28/32 \approx 88%$ us warp ka waste.

Better: dono ka multiple choose karo, yani $lcm (32, 64) = 64$ ka multiple. $128$ (ya 256) choose karne se AMD par full wavefronts milte hain ( $128 = 2 \times 64$ ) aur NVIDIA par full warps ( $128 = 4 \times 32$ ) — dono vendors par zero idle lanes. Issi liye portable code 64 se bade powers of two prefer karta hai.

Level 4 — Synthesis

Exercise 4.1 (L4)

Aapke paas ek CUDA saxpy kernel hai (compute karta hai $y [i] = a \cdot x [i] + y [i]$ ). Ise single source se AMD aur NVIDIA dono par run karne ke liye port karo. Concrete API substitutions aur compilation targets list karo, aur batao ki kernel body byte-for-byte identical kyun rehti hai.

Recall Solution

HIP use karo. Kernel body unchanged rahti hai kyunki HIP same built-ins define karta hai:

__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), almost 1:1 rename:

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

Compilation targets (dono HIP-Clang se, NVCC se nahi):

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

Body identical kyun hai: blockIdx, blockDim, threadIdx sirf names hain jo har backend apne intrinsics se map karta hai (AMD par workgroup/workitem IDs, NVIDIA par native CUDA registers). Arithmetic — index math aur a*x+y — plain C++ hai jo kisi bhi machine ke liye compile hoti hai. Sirf vendor-specific plumbing (allocation, launch) ko rename karna pada.

Exercise 4.2 (L4)

Same vector-add task lo aur count karo kitne explicit setup calls OpenCL aapko force karta hai jo CUDA/HIP hide karte hain. Phir trade-off ek sentence mein explain karo.

Recall Solution

OpenCL poora pipeline expose karta hai. Explicit stages (parent ke example se):

clGetPlatformIDs — ek platform choose karo.
clGetDeviceIDs — ek device choose karo.
clCreateContext — us device ka container.
clCreateCommandQueue — kaam submit karne ki queue.
clCreateBuffer — device memory.
clCreateProgramWithSource + clBuildProgram — kernel ko runtime par JIT-compile karo.
clCreateKernel — handle lo.
clSetKernelArg — har argument haath se bind karo.
clEnqueueWriteBuffer / clEnqueueNDRangeKernel — transfer aur launch.
clReleaseMemObject / clReleaseKernel / clReleaseCommandQueue — tear down karo.

Yeh ≈10 stages hain jo CUDA cudaMalloc / cudaMemcpy / <<<>>> / cudaFree mein compress karta hai.

Trade-off (ek sentence): OpenCL ki verbosity aapko kisi bhi vendor ke device par runtime portability deti hai boilerplate ki keemat par, jabki CUDA ki brevity NVIDIA lock-in ki keemat par convenience deti hai.

Level 5 — Mastery

Exercise 5.1 (L5)

Design decision. Aap 200,000-line ka CUDA physics codebase maintain karte ho. Management chahti hai ki yeh AMD MI300 GPUs par bhi run kare, sirf Linux par ship hoga, aur full rewrite afford nahi kar sakte. ROCm/HIP ya OpenCL choose karo, aur platform models se teen concrete facts ke saath justify karo. Phir apni choice ka sabse bada technical risk batao.

Recall Solution

ROCm / HIP choose karo. Justification:

Automatic translation: HIP-ify tooling CUDA ko HIP mein ~80–95% automatically convert karta hai, isliye 200k-line base ko mostly ~5–20% edge cases mein hi edits chahiye — rewrite nahi. OpenCL har ek kernel aur sara host boilerplate scratch se likhne ki demand karta.
Syntax parity: HIP CUDA built-ins (blockIdx, threadIdx) aur runtime API (hipMalloc ↔ cudaMalloc) mirror karta hai, isliye developer ka knowledge directly transfer hota hai. OpenCL ka platform/context/queue model ek alag mental model hai.
Target match: ROCm Linux-first hai aur MI300 ek first-class ROCm target hai — exactly wahi deployment constraint jo diya gaya hai. Aur HIP NVPTX ke through NVIDIA par bhi compile hoti hai, isliye existing NVIDIA fleet bani rehti hai.

Sabse bada technical risk: non-translatable slice. ~5–20% jo HIP-ify auto-convert nahi kar sakta woh usually performance-critical, hand-tuned code hota hai — warp-level intrinsics, PTX inline assembly, wavefront-size assumptions (NVIDIA warp 32 vs AMD wavefront 64). Koi bhi kernel jo shuffle ya reduction ke liye "32" hard-code karta hai woh 64-wide wavefront par silently wrong results produce karega jab tak manually audit nahi hota.

Exercise 5.2 (L5)

Numeric capstone. Aapke MI300 kernel mein har work-item ko matmul ke liye local memory mein $16 \times 16$ floats ka tile rakhna hai. LDS per compute unit 64 KB hai. (a) Ek work-group ko kitne bytes LDS chahiye agar uske paas 256 items hain jahan har item apna tile rakhta hai? (b) Kya yeh fit hota hai? (c) Agar nahi, to 64-wide wavefront ke multiple mein round down karte hue sabse bada local size kya hoga jo fit ho?

Recall Solution

Ek 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 sirf $64 KB = 65, 536$ B hai, isliye $256 KB > 64 KB$ → fit nahi hota (budget ka 4 guna chahiye). (c) Max items $= ⌊ 65, 536/1024 ⌋ = 64$ items. Yeh already exactly $1 \times 64$ hai, ek poora wavefront, isliye koi rounding nahi chahiye → local_size = 64.

Yeh real ceiling kyun hai: occupancy scarcest on-chip resource se bounded hoti hai. Yahan LDS, thread count nahi, wall hai — ek khubsoorat power-of-two 256 group impossible hai isliye nahi ki scheduler ko pasand nahi, balki isliye ki 64 KB scratchpad ki physics forbid karti hai.

Recall Quick self-check clozes

HIP HIP-Clang se compile hoti hai (NVCC se nahi), NVIDIA ke liye NVPTX emit karta hai aur AMD ke liye AMDGPU ISA. AMD wavefront 64 threads ka hota hai; NVIDIA warp 32 threads ka hota hai. Global index formula hai global_id = group_id × local_size + local_id. Woh formula undo kaise hota hai? ::: Integer division: group_id = floor(global_id / local_size), local_id = global_id mod local_size. Group count upar kyun round karte hain? ::: Ek partial work-group launch nahi ho sakta, isliye overhang cover karo aur if (i < n) se bounds-check karo.