diff --git a/src/compiler/reflection.jl b/src/compiler/reflection.jl
index 97511fa..a8bb714 100644
--- a/src/compiler/reflection.jl
+++ b/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)
diff --git a/test/Project.toml b/test/Project.toml
index 382d28b..ba8f907 100644
--- a/test/Project.toml
+++ b/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"
diff --git a/test/execution.jl b/test/execution.jl
new file mode 100644
index 0000000..a64b542
--- /dev/null
+++ b/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
diff --git a/test/runtests.jl b/test/runtests.jl
index c41a005..bacb223 100644
--- a/test/runtests.jl
+++ b/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