ROCm - OpenCL alternatives — Hinglish notes

6.2.15 · Hardware › GPU Architecture

Overview

Jab ki CUDA GPU computing mein dominant hai, ROCm (Radeon Open Compute) aur OpenCL (Open Computing Language) heterogeneous parallel computing ke liye vendor-agnostic alternatives provide karte hain. In platforms ko samajhna critical hai kyunki vendor lock-in innovation ko rok deta hai aur portability broader hardware access enable karti hai.

Core Concepts

First Principles Se: Vendor Abstraction Kaise Kaam Karta Hai

The Hardware Abstraction Stack

Kyun chahiye layers? Har GPU vendor ke paas alag hote hain:

Instruction sets (NVIDIA: PTX, AMD: GCN/RDNA ISA)
Memory hierarchies (L1/L2 cache sizes, shared memory)
Thread organization (NVIDIA warps vs AMD wavefronts)

Kya hai solution? Ek compilation pipeline jo high-level code → intermediate representation → vendor-specific machine code mein translate karta hai.

Kaise kaam karta hai step-by-step?

Step 1: Portable kernel code likho

// OpenCL kernel (vendor-agnostic)
__kernel void vector_add(__global float* A, 
                         __global float* B,
                         __global float* C) {
    int i = get_global_id(0);  // Abstract work-item ID
    C[i] = A[i] + B[i];
}

Yeh step kyun? Hum standardized built-ins (get_global_id) use karte hain hardware-specific registers ki jagah. OpenCL runtime is cheez ko jo bhi hardware provide karta hai usse map kar dega.

Step 2: Runtime compilation

Source code → LLVM IR → Device-specific ISA → Execution

Yeh step kyun? Just-in-time (JIT) compilation driver ko allow karta hai ki woh actual hardware present ke liye optimize kare. Agar tumhare paas AMD RX 7900 hai, toh woh RDNA 3 ke liye compile karta hai. Agar Intel Arc hai, toh Xe-HPG ke liye compile karta hai.

Step 3: Memory space mapping

OpenCL Memory	NVIDIA Mapping	AMD Mapping
`__global`	Global memory	VRAM
`__local`	Shared memory	LDS (Local Data Share)
`__private`	Registers	VGPRs (Vector GPUs)

Yeh kyun matter karta hai? Abstraction chhupa deta hai ki AMD ke paas 64 KB LDS per CU hai jabki NVIDIA ke paas 48-128 KB shared memory per SM hai (48 KB ya 96 KB on Fermi–Pascal, up to 128 KB on Ampere). Tumhara code automatically adapt kar leta hai.

Formula Work-Item Execution Model

OpenCL computation ko organize karta hai: $NDRange = Global Work Size = (G_{x}, G_{y}, G_{z})$ $Work-Group = Local Work Size = (L_{x}, L_{y}, L_{z})$ $Number of Work-Groups = (\frac{G _{x}}{L _{x}}, \frac{G _{y}}{L _{y}}, \frac{G _{z}}{L _{z}})$

Scratch se derivation:

Groups kyun? Hardware threads ko batches mein execute karta hai (NVIDIA: 32-thread warps, AMD: 64-thread wavefronts). Hume logical grouping chahiye jo physical execution units se map ho sake.
Total work se start karo: Maano hum 1,048,576 elements process karna chahte hain. $G_{x} = 1, 048, 576, G_{y} = G_{z} = 1$
Local size choose karo hardware se match karne ke liye (AMD wavefront = 64): $L_{x} = 64, L_{y} = L_{z} = 1$
Work-groups compute karo: $Num Groups_{x} = \frac{1 , 048 , 576}{64} = 16, 384$
Har work-item ko unique ID milta hai:

Global ID: get_global_id(0) returns 0 to 1,048,575
Local ID: get_local_id(0) returns 0 to 63 (within group)
Group ID: get_group_id(0) returns 0 to 16,383

Relationship: $global_id = group_id \times local_size + local_id$

