Skip to content
Earn a certificate
AssemblyBashCC#C++ClojureCOBOLCUDADartElixirErlangF#FortranGoGroovyHaskellJavaJavaScriptJuliaKotlinLispLuaMATLABMojoObjective-COCamlOpenCLPerlPHPPowerShellPrologPythonRRubyRustScalaSoliditySQLSwiftTritonTypeScriptVerilogVHDLVisual BasicWebAssemblyZig
Volume O · The TensorCore.si Library

OpenCL

Parallel computing on any GPU.

OpenCL (Open Computing Language) is an open standard for running parallel code on GPUs, CPUs, FPGAs and other accelerators from many vendors. Kernels are written in a C-based language and launched from a host program. It is the vendor-neutral alternative to NVIDIA's CUDA.

Prologue

Why learn OpenCL?

OpenCL teaches GPU programming without locking you into one vendor, and the same ideas carry over to CUDA, SYCL and Vulkan compute.

  • Portable: one kernel can run on many kinds of hardware.
  • Close to the metal: control work-groups and local memory directly.
  • Foundation: concepts map one-to-one onto CUDA blocks and threads.
Prologue

Set up and run your first program

Install your GPU vendor's OpenCL driver, plus the headers and the Python binding PyOpenCL for quick experiments.

terminal
sudo apt install ocl-icd-opencl-dev clinfo
clinfo | head
pip install pyopencl numpy
python hello.py
hello.py
import pyopencl as cl, numpy as np

ctx = cl.create_some_context(interactive=False)
queue = cl.CommandQueue(ctx)
print("Hello, AI! OpenCL device:", ctx.devices[0].name)
TipPyOpenCL handles the host-side setup so you can focus on kernels.
Chapter 1

Platforms, devices and work-items

A kernel runs once per work-item. Work-items are grouped into work-groups that share fast local memory. get_global_id(0) tells each work-item which element it owns (like CUDA's thread index).

add.cl
__kernel void add(__global const float *a,
                  __global const float *b,
                  __global float *out,
                  const int n) {
    int i = get_global_id(0);
    if (i < n) out[i] = a[i] + b[i];
}
Chapter 2

The host program

The host creates buffers, compiles the kernel at runtime, sets arguments and queues the kernel over a global range.

host.py
import pyopencl as cl, numpy as np
ctx = cl.create_some_context(interactive=False); q = cl.CommandQueue(ctx)
n = 1_000_000
a = np.random.rand(n).astype(np.float32); b = np.random.rand(n).astype(np.float32)
mf = cl.mem_flags
a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
out_g = cl.Buffer(ctx, mf.WRITE_ONLY, a.nbytes)
prg = cl.Program(ctx, open("add.cl").read()).build()
prg.add(q, (n,), None, a_g, b_g, out_g, np.int32(n))
out = np.empty_like(a); cl.enqueue_copy(q, out, out_g)
print(np.allclose(out, a + b))
Chapter 3

Memory spaces

OpenCL has four memory spaces: __global (large, slow), __constant (read-only, cached), __local (shared within a work-group) and __private (per work-item registers).

scale.cl
__kernel void scale(__global float *x,
                    __constant float *factor,   // read-only and cached
                    const int n) {
    int i = get_global_id(0);
    float v = x[i];          // private register
    if (i < n) x[i] = v * factor[0];
}
Chapter 4

Local memory and reductions

To sum a large array, each work-group loads data into local memory, synchronises with barrier, and halves the active work-items each step.

sum.cl
__kernel void sum(__global const float *in, __global float *partial,
                  __local float *tmp, const int n) {
    int gid = get_global_id(0), lid = get_local_id(0);
    tmp[lid] = gid < n ? in[gid] : 0.0f;
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int s = get_local_size(0) / 2; s > 0; s >>= 1) {
        if (lid < s) tmp[lid] += tmp[lid + s];
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    if (lid == 0) partial[get_group_id(0)] = tmp[0];
}
Chapter 5

Matrix multiplication

Each work-item computes one output element using a 2D range. Tiling with local memory (like CUDA shared memory) makes it much faster.

matmul.cl
__kernel void matmul(__global const float *A, __global const float *B,
                     __global float *C, const int N) {
    int row = get_global_id(1), col = get_global_id(0);
    if (row >= N || col >= N) return;
    float acc = 0.0f;
    for (int k = 0; k < N; k++)
        acc += A[row * N + k] * B[k * N + col];
    C[row * N + col] = acc;
}
// host: prg.matmul(q, (N, N), (16, 16), A_g, B_g, C_g, np.int32(N))
Part II · AI with OpenCL

OpenCL and AI

Most deep learning today targets CUDA or vendor libraries, but OpenCL still matters for portable inference on mobile GPUs (Android), embedded boards, Intel GPUs and FPGAs. llama.cpp has an OpenCL backend for Qualcomm Adreno GPUs, and frameworks such as TensorFlow Lite and Apache TVM have used OpenCL for on-device AI. Its successor ideas live on in SYCL and Vulkan compute.

The essential libraries

PyOpenCLPython host API with NumPy integration.
CLBlastTuned BLAS (matrix maths) for OpenCL.
Apache TVMCompiles models to OpenCL and other targets.
llama.cpp (OpenCL)LLM inference on mobile GPUs.

A ReLU + bias layer on any GPU

A fused neural network activation that runs on NVIDIA, AMD, Intel or Apple-era OpenCL devices.

layer.py
import pyopencl as cl, numpy as np
src = """
__kernel void bias_relu(__global float *x, __global const float *b, const int cols) {
    int i = get_global_id(0);
    float v = x[i] + b[i % cols];
    x[i] = v > 0.0f ? v : 0.0f;
}"""
ctx = cl.create_some_context(interactive=False); q = cl.CommandQueue(ctx)
x = np.random.randn(1024, 512).astype(np.float32); b = np.random.randn(512).astype(np.float32)
mf = cl.mem_flags
xg = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=x)
bg = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
cl.Program(ctx, src).build().bias_relu(q, (x.size,), None, xg, bg, np.int32(512))
out = np.empty_like(x); cl.enqueue_copy(q, out, xg)
print(np.allclose(out, np.maximum(x + b, 0)))
Part III · Build an AI agent

