ROCm - OpenCL alternatives — Hardware notes

Overview

While CUDA dominates GPU computing, ROCm (Radeon Open Compute) and OpenCL (Open Computing Language) provide vendor-agnostic alternatives for heterogeneous parallel computing. Understanding these platforms is critical because vendor lock-in stifles innovation and portability enables broader hardware access.

Core Concepts

From First Principles: How Vendor Abstraction Works

The Hardware Abstraction Stack

WHY do we need layers? Each GPU vendor has different:

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

WHAT is the solution? A compilation pipeline that translates high-level code → intermediate representation → vendor-specific machine code.

HOW does it work step-by-step?

Step 1: Write portable kernel code

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

Why this step? We use standardized built-ins (get_global_id) instead of hardware-specific registers. The OpenCL runtime will map this to whatever the hardware provides.

Step 2: Runtime compilation

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

Why this step? Just-in-time (JIT) compilation lets the driver optimize for the actual hardware present. If you have an AMD RX 7900, it compiles for RDNA 3. If you have an Intel Arc, it compiles for Xe-HPG.

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)

Why this matters? The abstraction hides that AMD has 64 KB LDS per CU while NVIDIA has 48-128 KB shared memory per SM (48 KB or 96 KB on Fermi–Pascal, up to 128 KB on Ampere). Your code adapts automatically.

Formula Work-Item Execution Model

OpenCL organizes computation into: $\text{NDRange} = \text{Global Work Size} = (G_x, G_y, G_z)$ $\text{Work-Group} = \text{Local Work Size} = (L_x, L_y, L_z)$ $\text{Number of Work-Groups} = \left(\frac{G_x}{L_x}, \frac{G_y}{L_y}, \frac{G_z}{L_z}\right)$

Derivation from scratch:

WHY groups? Hardware executes threads in batches (NVIDIA: 32-thread warps, AMD: 64-thread wavefronts). We need logical grouping that maps to physical execution units.
Start with total work: Say we want to process 1,048,576 elements. $G_x = 1{,}048{,}576, \quad G_y = G_z = 1$
Choose local size to match hardware (AMD wavefront = 64): $L_x = 64, \quad L_y = L_z = 1$
Compute work-groups: $\text{Num Groups}_x = \frac{1{,}048{,}576}{64} = 16{,}384$
Each work-item gets unique ID:

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: $\text{global\_id} = \text{group\_id} \times \text{local\_size} + \text{local\_id}$

WHY this formula? It's just chunk numbering. If you have 16,384 groups of 64 items each, item 4,160 is in group 65 (since $4160 = 65 \times 64 + 0$ ) at local position 0.

ROCm's HIP: CUDA Compatibility Layer

Derivation: How HIP achieves portability

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: Port to HIP (in modern HIP you use the same blockIdx.x/threadIdx.x built-ins)

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

Why this works? HIP defines the same built-in names (blockIdx, threadIdx, blockDim) and API functions (hipMalloc, hipMemcpy) that map to the appropriate backend:

// 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

WHY two paths? Different instruction sets require different code-generation targets, but the source code stays identical. Note HIP does not use NVCC — it uses HIP-Clang, which emits NVPTX (NVIDIA's virtual ISA) for NVIDIA GPUs and AMDGPU ISA for AMD GPUs.

Worked example Complete HIP vs CUDA vs OpenCL Comparison

Task: Add two vectors of 1 million floats 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);

Why these steps? Manual memory management gives control but requires explicit transfers.

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);

Why nearly identical? HIP mirrors the CUDA runtime API almost 1:1 with an hip prefix.

OpenCL approach (more 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);

Why so verbose? OpenCL exposes the entire execution model explicitly. You see platform → device → context → queue → kernel, the full pipeline CUDA hides.

Worked example Matrix Multiplication Performance Analysis

Problem: Multiply two 1024×1024 matrices 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: order of tens of GFLOPS on an RX 6800 (FP32 peak ≈ 16 TFLOPS) — a tiny fraction of peak.

WHY so slow? Each work-item reads:

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

For 1024×1024 outputs: $1{,}048{,}576 \times 2048 = 2.1$ billion memory transactions. Global memory is the bottleneck.

Optimized with local memory:

__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: typically several-fold to ~10× faster than the naive version in practice.

WHY this step? On-chip local memory (LDS on AMD) has roughly 5–10× higher bandwidth (and far lower latency) than global memory. By loading 16×16 tiles cooperatively:

Each tile is loaded once by the work-group
All 256 work-items reuse that tile
Global memory traffic is reduced by a factor of ~16 (each loaded element is reused TILE_SIZE times)

Where the speedup comes from: $\text{Global traffic reduction} \approx \text{TILE\_SIZE} = 16\times$ Since the naive kernel is memory-bound, cutting global traffic ~16× lifts the memory ceiling proportionally. The measured speedup is smaller than 16× because of barrier synchronization overhead and because the tiled kernel eventually approaches compute or LDS-bandwidth limits — so real gains are usually in the single-digit-to-low-double-digit range.

Common Mistakes & Misconceptions

Common mistake Mistake 2: "HIP just renames CUDA functions"

Why it feels right: The API looks like a simple search-and-replace of cuda → hip.

The steel-man: For simple runtime calls, it is mostly a prefix swap — cudaMalloc becomes hipMalloc, cudaMemcpy becomes hipMemcpy, and kernel built-ins like blockIdx.x are identical. So the "renaming" intuition genuinely captures the common case.