Yeh formula kyun? Yeh bas chunk numbering hai. Agar tumhare paas 16,384 groups hain of 64 items each, toh item 4,160, group 65 mein hai (kyunki $4160 = 65 \times 64 + 0$ ) local position 0 par.

ROCm ka HIP: CUDA Compatibility Layer

Derivation: HIP portability kaise achieve karta hai

Start: CUDA kernel

__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];
}

Step 1: HIP par port karo (modern HIP mein same blockIdx.x/threadIdx.x built-ins use hote hain)

// HIP kernel (nearly identical to CUDA)
__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];
}

Yeh kyun kaam karta hai? HIP same built-in names define karta hai (blockIdx, threadIdx, blockDim) aur API functions (hipMalloc, hipMemcpy) jo appropriate backend se map hote hain:

// When compiling for NVIDIA (HIP-Clang targeting NVPTX):
//   hipMalloc  -> cudaMalloc
//   blockIdx.x -> CUDA blockIdx.x
 
// When compiling for AMD (HIP-Clang targeting AMDGPU):
//   hipMalloc  -> native ROCm allocation
//   blockIdx.x -> AMDGPU workgroup/workitem intrinsics

Step 2: Compilation routes

HIP source → [HIP-Clang → NVPTX] → NVIDIA binary
HIP source → [HIP-Clang → AMDGPU] → ROCm binary

Do paths kyun? Alag instruction sets ke liye alag code-generation targets ki zaroorat hoti hai, lekin source code identical rehta hai. Note karo ki HIP NVCC use nahi karta — yeh HIP-Clang use karta hai, jo NVIDIA GPUs ke liye NVPTX (NVIDIA ka virtual ISA) emit karta hai aur AMD GPUs ke liye AMDGPU ISA.

Worked example Complete HIP vs CUDA vs OpenCL Comparison

Task: 1 million floats ke do vectors add karo CUDA approach:

// Allocate
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);
 
// Copy
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
// Launch
vector_add<<<blocks, threads>>>(d_A, d_B, d_C);
// Cleanup
cudaFree(d_A);

Yeh steps kyun? Manual memory management control deta hai lekin explicit transfers ki zaroorat padti hai.

HIP approach (nearly identical syntax):

hipMalloc(&d_A, size);
hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice);
hipLaunchKernelGGL(vector_add, blocks, threads, 0, 0, d_A, d_B, d_C);
hipFree(d_A);

Nearly identical kyun? HIP, CUDA runtime API ko almost 1:1 mirror karta hai ek hip prefix ke saath.

OpenCL approach (zyada verbose):

// 1. Get platform and device
clGetPlatformIDs(1, &platform, NULL);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
 
// 2. Create context and queue
context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
queue = clCreateCommandQueue(context, device, 0, &err);
// 3. Create buffers
d_A = clCreateBuffer(context, CL_MEM_READ_ONLY, size, NULL, &err);
// 4. Compile kernel
program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);
clBuildProgram(program, 1, &device, NULL, NULL, NULL);
kernel = clCreateKernel(program, "vector_add", &err);
 
// 5. Set arguments
clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_A);
// 6. Enqueue operations
clEnqueueWriteBuffer(queue, d_A, CL_TRUE, 0, size, h_A, 0, NULL, NULL);
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, NULL);
// 7. Cleanup
clReleaseMemObject(d_A);
clReleaseKernel(kernel);
clReleaseCommandQueue(queue);

Itna verbose kyun? OpenCL poore execution model ko explicitly expose karta hai. Tum platform → device → context → queue → kernel, poori pipeline dekh sakte ho jo CUDA chhupa deta hai.

Worked example Matrix Multiplication Performance Analysis

Problem: Do 1024×1024 matrices multiply karo Naive OpenCL kernel:

__kernel void matmul_naive(__global float* A,
                          __global float* B,
                          __global float* C,
                          int N) {
   int row = get_global_id(0);
   int col = get_global_id(1);
   
   float sum = 0.0f;
   for (int k = 0; k < N; k++) {
       sum += A[row * N + k] * B[k * N + col];
   }
   C[row * N + col] = sum;
}

Performance: RX 6800 par tens of GFLOPS ke order mein (FP32 peak ≈ 16 TFLOPS) — peak ka ek tiny fraction.

Itna slow kyun? Har work-item read karta hai:

A se 1024 elements (row)
B se 1024 elements (column, non-contiguous!)
Total: 2048 global memory accesses per output

1024×1024 outputs ke liye: $1, 048, 576 \times 2048 = 2.1$ billion memory transactions. Global memory bottleneck hai.

Local memory se optimized:

__kernel void matmul_tiled(__global float* A,
                          __global float* B,
                          __global float* C,
                          int N,
                          __local float* A_tile,
                          __local float* B_tile) {
   int TILE_SIZE = 16;
   int row = get_global_id(0);
   int col = get_global_id(1);
   int local_row = get_local_id(0);
   int local_col = get_local_id(1);
   float sum = 0.0f;
   // Process in tiles
   for (int t = 0; t < N / TILE_SIZE; t++) {
       // Cooperatively load tile into local memory
       A_tile[local_row * TILE_SIZE + local_col] = 
           A[row * N + (t * TILE_SIZE + local_col)];
       B_tile[local_row * TILE_SIZE + local_col] = 
           B[(t * TILE_SIZE + local_row) * N + col];
       
       barrier(CLK_LOCAL_MEM_FENCE);  // Synchronize work-group
       
       // Compute using local memory
       for (int k = 0; k < TILE_SIZE; k++) {
           sum += A_tile[local_row * TILE_SIZE + k] * 
                  B_tile[k * TILE_SIZE + local_col];
       }
       
       barrier(CLK_LOCAL_MEM_FENCE);
   }
   
   C[row * N + col] = sum;
}

Performance: practice mein naive version se typically kaafi-fold se ~10× faster.

Yeh step kyun? On-chip local memory (AMD par LDS) ki roughly 5–10× zyada bandwidth hoti hai (aur global memory se far lower latency). 16×16 tiles cooperatively load karke:

Har tile work-group dwara ek baar load hoti hai
Saare 256 work-items us tile ko reuse karte hain
Global memory traffic roughly ~16 factor se reduce hoti hai (har loaded element TILE_SIZE times reuse hota hai)

Speedup kahaan se aata hai: $Global traffic reduction \approx TILE_SIZE = 16 \times$ Kyunki naive kernel memory-bound hai, global traffic ~16× cut karna memory ceiling ko proportionally lift karta hai. Measured speedup 16× se kam hota hai barrier synchronization overhead ki wajah se aur isliye bhi ki tiled kernel eventually compute ya LDS-bandwidth limits approach karta hai — toh real gains usually single-digit-to-low-double-digit range mein hote hain.

Common Mistakes & Misconceptions

Common mistake Mistake 2: "HIP sirf CUDA functions rename karta hai"

Kyun sahi lagta hai: API bilkul ek simple search-and-replace jaisi lagti hai cuda → hip ki.

Steel-man: Simple runtime calls ke liye, yeh mostly prefix swap hi hai — cudaMalloc ban jaata hai hipMalloc, cudaMemcpy ban jaata hai hipMemcpy, aur kernel built-ins jaise blockIdx.x identical hain. Toh "renaming" wali intuition genuinely common case capture karti hai.

Sach: Lekin real semantic aur architectural differences hain:

Warp vs wavefront width: CUDA warps 32 threads ke hain; AMD wavefronts 64 ke hain (ya RDNA wave32 mode mein 32). Code jo 32 hard-code karta hai warp-level primitives mein (__shfl, ballot masks) break ho sakta hai.
Occupancy tuning: Optimal block sizes aur register/LDS budgets architectures ke beech differ karte hain, toh ek CUDA-tuned launch config AMD par suboptimal ho sakti hai.
Kuch CUDA features ka HIP equivalent nahi hota (jaise kuch tensor-core / PTX inline assembly paths), jinhein manual rewrites chahiye.

