Add a native compiler, using SPIR-V IL. (#222)

.buildkite/pipeline.yml

@@ -3,9 +3,19 @@ steps:
     plugins:
       - JuliaCI/julia#v1:
           version: "1.10"
-      - JuliaCI/julia-test#v1: ~
       - JuliaCI/julia-coverage#v1:
           codecov: true
+    commands: |
+      julia --project -e '
+        using Pkg
+
+        println("--- :julia: Instantiating project")
+        Pkg.develop(path="lib/intrinsics")
+        Pkg.instantiate()
+
+
+        println("+++ :julia: Running tests")
+        Pkg.test(; coverage=true)'
     agents:
       queue: "juliagpu"
       cuda: "*"

codecov.yml → .github/codecov.yml

@@ -1,7 +1,8 @@
 coverage:
   ignore:
-    - "lib/lib*.jl"
-    - "src/kernels"
+    - "lib/*/lib*.jl"
+    - "src/kernels/"
+    - "src/device/"
   status:
     patch: false
     project: false

.github/workflows/CI.yml

@@ -28,7 +28,7 @@ jobs:
           version: ${{ matrix.version }}
           arch: ${{ matrix.arch }}
       - uses: julia-actions/cache@v2
-      - uses: julia-actions/julia-buildpkg@v1
+      - run: julia --project -e 'using Pkg; Pkg.develop(path="lib/intrinsics")'
       - uses: julia-actions/julia-runtest@v1
       - uses: julia-actions/julia-processcoverage@v1
       - uses: codecov/codecov-action@v4

Project.toml

@@ -3,11 +3,17 @@ uuid = "08131aa3-fb12-5dee-8b74-c09406e224a2"
 version = "0.10.0"
 
 [deps]
+Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
+GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
+LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 OpenCL_jll = "6cb37087-e8b6-5417-8430-1f242f1e46e4"
 Printf = "de0858da-6303-5e67-8744-51eddeeeb8d7"
 Reexport = "189a3867-3050-52da-a836-e630ba90ab69"
+SPIRVIntrinsics = "71d1d633-e7e8-4a92-83a1-de8814b09ba8"
+SPIRV_LLVM_Translator_unified_jll = "85f0d8ed-5b39-5caa-b1ae-7472de402361"
 
 [compat]
+LLVM = "9.1"
 OpenCL_jll = "2024.5.8"
 julia = "1.10"

README.md

@@ -58,38 +58,64 @@ Available platforms: 3
 
 ## Basic example: vector add
 
+The traditional way of using OpenCL is by writing kernel source code in OpenCL C. For
+example, a simple vector addition:
+
 ```julia
-using LinearAlgebra
 using OpenCL, pocl_jll
 
-const sum_kernel = "
-   __kernel void sum(__global const float *a,
-                     __global const float *b,
-                     __global float *c)
-    {
+const source = """
+   __kernel void vadd(__global const float *a,
+                      __global const float *b,
+                      __global float *c) {
       int gid = get_global_id(0);
       c[gid] = a[gid] + b[gid];
-    }
-"
+    }"""
+
 a = rand(Float32, 50_000)
 b = rand(Float32, 50_000)
 
-a_buff = cl.Buffer(Float32, length(a), (:r, :copy), hostbuf=a)
-b_buff = cl.Buffer(Float32, length(b), (:r, :copy), hostbuf=b)
-c_buff = cl.Buffer(Float32, length(a), :w)
+d_a = CLArray(a; access=:r)
+d_b = CLArray(b; access=:r)
+d_c = similar(d_a; access=:w)
+
+p = cl.Program(; source) |> cl.build!
+k = cl.Kernel(p, "vadd")
 
-p = cl.Program(source=sum_kernel) |> cl.build!
-k = cl.Kernel(p, "sum")
+clcall(k, Tuple{Ptr{Float32}, Ptr{Float32}, Ptr{Float32}},
+       d_a, d_b, d_c; global_size=size(a))
 
-cl.call(k, a_buff, b_buff, c_buff; global_size=size(a))
+c = Array(d_c)
+
+@assert a + b ≈ c
+```
 
-r = cl.read(c_buff)
 
-if isapprox(norm(r - (a+b)), zero(Float32))
-    @info "Success!"
-else
-    @error "Norm should be 0.0f"
+## Native example: vector add
+
+If your platform supports SPIR-V, it's possible to use Julia functions as kernels:
+
+```julia
+using OpenCL, pocl_jll
+
+function vadd(a, b, c)
+    gid = get_global_id(1)
+    @inbounds c[gid] = a[gid] + b[gid]
+    return
 end
+
+a = rand(Float32, 50_000)
+b = rand(Float32, 50_000)
+
+d_a = CLArray(a; access=:r)
+d_b = CLArray(b; access=:r)
+d_c = similar(d_a; access=:w)
+
+@opencl global_size=size(a) vadd(d_a, d_b, d_c)
+
+c = Array(d_c)
+
+@assert a + b ≈ c
 ```
 
 

examples/vadd.jl

@@ -0,0 +1,27 @@
+using OpenCL, pocl_jll, Test
+
+const source = """
+   __kernel void vadd(__global const float *a,
+                      __global const float *b,
+                      __global float *c) {
+      int i = get_global_id(0);
+      c[i] = a[i] + b[i];
+    }"""
+
+dims = (2,)
+a = round.(rand(Float32, dims) * 100)
+b = round.(rand(Float32, dims) * 100)
+c = similar(a)
+
+d_a = CLArray(a)
+d_b = CLArray(b)
+d_c = CLArray(c)
+
+prog = cl.Program(; source) |> cl.build!
+kern = cl.Kernel(prog, "vadd")
+
+len = prod(dims)
+clcall(kern, Tuple{Ptr{Float32}, Ptr{Float32}, Ptr{Float32}},
+       d_a, d_b, d_c; global_size=(len,))
+c = Array(d_c)
+@test a+b ≈ c

examples/vadd_native.jl

@@ -0,0 +1,21 @@
+using OpenCL, pocl_jll, Test
+
+function vadd(a, b, c)
+    i = get_global_id()
+    @inbounds c[i] = a[i] + b[i]
+    return
+end
+
+dims = (2,)
+a = round.(rand(Float32, dims) * 100)
+b = round.(rand(Float32, dims) * 100)
+c = similar(a)
+
+d_a = CLArray(a)
+d_b = CLArray(b)
+d_c = CLArray(c)
+
+len = prod(dims)
+@opencl global_size=len vadd(d_a, d_b, d_c)
+c = Array(d_c)
+@test a+b ≈ c

lib/api.jl → lib/cl/api.jl

@@ -69,7 +69,7 @@ function retry_reclaim(f, isfailed)
     ret
 end
 
-include("../lib/libopencl.jl")
+include("libopencl.jl")
 
 # lazy initialization
 const initialized = Ref{Bool}(false)

lib/kernel.jl → lib/cl/kernel.jl

@@ -181,8 +181,13 @@ function enqueue_kernel(k::Kernel, global_work_size, local_work_size=nothing;
 end
 
 function call(k::Kernel, args...; global_size=(1,), local_size=nothing,
-              global_work_offset=nothing, wait_on::Vector{Event}=Event[])
+              global_work_offset=nothing, wait_on::Vector{Event}=Event[],
+              svm_pointers::Vector{Ptr{Cvoid}}=Ptr{Cvoid}[])
     set_args!(k, args...)
+    if !isempty(svm_pointers)
+        clSetKernelExecInfo(k, CL_KERNEL_EXEC_INFO_SVM_PTRS,
+                            sizeof(svm_pointers), svm_pointers)
+    end
     enqueue_kernel(k, global_size, local_size; global_work_offset, wait_on)
 end
 
@@ -215,8 +220,8 @@ clcall(f::F, types::Tuple, args::Vararg{Any,N}; kwargs...) where {N,F} =
     clcall(f, _to_tuple_type(types), args...; kwargs...)
 
 function clcall(k::Kernel, types::Type{T}, args::Vararg{Any,N}; kwargs...) where {T,N}
-    call_closure = function (pointers::Vararg{Any,N})
-        call(k, pointers...; kwargs...)
+    call_closure = function (converted_args::Vararg{Any,N})
+        call(k, converted_args...; kwargs...)
     end
     convert_arguments(call_closure, types, args...)
 end

lib/intrinsics/Project.toml

@@ -0,0 +1,9 @@
+name = "SPIRVIntrinsics"
+uuid = "71d1d633-e7e8-4a92-83a1-de8814b09ba8"
+authors = ["Tim Besard <[email protected]>"]
+version = "0.1.0"
+
+[deps]
+ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04"
+LLVM = "929cbde3-209d-540e-8aea-75f648917ca0"
+SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b"

lib/intrinsics/src/SPIRVIntrinsics.jl

@@ -0,0 +1,25 @@
+module SPIRVIntrinsics
+
+using LLVM, LLVM.Interop
+using Core: LLVMPtr
+
+import ExprTools
+
+import SpecialFunctions
+
+include("pointer.jl")
+include("utils.jl")
+
+# OpenCL intrinsics
+#
+# we currently don't implement SPIR-V intrinsics directly, but rely on
+# the SPIR-V to LLVM translator supporting OpenCL intrinsics
+include("work_item.jl")
+include("synchronization.jl")
+include("memory.jl")
+include("printf.jl")
+include("math.jl")
+include("integer.jl")
+include("atomic.jl")
+
+end