GPU Programming Without CUDA: OpenCL, ROCm, Vulkan, Metal, and WebGPU

April 12, 2026 21 min read The Physics of Computing

gpuopenclrocmvulkanmetalwebgpucomputeperformanceamd

Contents

GPU Programming Without CUDA: OpenCL, ROCm, Vulkan, Metal, and WebGPU

CUDA is the dominant GPU compute platform, and for good reason — NVIDIA’s toolchain, libraries (cuBLAS, cuDNN, NCCL), and ecosystem are genuinely excellent. But CUDA only runs on NVIDIA hardware. If you’re targeting AMD GPUs, Apple Silicon, Intel Arc, mobile devices, or the browser, you need something else.

This post covers the full non-CUDA landscape: what each platform is, what hardware it targets, where it excels, and enough working code to understand the programming model. We’ll also cover abstraction layers that let you write once and run on multiple backends.

The GPU Compute Landscape

A quick orientation before diving in:

Platform	Vendor	Hardware	Primary Use
CUDA	NVIDIA	NVIDIA GPUs only	ML training, HPC, scientific compute
ROCm / HIP	AMD	AMD GPUs (+ NVIDIA via translation)	ML training, HPC
OpenCL	Khronos	NVIDIA, AMD, Intel, CPUs, FPGAs	Portable compute, embedded
Metal Compute	Apple	Apple Silicon, AMD in older Macs	iOS/macOS apps, ML on Apple
Vulkan Compute	Khronos	Any Vulkan-capable GPU	Games, graphics pipelines + compute
WebGPU	W3C/browsers	Any GPU via browser	Browser apps, portable compute
SYCL / oneAPI	Intel/Khronos	Intel GPUs, AMD, NVIDIA, CPUs	HPC portability

OpenCL: The Portable Option

OpenCL (Open Computing Language) is the ISO-standardized, cross-vendor GPU compute API. It runs on NVIDIA, AMD, Intel GPUs, Apple Silicon, CPUs, FPGAs, and DSPs. If you need code that runs on hardware you don’t control, OpenCL is the most universal option.

The tradeoff: OpenCL is more verbose than CUDA, vendor implementations vary in quality, and it doesn’t have an ecosystem of pre-built libraries comparable to cuDNN. It’s also perpetually one major version behind what any single vendor’s hardware could support.

The OpenCL Programming Model

OpenCL has a host-device model similar to CUDA:

Host: your CPU program, manages memory and launches kernels
Device: the GPU (or other accelerator)
Kernel: a function that runs in parallel on the device, written in OpenCL C (a C99 subset with extensions)

The key objects:

cl_platform_id — a vendor’s OpenCL implementation (e.g., “AMD”, “Intel”)
cl_device_id — a specific compute device
cl_context — groups devices and manages memory
cl_command_queue — serializes operations on a device
cl_program — compiled kernels
cl_kernel — a single kernel entry point
cl_mem — a buffer on the device

A Complete OpenCL Example: Vector Addition

Kernel (OpenCL C):

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12


// vadd.cl
__kernel void vector_add(
    __global const float* a,
    __global const float* b,
    __global float* c,
    const int n)
{
    int gid = get_global_id(0);
    if (gid < n) {
        c[gid] = a[gid] + b[gid];
    }
}

Host program (C):

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86


#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

#define N 1024 * 1024

const char *kernel_source =
"__kernel void vector_add("
"    __global const float* a,"
"    __global const float* b,"
"    __global float* c,"
"    const int n)"
"{"
"    int gid = get_global_id(0);"
"    if (gid < n) c[gid] = a[gid] + b[gid];"
"}";

int main() {
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;
    cl_mem buf_a, buf_b, buf_c;
    cl_int err;

    // 1. Get platform and device
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);

    // 2. Create context and command queue
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    queue = clCreateCommandQueueWithProperties(context, device, 0, &err);

    // 3. Compile kernel
    program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);
    clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    kernel = clCreateKernel(program, "vector_add", &err);

    // 4. Allocate and fill host data
    float *h_a = malloc(N * sizeof(float));
    float *h_b = malloc(N * sizeof(float));
    float *h_c = malloc(N * sizeof(float));
    for (int i = 0; i < N; i++) { h_a[i] = i; h_b[i] = i * 2; }

    // 5. Create device buffers and copy data
    buf_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                           N * sizeof(float), h_a, &err);
    buf_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                           N * sizeof(float), h_b, &err);
    buf_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                           N * sizeof(float), NULL, &err);

    // 6. Set kernel arguments and launch
    int n = N;
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_a);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_b);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_c);
    clSetKernelArg(kernel, 3, sizeof(int), &n);

    size_t global_size = N;
    size_t local_size = 256;
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                           &global_size, &local_size, 0, NULL, NULL);

    // 7. Read result back
    clEnqueueReadBuffer(queue, buf_c, CL_TRUE, 0,
                        N * sizeof(float), h_c, 0, NULL, NULL);

    printf("c[0] = %.1f (expected %.1f)\n", h_c[0], h_a[0] + h_b[0]);
    printf("c[1023] = %.1f (expected %.1f)\n", h_c[1023], h_a[1023] + h_b[1023]);

    // Cleanup
    clReleaseMemObject(buf_a); clReleaseMemObject(buf_b); clReleaseMemObject(buf_c);
    clReleaseKernel(kernel); clReleaseProgram(program);
    clReleaseCommandQueue(queue); clReleaseContext(context);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

1
2
3
4


# Compile
gcc -o vadd vadd.c -lOpenCL
# macOS
gcc -o vadd vadd.c -framework OpenCL

OpenCL Memory Model

OpenCL defines a hierarchy of memory spaces (similar to CUDA):

Memory Type	OpenCL Qualifier	Scope	Speed
Global	`__global`	All work-items	Slowest
Local	`__local`	Work-group (like CUDA shared memory)	Fast
Private	(default)	Single work-item (like CUDA registers)	Fastest
Constant	`__constant`	Read-only global, cached	Fast reads

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27


__kernel void matrix_multiply(
    __global const float* A,
    __global const float* B,
    __global float* C,
    __local float* tile_A,   // work-group shared memory
    __local float* tile_B,
    const int N)
{
    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);
    int tile_size = get_local_size(0);

    float sum = 0.0f;
    for (int t = 0; t < N / tile_size; t++) {
        // Load tile into local memory (fast, shared within work-group)
        tile_A[local_row * tile_size + local_col] = A[row * N + t * tile_size + local_col];
        tile_B[local_row * tile_size + local_col] = B[(t * tile_size + local_row) * N + col];
        barrier(CLK_LOCAL_MEM_FENCE);  // sync before use

        for (int k = 0; k < tile_size; k++)
            sum += tile_A[local_row * tile_size + k] * tile_B[k * tile_size + local_col];
        barrier(CLK_LOCAL_MEM_FENCE);  // sync before next tile
    }
    C[row * N + col] = sum;
}

OpenCL 3.0 and the Current State

OpenCL 3.0 (released 2020) took a pragmatic turn: it made most features optional and allowed vendors to implement only what their hardware supports. This solved the compliance problem but created a new one — you need to query at runtime which features are available.

The practical reality in 2026: OpenCL is well-supported on AMD (via ROCm), Intel (Arc GPUs and integrated graphics), and Apple Silicon (Metal-backed OpenCL). NVIDIA’s OpenCL support is frozen at OpenCL 3.0 with only mandatory features — they have no incentive to improve the competition.

ROCm and HIP: AMD’s CUDA Alternative

ROCm (Radeon Open Compute) is AMD’s open-source compute platform. HIP (Heterogeneous-compute Interface for Portability) is the CUDA-like programming API within ROCm.

HIP’s killer feature: it’s syntactically very similar to CUDA, and AMD provides a translation tool (hipify) that converts CUDA code to HIP automatically. The same HIP code compiles for AMD GPUs (via ROCm/HCC) and NVIDIA GPUs (via CUDA as a backend).