Fix: 1:1 semantics assume mat karo. Automatic hipify pass ke baad, launch parameters ko per architecture re-tune karo aur koi bhi warp-size-dependent code audit karo. Architectural gotchas ke liye HIP porting guide padho.

Common mistake Mistake 3: "Work-group size matter nahi karta"

Kyun sahi lagta hai: Runtime saare work-items execute karega chahे grouping kuch bhi ho.

Steel-man: Haan, correctness unaffected hai. Lekin performance dramatically change hoti hai.

Bad choice (AMD RX 6800):

size_t local_size = 13;  // Prime number

Result: Hardware 64 ke wavefronts mein execute karta hai. 13 ke groups ke saath, har work-group sirf ek wavefront ka part fill karta hai:

13 active threads, us wavefront mein 51 idle lanes
13/64 ≈ 20% lane utilization
Wasted execution slots se bada slowdown

Good choice:

size_t local_size = 64;  // Matches wavefront size

Result: full wavefront occupancy, bahut better throughput.

Rule: Local size choose karo multiple of:

NVIDIA: 32 (warp size)
AMD: 64 (wavefront size; RDNA wave32 mode mein 32)
Intel: 8-32 (subgroup size, GPU ke hisaab se vary karta hai)

Kyun? Hardware SIMD units threads ko lockstep mein execute karte hain. Misaligned groups execution slots waste karte hain.

Architecture Comparison

Formula Performance Model Derivation

Goal: Predict karo ki code CUDA par faster chalega ya ROCm par.

Step 1: Theoretical peak performance define karo $TFLOPS = Cores \times Clock (GHz) \times \frac{FLOPs}{cycle}$

Example hardware:

NVIDIA RTX 4090: 16,384 CUDA cores @ 2.52 GHz, 2 FP32/cycle $Peak = 16, 384 \times 2.52 \times 2 = 82.6 TFLOPS$
AMD RX 7900 XTX: 6,144 stream processors @ 2.5 GHz, 2 FP32/cycle $Peak = 6, 144 \times 2.5 \times 2 = 30.7 TFLOPS$

Step 2: Memory bandwidth bottleneck account karo $Compute Intensity = \frac{FLOPs}{Bytes accessed}$

Example: Vector addition ( $C [i] = A [i] + B [i]$ )

Ops: 1 addition = 1 FLOP
Memory: Read A (4 bytes) + Read B (4 bytes) + Write C (4 bytes) = 12 bytes
Intensity: $\frac{1}{12} = 0.083$ FLOP/byte

Step 3: Bottleneck determine karo $Achieved TFLOPS = min (Peak Compute, \frac{Bandwidth \times Intensity}{1000})$

RTX 4090: Bandwidth = 1,008 GB/s $Memory-limited = \frac{1008 \times 0.083}{1000} = 0.084 TFLOPS$ $min (82.6, 0.084) = 0.084 TFLOPS (memory-bound!)$

RX 7900 XTX: Bandwidth = 960 GB/s $Memory-limited = \frac{960 \times 0.083}{1000} = 0.080 TFLOPS$ Conclusion: Vector addition ke liye, dono GPUs nearly identically perform karte hain kyunki memory bandwidth bottleneck hai, compute nahi. Yahaan CUDA vs ROCm ka choice matter nahi karta.

Counterexample: Matrix multiply (tiling ke saath, intensity ~100 FLOP/byte) $RTX 4090 : min (82.6, 100.8) = 82.6 TFLOPS (compute-bound)$ $RX 7900 XTX : min (30.7, 96.0) = 30.7 TFLOPS (compute-bound)$

Ab NVIDIA ~2.7× jeet jaata hai zyada compute throughput ki wajah se.