OpenCL Compiler

openclc is an experimental compiler for opencl compute kernels that supports

1. Colocation of OpenCL kernel source code with c/cpp source code.
2. Build time compilation of kernels to SPIR-V (or PTX/Metal IR in the future)
3. Automatic device discovery
4. CUDA kernel invocation syntax

These features make for far easier use of data parallel hardware shipped with most devices.

Roadmap

• Vulkan Target: GLCompute SPIR-V for Vulkan. Code Generation passes already exist in clspv. This exchanges some OpenCL features for increased portability.
• Metal Target: SPIR-V Cross can reflect GLCompute SPIR-V up to MSL source which can be used in the Metal API. This sounds brittle but is the crux of WebGPU shaders running correctly on MacOS.
• CUDA: OpenCL kernels can be compiled to PTX with the LLVM PTX backend. It should also be possible to support many CUDA built-ins, synchronization primitives, and inline PTX.

Usage

vec_add.cl

#include <assert.h>
#include <openclc_rt.h>
#include <stdio.h>

kernel void add(constant float *A, constant float *B, global float *C)
{
    size_t gid = get_global_id(0);
    C[gid] = A[gid] + B[gid];
}

int main()
{
    oclcInit();

    int n = 256;
    size_t sz = n * sizeof(float);

    // Allocate Host Memory
    float* A = (float*)malloc(sz);
    float* B = (float*)malloc(sz);
    float* C = (float*)malloc(sz);

    // Initialize Host Memory
    for (int i = 0; i < n; i++) {
        A[i] = i;
        B[i] = i;
        C[i] = 0;
    }

    // Allocate Device Memory
    float* dA = (float*)oclcMalloc(sz);
    float* dB = (float*)oclcMalloc(sz);
    float* dC = (float*)oclcMalloc(sz);

    // Copy Host Values to the Device Memory
    oclcMemcpy(dA, A, sz, oclcMemcpyHostToDevice);
    oclcMemcpy(dB, B, sz, oclcMemcpyHostToDevice);
    oclcMemcpy(dC, C, sz, oclcMemcpyHostToDevice);

    dim3 gridDim = { n / 32 };
    dim3 blockDim = { 32 };

    // Invoke the Kernel!
    add<<<gridDim, blockDim>>>(dA, dB, dC);

    oclcMemcpy(C, dC, sz, oclcMemcpyDeviceToHost);

    oclcDeviceSynchronize();

    for (int i = 0; i < n; i++) {
        assert(C[i] == 2 * i);
    }

    puts("\n---------------------------\n");
    puts("Passed");
    puts("\n---------------------------\n");

    oclcFree(dA);
    oclcFree(dB);
    oclcFree(dC);
}

openclc vec_add.cl -o vadd

Installation

install.sh will download the lastest release for your os/arch and extract the contents to $HOME/.openclc.

curl -fsSL https://raw.githubusercontent.com/e253/openclc/main/install.sh | bash

Alternatively, download the release archive and extract it anywhere you'd like.

OpenCLC does not need to be installed in any specific location.

Dependencies

• zig-bootstrap : Rev ec2dca85a340f134d2fcfdc9007e91f9abed6996
• SPIRV-LLVM-Translator : Release v18.1.2
• SPIRV-Headers : Rev 4f7b471f1a66b6d06462cd4ba57628cc0cd087d7
• SPIRV-Tools : Release v2024.2
• fmt
• LLVM-18 (Build Dependency) Use llvm.sh to install it with apt. Any 18.x should work.

Build

Openclc only builds on Debian, but can cross compile to many targets with the help of zig.

1. Get LLVM 18
2. Get ZVM, then zig 0.13.0
3. ./build-deps.sh <zig-target-triple> <mcpu>
4. ./build-openclc.sh <zig-target-triple> <mcpu>
5. ./<zig-target-triple>-<mcpu>/bin/openclc

Credits

1. Khronos Group, for the SPIRV-Tools, SPIRV-LLVM-Translator dependencies
2. Google for the clspv project. I couldn't have figured out the clang frontend calls otherwise.