An agent with a GPU tool

Agents can offload number crunching to OpenCL. Here a local model calls a GPU vector-sum tool built on the reduction idea.

Your goal→Model plans→Tool call→Your code runs the tool→Result back to model→Final answer
agent.py
import json, numpy as np, pyopencl as cl, pyopencl.array as cla
from openai import OpenAI

ctx = cl.create_some_context(interactive=False); q = cl.CommandQueue(ctx)
def gpu_stats(values):
    a = cla.to_device(q, np.array(values, dtype=np.float32))
    return {"sum": float(cla.sum(a).get()), "max": float(cla.max(a).get()), "count": len(values)}

client = OpenAI(base_url="http://localhost:11434/v1", api_key="ollama")
tools = [{"type": "function", "function": {"name": "gpu_stats", "description": "Sum, max and count of numbers, computed on the GPU",
    "parameters": {"type": "object", "properties": {"values": {"type": "array", "items": {"type": "number"}}}, "required": ["values"]}}}]
messages = [{"role": "user", "content": "Total and peak of these GPU hours: 12, 40.5, 33, 8, 61"}]
for _ in range(5):
    msg = client.chat.completions.create(model="llama3.1", messages=messages, tools=tools).choices[0].message
    messages.append(msg)
    if not msg.tool_calls:
        print(msg.content); break
    for c in msg.tool_calls:
        messages.append({"role": "tool", "tool_call_id": c.id, "content": json.dumps(gpu_stats(**json.loads(c.function.arguments)))})

pyopencl.array provides GPU reductions such as sum and max out of the box.

Run it locallyInstall Ollama, run ollama pull llama3.1, and keep it running. The code talks to its OpenAI-compatible API at http://localhost:11434/v1, so the same program also works with vLLM, llama.cpp server or a cloud API by changing the URL, model name and adding an API key.
Part IV · Your path

OpenCL learning roadmap

Tick items off as you go. Your progress is saved in this browser.

Beginner

Intermediate

Advanced

Part IV · Your path

Projects to build

Image blurA Gaussian blur kernel on photos.
MandelbrotRender fractals on any GPU.
Mini NN layerA matmul + activation layer, checked against NumPy.
Reference

OpenCL cheat sheet

TaskSyntax
Kernel__kernel void f(__global float *x)
Global idint i = get_global_id(0);
Local idget_local_id(0)
Sync groupbarrier(CLK_LOCAL_MEM_FENCE);
Launch (Py)prg.f(q, (n,), None, buf)
Copy backcl.enqueue_copy(q, out, buf)
Reference

Common mistakes

  • The global size must be a multiple of the local size in OpenCL 1.x.
  • Kernel compile errors appear at runtime; print the build log.
  • Moving data between host and device often costs more than the computation.
Reference

Where to go next

On TensorCore.si: CUDA volume, C volume.

Keyboard: [ previous language · ] next language · / search the library