HIP Programming Model

HIP mirrors CUDA almost exactly:

CUDA	HIP
`cudaMalloc`	`hipMalloc`
`cudaMemcpy`	`hipMemcpy`
`__global__`	`__global__`
`threadIdx.x`	`hipThreadIdx_x`
`blockIdx.x`	`hipBlockIdx_x`
`__syncthreads()`	`__syncthreads()`
`cudaStream_t`	`hipStream_t`

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47


// vector_add.hip.cpp
#include <hip/hip_runtime.h>
#include <stdio.h>

#define N (1024 * 1024)

__global__ void vector_add(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;
    size_t size = N * sizeof(float);

    // Allocate host memory
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);
    for (int i = 0; i < N; i++) { h_a[i] = i; h_b[i] = i * 2.0f; }

    // Allocate device memory
    hipMalloc(&d_a, size);
    hipMalloc(&d_b, size);
    hipMalloc(&d_c, size);

    // Copy to device
    hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice);

    // Launch kernel
    int threads = 256;
    int blocks = (N + threads - 1) / threads;
    hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threads), 0, 0,
                       d_a, d_b, d_c, N);

    // Copy result back
    hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost);
    printf("c[0] = %.1f\n", h_c[0]);

    hipFree(d_a); hipFree(d_b); hipFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

1
2
3
4
5


# Compile for AMD GPU
hipcc -o vector_add vector_add.hip.cpp

# Compile for NVIDIA GPU (using CUDA backend)
hipcc --platform nvidia -o vector_add vector_add.hip.cpp

Converting CUDA to HIP

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11


# Install hipify tools
sudo apt install hipify-clang  # from ROCm repo

# Convert a single file
hipify-clang cuda_code.cu -o hip_code.cpp

# Convert an entire project
hipify-perl --inplace project/src/*.cu

# Most CUDA patterns translate 1:1
# Exceptions: PTX inline assembly, some cuBLAS internals, NVLink

ROCm ML Libraries

ROCm ships equivalents of NVIDIA’s ML libraries:

NVIDIA	AMD ROCm
cuBLAS	rocBLAS
cuDNN	MIOpen
cuFFT	rocFFT
cuSPARSE	rocSPARSE
NCCL	RCCL
Thrust	rocThrust

PyTorch and TensorFlow both have ROCm backends:

1
2
3
4
5
6
7


# Install PyTorch for ROCm
pip install torch torchvision torchaudio \
  --index-url https://download.pytorch.org/whl/rocm6.1

# Verify
python -c "import torch; print(torch.cuda.is_available())"  # Returns True on ROCm
python -c "import torch; print(torch.version.hip)"

PyTorch on ROCm uses the same CUDA API from Python (torch.cuda.*) — the ROCm backend is transparent to Python code.

Supported Hardware

ROCm supports AMD’s RDNA 2 (RX 6000 series), RDNA 3 (RX 7000 series), RDNA 4 (RX 9000 series), and CDNA (Instinct MI series) GPUs. The Instinct MI300X is AMD’s flagship for ML training — it’s competitive with the H100 for many workloads and has 192 GB of HBM3 on the MI300X.

1
2
3
4
5
6


# Check ROCm device info
rocm-smi
rocminfo | grep -A5 "Agent [0-9]"

# Check GPU utilization
watch -n 1 rocm-smi

Vulkan Compute: Graphics Pipeline + Compute

Vulkan is a low-level graphics API from Khronos that also exposes GPU compute through compute shaders. It runs on any GPU with a Vulkan driver — NVIDIA, AMD, Intel, ARM Mali, Qualcomm Adreno, and more.

Vulkan Compute makes the most sense when:

You’re already using Vulkan for graphics and want to integrate compute
You need the exact same code path for iOS/Android, Windows, Linux, and macOS (via MoltenVK)
You want to avoid the overhead of CUDA/OpenCL context management in a render loop
You need tight interop between graphics and compute (no PCIe copy needed)

The tradeoff: Vulkan has massive boilerplate. A “hello world” compute shader in Vulkan is ~300 lines of C. Most production use wraps it in a higher-level framework.

Compute Shaders in GLSL

Vulkan shaders are written in GLSL (or HLSL) and compiled to SPIR-V bytecode:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25


// vadd.comp — GLSL compute shader
#version 450

layout(local_size_x = 256) in;

layout(set = 0, binding = 0) readonly buffer InputA {
    float a[];
};
layout(set = 0, binding = 1) readonly buffer InputB {
    float b[];
};
layout(set = 0, binding = 2) writeonly buffer Output {
    float c[];
};

layout(push_constant) uniform PushConstants {
    uint n;
} pc;

void main() {
    uint idx = gl_GlobalInvocationID.x;
    if (idx < pc.n) {
        c[idx] = a[idx] + b[idx];
    }
}

1
2
3
4
5


# Compile to SPIR-V
glslc -fshader-stage=compute vadd.comp -o vadd.spv

# Or with glslangValidator
glslangValidator -V vadd.comp -o vadd.spv

Vulkan Compute Dispatch (Simplified)

The actual Vulkan boilerplate to dispatch this shader is substantial, but the high-level steps are:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28


// Pseudocode — real Vulkan requires ~300 lines
VkInstance instance = createInstance();
VkPhysicalDevice physDevice = pickPhysicalDevice(instance);
VkDevice device = createLogicalDevice(physDevice);

// Create buffers
VkBuffer bufA = createBuffer(device, N * sizeof(float), VK_BUFFER_USAGE_STORAGE_BUFFER_BIT);
VkBuffer bufB = createBuffer(device, N * sizeof(float), VK_BUFFER_USAGE_STORAGE_BUFFER_BIT);
VkBuffer bufC = createBuffer(device, N * sizeof(float), VK_BUFFER_USAGE_STORAGE_BUFFER_BIT);

// Upload data
copyToBuffer(bufA, h_a, N * sizeof(float));
copyToBuffer(bufB, h_b, N * sizeof(float));

// Create compute pipeline
VkShaderModule shader = loadSPIRV(device, "vadd.spv");
VkPipeline pipeline = createComputePipeline(device, shader, layout);

// Record and submit command buffer
VkCommandBuffer cmd = beginCommandBuffer();
vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 0, 1, &descriptorSet, 0, NULL);
vkCmdPushConstants(cmd, layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &n);
vkCmdDispatch(cmd, (N + 255) / 256, 1, 1);
endAndSubmitCommandBuffer(cmd, queue);

// Read result
copyFromBuffer(bufC, h_c, N * sizeof(float));

Higher-Level Vulkan Compute: kompute

kompute is a framework that wraps Vulkan compute to eliminate boilerplate:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26


#include <kompute/Kompute.hpp>

int main() {
    kp::Manager mgr;  // initializes Vulkan automatically

    std::vector<float> h_a(N), h_b(N), h_c(N);
    // ... fill h_a and h_b

    // Create tensors (GPU buffers)
    auto t_a = mgr.tensor(h_a);
    auto t_b = mgr.tensor(h_b);
    auto t_c = mgr.tensor(h_c);

    // Load SPIR-V shader
    auto spirv = readSPIRV("vadd.spv");
    auto algo = mgr.algorithm({t_a, t_b, t_c}, spirv, kp::Workgroup({N/256, 1, 1}));

    // Build and run sequence
    mgr.sequence()
        ->record<kp::OpTensorSyncDevice>({t_a, t_b})
        ->record<kp::OpAlgoDispatch>(algo)
        ->record<kp::OpTensorSyncLocal>({t_c})
        ->eval();

    printf("c[0] = %.1f\n", t_c->data<float>()[0]);
}

Metal Compute: Apple Silicon

Metal is Apple’s GPU API, introduced in 2014. On Apple Silicon (M-series Macs, A-series iPhones/iPads), Metal Compute is the only first-party path to the GPU — there’s no CUDA, no ROCm, no Vulkan (MoltenVK translates Vulkan to Metal, but at overhead cost).

Metal Compute is genuinely excellent. Apple Silicon has a unified memory architecture (GPU and CPU share the same physical RAM), which eliminates the PCIe bottleneck that limits discrete GPU compute. On an M4 Pro with 48 GB unified memory, you can run 30B parameter models that simply don’t fit on a 24 GB VRAM discrete card.

Metal Shading Language

Metal shaders are written in Metal Shading Language (MSL), a C++14 subset:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15


// vadd.metal
#include <metal_stdlib>
using namespace metal;

kernel void vector_add(
    device const float* a [[buffer(0)]],
    device const float* b [[buffer(1)]],
    device float* c       [[buffer(2)]],
    constant uint& n      [[buffer(3)]],
    uint gid              [[thread_position_in_grid]])
{
    if (gid < n) {
        c[gid] = a[gid] + b[gid];
    }
}

Metal Compute in Swift

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46


import Metal
import Foundation

let N = 1024 * 1024

// 1. Get the default GPU
guard let device = MTLCreateSystemDefaultDevice() else { fatalError("No Metal device") }

// 2. Load shader from the default library (compiled at build time)
let library = device.makeDefaultLibrary()!
let function = library.makeFunction(name: "vector_add")!
let pipeline = try! device.makeComputePipelineState(function: function)

// 3. Create command queue
let queue = device.makeCommandQueue()!

// 4. Create buffers (on unified memory - no explicit copy needed!)
let bufA = device.makeBuffer(length: N * MemoryLayout<Float>.size, options: .storageModeShared)!
let bufB = device.makeBuffer(length: N * MemoryLayout<Float>.size, options: .storageModeShared)!
let bufC = device.makeBuffer(length: N * MemoryLayout<Float>.size, options: .storageModeShared)!

// Fill input buffers (direct pointer access - no copy!)
let pA = bufA.contents().bindMemory(to: Float.self, capacity: N)
let pB = bufB.contents().bindMemory(to: Float.self, capacity: N)
for i in 0..<N { pA[i] = Float(i); pB[i] = Float(i) * 2.0 }

// 5. Encode and dispatch
var n = UInt32(N)
let cmdBuf = queue.makeCommandBuffer()!
let encoder = cmdBuf.makeComputeCommandEncoder()!
encoder.setComputePipelineState(pipeline)
encoder.setBuffer(bufA, offset: 0, index: 0)
encoder.setBuffer(bufB, offset: 0, index: 1)
encoder.setBuffer(bufC, offset: 0, index: 2)
encoder.setBytes(&n, length: MemoryLayout<UInt32>.size, index: 3)

let threadsPerGroup = MTLSize(width: pipeline.maxTotalThreadsPerThreadgroup, height: 1, depth: 1)
let numGroups = MTLSize(width: (N + threadsPerGroup.width - 1) / threadsPerGroup.width, height: 1, depth: 1)
encoder.dispatchThreadgroups(numGroups, threadsPerThreadgroup: threadsPerGroup)
encoder.endEncoding()
cmdBuf.commit()
cmdBuf.waitUntilCompleted()

// 6. Read result (no copy needed - unified memory!)
let pC = bufC.contents().bindMemory(to: Float.self, capacity: N)
print("c[0] = \(pC[0])")

MLX: ML on Apple Silicon

Apple’s MLX framework (open-sourced in late 2023) is purpose-built for ML on Apple Silicon. It provides a NumPy-like API with lazy evaluation, unified memory, and Metal acceleration:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22


import mlx.core as mx
import mlx.nn as nn

# Arrays live in unified memory — no explicit GPU copies
a = mx.array([1.0, 2.0, 3.0])
b = mx.array([4.0, 5.0, 6.0])
c = a + b  # lazy — not computed yet
mx.eval(c)  # trigger computation on GPU
print(c)  # array([5, 7, 9], dtype=float32)

# Matrix multiplication (dispatches to Metal)
A = mx.random.normal((4096, 4096))
B = mx.random.normal((4096, 4096))
C = A @ B
mx.eval(C)

# Simple neural network
model = nn.Sequential(
    nn.Linear(784, 256),
    nn.ReLU(),
    nn.Linear(256, 10),
)

llama.cpp, Ollama, and LM Studio all use Metal for inference on Apple Silicon. The M4 Max achieves competitive inference speeds with discrete GPUs for 7B–30B models.

WebGPU: GPU Compute in the Browser (and Beyond)

WebGPU is the W3C standard for GPU access from JavaScript (and now Rust, Python, and C++ via native ports). It launched in Chrome in 2023, Firefox in 2024, and Safari/WebKit shortly after.

WebGPU’s compute model is based on Vulkan/Metal/D3D12 — it’s intentionally modern and explicitly excludes the legacy baggage of WebGL (which was based on OpenGL ES 2.0).

WebGPU Compute in JavaScript

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74


// WebGPU vector addition
async function vectorAdd() {
  const N = 1024 * 1024;

  // 1. Get GPU adapter and device
  const adapter = await navigator.gpu.requestAdapter();
  const device = await adapter.requestDevice();

  // 2. Define the compute shader (WGSL)
  const shaderModule = device.createShaderModule({
    code: `
      @group(0) @binding(0) var<storage, read> a: array<f32>;
      @group(0) @binding(1) var<storage, read> b: array<f32>;
      @group(0) @binding(2) var<storage, read_write> c: array<f32>;

      @compute @workgroup_size(256)
      fn main(@builtin(global_invocation_id) gid: vec3u) {
        let idx = gid.x;
        if (idx < arrayLength(&a)) {
          c[idx] = a[idx] + b[idx];
        }
      }
    `
  });

  // 3. Create buffers
  const bufSize = N * Float32Array.BYTES_PER_ELEMENT;

  const gpuBufA = device.createBuffer({ size: bufSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
  const gpuBufB = device.createBuffer({ size: bufSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
  const gpuBufC = device.createBuffer({ size: bufSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
  const stagingBuf = device.createBuffer({ size: bufSize, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST });

  // 4. Upload data
  const h_a = new Float32Array(N).map((_, i) => i);
  const h_b = new Float32Array(N).map((_, i) => i * 2);
  device.queue.writeBuffer(gpuBufA, 0, h_a);
  device.queue.writeBuffer(gpuBufB, 0, h_b);

  // 5. Create pipeline and bind group
  const pipeline = device.createComputePipeline({
    layout: 'auto',
    compute: { module: shaderModule, entryPoint: 'main' }
  });

  const bindGroup = device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: gpuBufA } },
      { binding: 1, resource: { buffer: gpuBufB } },
      { binding: 2, resource: { buffer: gpuBufC } },
    ]
  });

  // 6. Encode and submit
  const encoder = device.createCommandEncoder();
  const pass = encoder.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(Math.ceil(N / 256));
  pass.end();

  encoder.copyBufferToBuffer(gpuBufC, 0, stagingBuf, 0, bufSize);
  device.queue.submit([encoder.finish()]);

  // 7. Read result
  await stagingBuf.mapAsync(GPUMapMode.READ);
  const result = new Float32Array(stagingBuf.getMappedRange());
  console.log('c[0]:', result[0]);  // 0
  console.log('c[1]:', result[1]);  // 3
  stagingBuf.unmap();
}

vectorAdd();

WGSL: WebGPU Shading Language

WGSL (WebGPU Shading Language) was designed from scratch for WebGPU. It’s memory-safe, strongly typed, and doesn’t have pointers — by design, to make security analysis tractable in a browser environment.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41


// matrix_multiply.wgsl
struct Matrices {
  data: array<f32>,
}

@group(0) @binding(0) var<storage, read> A: Matrices;
@group(0) @binding(1) var<storage, read> B: Matrices;
@group(0) @binding(2) var<storage, read_write> C: Matrices;

struct Params {
  N: u32,
}
@group(0) @binding(3) var<uniform> params: Params;

var<workgroup> tileA: array<f32, 256>;  // 16x16 tile
var<workgroup> tileB: array<f32, 256>;

@compute @workgroup_size(16, 16)
fn main(
  @builtin(global_invocation_id) gid: vec3u,
  @builtin(local_invocation_id) lid: vec3u,
  @builtin(workgroup_id) wid: vec3u,
) {
  let row = gid.y;
  let col = gid.x;
  let N = params.N;
  let tile = 16u;

  var sum = 0.0f;
  for (var t = 0u; t < N / tile; t++) {
    tileA[lid.y * tile + lid.x] = A.data[row * N + t * tile + lid.x];
    tileB[lid.y * tile + lid.x] = B.data[(t * tile + lid.y) * N + col];
    workgroupBarrier();

    for (var k = 0u; k < tile; k++) {
      sum += tileA[lid.y * tile + k] * tileB[k * tile + lid.x];
    }
    workgroupBarrier();
  }
  C.data[row * N + col] = sum;
}

wgpu: Native WebGPU

wgpu is a Rust implementation of the WebGPU API that runs natively (not just in browsers). It translates to Vulkan, Metal, D3D12, and OpenGL ES depending on the platform — the same WGSL shader runs everywhere.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29


// Cargo.toml
[dependencies]
wgpu = "22"
pollster = "0.3"

// src/main.rs
use wgpu::util::DeviceExt;

async fn run() {
    let instance = wgpu::Instance::default();
    let adapter = instance.request_adapter(&wgpu::RequestAdapterOptions {
        power_preference: wgpu::PowerPreference::HighPerformance,
        ..Default::default()
    }).await.unwrap();

    let (device, queue) = adapter.request_device(&Default::default(), None)
        .await.unwrap();

    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
        label: None,
        source: wgpu::ShaderSource::Wgsl(include_str!("vadd.wgsl").into()),
    });

    // ... buffer creation, pipeline, dispatch (same pattern as JS API)
}

fn main() {
    pollster::block_on(run());
}

SYCL and Intel oneAPI

SYCL is a Khronos standard for heterogeneous compute built on modern C++. Intel’s oneAPI implements it and targets Intel GPUs (Arc, Xe), CPUs, and FPGAs. It also has backends for AMD (via HIP) and NVIDIA (via CUDA).

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30


// SYCL vector addition
#include <sycl/sycl.hpp>
#include <vector>

int main() {
    const int N = 1024 * 1024;
    std::vector<float> h_a(N), h_b(N), h_c(N);
    for (int i = 0; i < N; i++) { h_a[i] = i; h_b[i] = i * 2.0f; }

    sycl::queue q(sycl::gpu_selector_v);
    std::cout << "Using: " << q.get_device().get_info<sycl::info::device::name>() << "\n";

    {
        sycl::buffer<float> buf_a(h_a.data(), N);
        sycl::buffer<float> buf_b(h_b.data(), N);
        sycl::buffer<float> buf_c(h_c.data(), N);

        q.submit([&](sycl::handler& h) {
            auto a = buf_a.get_access<sycl::access::mode::read>(h);
            auto b = buf_b.get_access<sycl::access::mode::read>(h);
            auto c = buf_c.get_access<sycl::access::mode::write>(h);

            h.parallel_for(sycl::range<1>(N), [=](sycl::id<1> i) {
                c[i] = a[i] + b[i];
            });
        });
    }  // buffers go out of scope, result copied back

    printf("c[0] = %.1f\n", h_c[0]);
}

1
2
3
4
5


# Compile with Intel oneAPI
icpx -fsycl -o vadd vadd.cpp

# For AMD target
icpx -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx1100 -o vadd vadd.cpp

Abstraction Layers: Write Once, Run Everywhere

For most applications, you don’t want to target a single GPU API. Several frameworks abstract over the underlying APIs:

PyTorch (Python/C++)

PyTorch is the de facto standard for ML and has backends for CUDA, ROCm, Metal (MPS), and CPU. From Python, the API is identical regardless of backend:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13


import torch

# Automatically uses CUDA, ROCm, or MPS based on availability
device = (
    "cuda" if torch.cuda.is_available()
    else "mps" if torch.backends.mps.is_available()
    else "cpu"
)
print(f"Using: {device}")

x = torch.randn(4096, 4096, device=device)
y = torch.randn(4096, 4096, device=device)
z = x @ y  # dispatches to cuBLAS, rocBLAS, or MPS depending on device

ArrayFire

ArrayFire provides a high-level array API with OpenCL, CUDA, and CPU backends:

1
2
3
4
5
6
7
8


#include <arrayfire.h>

int main() {
    af::array a = af::randu(1024, 1024);  // on GPU automatically
    af::array b = af::randu(1024, 1024);
    af::array c = af::matmul(a, b);       // uses best available backend
    af::print("result", c(af::seq(3), af::seq(3)));
}

Taichi

Taichi is a Python-embedded DSL for compute-intensive workloads with backends for CUDA, Metal, Vulkan, OpenCL, and CPU:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17


import taichi as ti
ti.init(arch=ti.gpu)  # auto-selects best available GPU backend

@ti.kernel
def vector_add(a: ti.template(), b: ti.template(), c: ti.template()):
    for i in a:
        c[i] = a[i] + b[i]

N = 1024 * 1024
a = ti.field(dtype=ti.f32, shape=N)
b = ti.field(dtype=ti.f32, shape=N)
c = ti.field(dtype=ti.f32, shape=N)

a.from_numpy(np.arange(N, dtype=np.float32))
b.from_numpy(np.arange(N, dtype=np.float32) * 2)
vector_add(a, b, c)
print(c[0])

Choosing the Right Platform

Situation	Recommendation
AMD GPU, ML/HPC workloads	ROCm + HIP, PyTorch ROCm
Apple Silicon, any workload	Metal Compute / MLX
Must run on any GPU (including integrated)	OpenCL or WebGPU (wgpu)
Graphics + compute in same pipeline	Vulkan Compute
Browser-based GPU compute	WebGPU (JavaScript/WGSL)
Intel Arc GPU	SYCL / oneAPI
Python ML, don’t care about backend	PyTorch (auto-selects CUDA/ROCm/MPS)
Cross-platform desktop app	wgpu (Rust) or Taichi (Python)
Existing CUDA code, move to AMD	hipify-clang → HIP

The good news for 2026: the ecosystem is much more mature than it was five years ago. ROCm is production-quality for training. Metal performance is genuinely competitive. WebGPU is shipping in all major browsers. You’re no longer forced to use NVIDIA hardware for serious GPU compute — you have real options, and they’re actually good.

GPU Programming Without CUDA: OpenCL, ROCm, Vulkan, Metal, and WebGPU

GPU Programming Without CUDA: OpenCL, ROCm, Vulkan, Metal, and WebGPU

The GPU Compute Landscape

OpenCL: The Portable Option

The OpenCL Programming Model

A Complete OpenCL Example: Vector Addition

OpenCL Memory Model

OpenCL 3.0 and the Current State

ROCm and HIP: AMD’s CUDA Alternative

HIP Programming Model

Converting CUDA to HIP

ROCm ML Libraries

Supported Hardware

Vulkan Compute: Graphics Pipeline + Compute

Compute Shaders in GLSL

Vulkan Compute Dispatch (Simplified)

Higher-Level Vulkan Compute: kompute

Metal Compute: Apple Silicon

Metal Shading Language

Metal Compute in Swift

MLX: ML on Apple Silicon

WebGPU: GPU Compute in the Browser (and Beyond)

WebGPU Compute in JavaScript

WGSL: WebGPU Shading Language

wgpu: Native WebGPU

SYCL and Intel oneAPI

Abstraction Layers: Write Once, Run Everywhere

PyTorch (Python/C++)

ArrayFire

Taichi

Choosing the Right Platform

Further Reading

Comments