Add some tests.

src/compiler/reflection.jl

@@ -28,7 +28,7 @@ for method in (:code_typed, :code_warntype, :code_llvm, :code_native)
  kernel::Bool=false, kwargs...)
  compiler_kwargs, kwargs = split_kwargs_runtime(kwargs, COMPILER_KWARGS)
  source = methodinstance(typeof(func), Base.to_tuple_type(types))
- config = compiler_config(device(); kernel, compiler_kwargs...)
+ config = compiler_config(cl.device(); kernel, compiler_kwargs...)
  job = CompilerJob(source, config)
  GPUCompiler.$method($(args...); kwargs...)
  end
@@ -66,7 +66,7 @@ Return a type `r` such that `f(args...)::r` where `args::tt`.
 """
 function return_type(@nospecialize(func), @nospecialize(tt))
  source = methodinstance(typeof(func), tt)
- config = compiler_config(device())
+ config = compiler_config(cl.device())
  job = CompilerJob(source, config)
  interp = GPUCompiler.get_interpreter(job)
  sig = Base.signature_type(func, tt)

test/Project.toml

@@ -1,4 +1,5 @@
 [deps]
+IOCapture = "b5f81e59-6552-4d32-b1f0-c071b021bf89"
 LinearAlgebra = "37e2e46d-f89d-539d-b4ee-838fcccc9c8e"
 Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
 pocl_jll = "627d6b7a-bbe6-5189-83e7-98cc0a5aeadd"

test/execution.jl

@@ -0,0 +1,124 @@
+if !in("cl_khr_il_program", cl.device().extensions)
+@warn "Skipping execution tests on $(cl.platform().name)"
+else
+
+@testset "execution" begin
+
+@testset "@opencl" begin
+
+dummy() = nothing
+
+@test_throws UndefVarError @opencl undefined()
+@test_throws MethodError @opencl dummy(1)
+
+
+@testset "launch configuration" begin
+ @opencl dummy()
+
+ global_size = 1
+ @opencl global_size dummy()
+ @opencl global_size=1 dummy()
+ @opencl global_size=(1,1) dummy()
+ @opencl global_size=(1,1,1) dummy()
+
+ local_size = 1
+ @opencl global_size local_size dummy()
+ @opencl global_size=1 local_size=1 dummy()
+ @opencl global_size=(1,1) local_size=(1,1) dummy()
+ @opencl global_size=(1,1,1) local_size=(1,1,1) dummy()
+
+ @test_throws ArgumentError @opencl global_size=(1,) local_size=(1,1) dummy()
+ @test_throws InexactError @opencl global_size=(-2) dummy()
+ @test_throws InexactError @opencl local_size=(-2) dummy()
+end
+
+@testset "launch=false" begin
+ # XXX: how are svm_pointers handled here?
+ k = @opencl launch=false dummy()
+ k()
+ k(; global_size=1)
+end
+
+@testset "inference" begin
+ foo() = @opencl dummy()
+ @inferred foo()
+
+ # with arguments, we call clconvert
+ kernel(a) = return
+ bar(a) = @opencl kernel(a)
+ @inferred bar(CLArray([1]))
+end
+
+
+@testset "reflection" begin
+ OpenCL.code_lowered(dummy, Tuple{})
+ OpenCL.code_typed(dummy, Tuple{})
+ OpenCL.code_warntype(devnull, dummy, Tuple{})
+ OpenCL.code_llvm(devnull, dummy, Tuple{})
+ OpenCL.code_native(devnull, dummy, Tuple{})
+
+ @device_code_lowered @opencl dummy()
+ @device_code_typed @opencl dummy()
+ @device_code_warntype io=devnull @opencl dummy()
+ @device_code_llvm io=devnull @opencl dummy()
+ @device_code_native io=devnull @opencl dummy()
+
+ mktempdir() do dir
+ @device_code dir=dir @opencl dummy()
+ end
+
+ @test_throws ErrorException @device_code_lowered nothing
+
+ # make sure kernel name aliases are preserved in the generated code
+ @test occursin("dummy", sprint(io->(@device_code_llvm io=io optimize=false @opencl dummy())))
+ @test occursin("dummy", sprint(io->(@device_code_llvm io=io @opencl dummy())))
+ @test occursin("dummy", sprint(io->(@device_code_native io=io @opencl dummy())))
+
+ # make sure invalid kernels can be partially reflected upon
+ let
+ invalid_kernel() = throw()
+ @test_throws OpenCL.InvalidIRError @opencl invalid_kernel()
+ @test_throws OpenCL.InvalidIRError IOCapture.capture() do
+ @device_code_warntype @opencl invalid_kernel()
+ end
+ c = IOCapture.capture() do
+ try
+ @device_code_warntype @opencl invalid_kernel()
+ catch
+ end
+ end
+ @test occursin("Body::Union{}", c.output)
+ end
+
+ # set name of kernel
+ @test occursin("mykernel", sprint(io->(@device_code_llvm io=io begin
+ @opencl name="mykernel" dummy()
+ end)))
+
+ @test OpenCL.return_type(identity, Tuple{Int}) === Int
+ @test OpenCL.return_type(sin, Tuple{Float32}) === Float32
+ @test OpenCL.return_type(getindex, Tuple{CLDeviceArray{Float32,1,AS.Global},Int32}) === Float32
+ @test OpenCL.return_type(getindex, Tuple{Base.RefValue{Integer}}) === Integer
+end
+
+end
+
+###############################################################################
+
+@testset "argument passing" begin
+
+function memset(a, val)
+ gid = get_global_id(1)
+ @inbounds a[gid] = val
+ return
+end
+
+a = CLArray{Int}(undef, 10)
+@opencl global_size=length(a) memset(a, 42)
+@test all(Array(a) .== 42)
+
+end
+
+end
+
+end

test/runtests.jl

@@ -2,16 +2,19 @@ using Test
 using OpenCL
 using pocl_jll
 
+using IOCapture
+
 @info "System information:\n" * sprint(io->OpenCL.versioninfo(io))
 
 @testset "OpenCL.jl" begin
 
 @testset "$(platform.name): $(device.name)" for platform in cl.platforms(),
-  device in cl.devices(platform)
+ device in cl.devices(platform)
 
 cl.platform!(platform)
 cl.device!(device)
 
+# libopencl wrappers
 include("platform.jl")
 include("context.jl")
 include("device.jl")
@@ -22,7 +25,9 @@ include("kernel.jl")
 include("behaviour.jl")
 include("memory.jl")
 include("buffer.jl")
+
 include("array.jl")
+include("execution.jl")
 
 end