The truth: But there are real semantic and architectural differences:

Warp vs wavefront width: CUDA warps are 32 threads; AMD wavefronts are 64 (or 32 on RDNA in wave32 mode). Code that hard-codes 32 in warp-level primitives (__shfl, ballot masks) can break.
Occupancy tuning: Optimal block sizes and register/LDS budgets differ between architectures, so a CUDA-tuned launch config may be suboptimal on AMD.
Some CUDA features have no HIP equivalent (e.g., certain tensor-core / PTX inline assembly paths), requiring manual rewrites.

The fix: Don't assume 1:1 semantics. After the automatic hipify pass, re-tune launch parameters per architecture and audit any warp-size-dependent code. Read the HIP porting guide for the architectural gotchas.

Common mistake Mistake 3: "Work-group size doesn't matter"

Why it feels right: The runtime will execute all work-items regardless of grouping.

The steel-man: Yes, correctness is unaffected. But performance changes dramatically.

Bad choice (AMD RX 6800):

size_t local_size = 13;  // Prime number

Result: Hardware executes in wavefronts of 64. With groups of 13, each work-group only fills part of a wavefront:

13 active threads, 51 idle lanes in that wavefront
13/64 ≈ 20% lane utilization
Large slowdown from wasted execution slots

Good choice:

size_t local_size = 64;  // Matches wavefront size

Result: full wavefront occupancy, far better throughput.

The rule: Choose local size as a multiple of:

NVIDIA: 32 (warp size)
AMD: 64 (wavefront size; 32 in RDNA wave32 mode)
Intel: 8-32 (subgroup size, varies by GPU)

WHY? Hardware SIMD units execute threads in lockstep. Misaligned groups waste execution slots.

Architecture Comparison

Formula Performance Model Derivation

Goal: Predict whether code will run faster on CUDA vs ROCm.

Step 1: Define theoretical peak performance $\text{TFLOPS} = \text{Cores} \times \text{Clock (GHz)} \times \frac{\text{FLOPs}}{\text{cycle}}$

Example hardware:

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

Step 2: Account for memory bandwidth bottleneck $\text{Compute Intensity} = \frac{\text{FLOPs}}{\text{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: Determine bottleneck $\text{Achieved TFLOPS} = \min\left(\text{Peak Compute},\; \frac{\text{Bandwidth} \times \text{Intensity}}{1000}\right)$

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

RX 7900 XTX: Bandwidth = 960 GB/s $\text{Memory-limited} = \frac{960 \times 0.083}{1000} = 0.080 \text{ TFLOPS}$ Conclusion: For vector addition, both GPUs perform nearly identically because memory bandwidth is the bottleneck, not compute. The CUDA vs ROCm choice doesn't matter here.

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

Now NVIDIA wins ~2.7× due to more compute thro

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, yaha core intuition ye samajhna hai ki CUDA bahut powerful hai lekin wo sirf NVIDIA ke GPUs pe hi chalta hai — matlab agar tum hazaaron lines CUDA code likhte ho, toh tum permanently NVIDIA hardware ke saath "lock" ho jaate ho. Isko vendor lock-in kehte hain. Yahi problem solve karne ke liye aaye ROCm (AMD ka open-source platform) aur OpenCL (ek cross-platform standard). Simple example samajh lo — jaise USB-C cable har device mein chal jaati hai, waise OpenCL har vendor ke GPU (AMD, Intel, NVIDIA), aur CPU-FPGA pe bhi chalta hai, jabki CUDA sirf Lightning cable ki tarah ek hi company tak limited hai.

Ab technical intuition ye hai ki har vendor ka hardware alag hota hai — instruction set alag, memory hierarchy alag, thread organization bhi alag (NVIDIA "warps" use karta hai jismein 32 threads hoti hain, AMD "wavefronts" use karta hai jismein 64 threads hoti hain). Toh portability kaise possible hoti hai? Iska trick hai abstraction layers aur JIT (just-in-time) compilation. Tum ek high-level portable kernel likhte ho jismein hardware-specific cheezein directly use nahi karte — jaise get_global_id() jaise standardized built-ins use karte ho. Phir runtime pe ye code LLVM IR (intermediate representation) mein convert hota hai, aur uske baad jo actual hardware present hai uske lihaz se machine code banta hai. Matlab tumhara same code AMD RX 7900 pe RDNA 3 ke liye compile hoga, aur Intel Arc pe Xe-HPG ke liye — automatic!

Ye cheez kyun matter karti hai? Kyunki competition price kam karti hai, innovation badhti hai, aur tumhe broader hardware access milta hai — tum ek hi company ke mehenge GPUs khareedne ke liye majboor nahi hote. Memory spaces (jaise __global, __local, __private) bhi automatically map ho jaate hain — AMD ka LDS ya NVIDIA ka shared memory, tumhe worry karne ki zaroorat nahi. Toh future ke developer ke liye ye concepts samajhna zaroori hai, kyunki real-world mein cost aur flexibility dono cheezein matter karti hain, aur ek achha engineer wahi hai jo vendor lock-in se aware ho aur portable solutions choose kar sake.

ROCm - OpenCL alternatives

Overview

Core Concepts

From First Principles: How Vendor Abstraction Works

The Hardware Abstraction Stack

ROCm's HIP: CUDA Compatibility Layer

Common Mistakes & Misconceptions

Architecture Comparison

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture