get cu string compilation working

Author: SimonDanisch

src/backends/cudanative/compilation.jl

@@ -1,24 +1,34 @@
-# EXCLUDE FROM TESTING
+#TODO tag CUDArt#103 and use that
 import CUDArt
+export CompileError
 
 # Generate a temporary file with specific suffix
 # NOTE: mkstemps is glibc 2.19+, so emulate its behavior
 function mkstemps(suffix::AbstractString)
     base = tempname()
     filename = base * suffix
+
     # make sure the filename is unique
     i = 0
     while isfile(filename)
         i += 1
         filename = base * ".$i" * suffix
     end
+
     return (filename, Base.open(filename, "w"))
 end
 
+macro compile(dev, kernel, code)
+    kernel_name = string(kernel)
+    containing_file = @__FILE__
+
+    return Expr(:toplevel,
+        Expr(:export,esc(kernel)),
+        :($(esc(kernel)) = _compile($(esc(dev)), $kernel_name, $code, $containing_file)))
+end
 
-type CompileError <: Base.WrappedException
+immutable CompileError <: Exception
     message::String
-    error
 end
 
 const builddir = joinpath(@__DIR__, ".cache")
@@ -31,7 +41,7 @@ function _compile(dev, kernel, code, containing_file)
         mkpath(builddir)
     end
 
-    # Check if we need to compile
+    # check if we need to compile
     codehash = hex(hash(code))
     output = "$builddir/$(kernel)_$(codehash)-$(arch).ptx"
     if isfile(output)
@@ -40,47 +50,44 @@ function _compile(dev, kernel, code, containing_file)
         need_compile = true
     end
 
-    # Compile the source, if necessary
+    # compile the source, if necessary
     if need_compile
-        # Write the source into a compilable file
+        # write the source to a compilable file
         (source, io) = mkstemps(".cu")
         write(io, """
 extern "C"
 {
 $code
 }
 """)
-        close(io)
+        Base.close(io)
 
         compile_flags = vcat(CUDArt.toolchain_flags, ["--gpu-architecture", arch])
-        try
-            # TODO: capture STDERR
-            run(pipeline(`$(CUDArt.toolchain_nvcc) $(compile_flags) -ptx -o $output $source`, stderr=DevNull))
-        catch ex
-            isa(ex, ErrorException) || rethrow(ex)
-            rethrow(CompileError("compilation of kernel $kernel failed (typo in C++ source?)", ex))
-        finally
-            rm(source)
+        err = Pipe()
+        cmd = `$(CUDArt.toolchain_nvcc) $(compile_flags) -ptx -o $output $source`
+        result = success(pipeline(cmd; stdout=DevNull, stderr=err))
+        Base.close(err.in)
+        rm(source)
+
+        errors = readstring(err)
+        if !result
+            throw(CompileError("compilation of kernel $kernel failed\n$errors"))
+        elseif !isempty(errors)
+            warn("during compilation of kernel $kernel:\n$errors")
         end
 
         if !isfile(output)
             error("compilation of kernel $kernel failed (no output generated)")
         end
     end
 
-    # Pass the module to the CUDA driver
-    mod = try
-        CUDAdrv.CuModuleFile(output)
-    catch ex
-        rethrow(CompileError("loading of kernel $kernel failed (invalid CUDA code?)", ex))
-    end
+    mod = CUDAdrv.CuModuleFile(output)
+    return CUDAdrv.CuFunction(mod, kernel)
+end
 
-    # Load the function pointer
-    func = try
-        CUDAdrv.CuFunction(mod, kernel)
-    catch ex
-        rethrow(CompileError("could not find kernel $kernel in the compiled binary (wrong function name?)", ex))
+function clean_cache()
+    if ispath(builddir)
+        @assert isdir(builddir)
+        rm(builddir; recursive=true)
     end
-
-    return func
 end

src/backends/cudanative/cudanative.jl

@@ -209,7 +209,7 @@ immutable CUFunction{T}
     kernel::T
 end
 
-if try success(`nvcc --version`); catch false; end
+if try success(`$(CUDArt.toolchain_nvcc) --version`); catch false; end
     include("compilation.jl")
     hasnvcc() = true
 else
@@ -236,16 +236,20 @@ function (f::CUFunction{F}){F <: Function, T, N}(A::CUArray{T, N}, args...)
         f.kernel, map(unpack_cu_array, args)...
     )
 end
-function cu_convert{T, N}(x::CUArray{T, N})
-    pointer(buffer(x))
-end
-cu_convert(x) = x
+
+cudacall_types(x::CUArray{T, N}) where {T, N} = Ptr{T}
+cudacall_types(x::T) where T = T
+
+cudacall_convert(x) = x
+cudacall_convert(x::CUArray{T, N}) where {T, N} = pointer(buffer(x))
 
 function (f::CUFunction{F}){F <: CUDAdrv.CuFunction, T, N}(A::CUArray{T, N}, args)
     griddim, blockdim = thread_blocks_heuristic(A)
-    CUDAdrv.launch(
-        f.kernel, CUDAdrv.CuDim3(griddim...), CUDAdrv.CuDim3(blockdim...), 0, CuDefaultStream(),
-        map(cu_convert, args)
+    typs = Tuple{cudacall_types.(args)...}
+    cuargs = cudacall_convert.(args)
+    CUDAdrv.cudacall(
+        f.kernel, CUDAdrv.CuDim3(griddim...), CUDAdrv.CuDim3(blockdim...),
+        typs, cuargs...
     )
 end
 

test/cudanative.jl

@@ -73,20 +73,20 @@ end
     jy = Array(y)
     @test map!(sin, jy, jy) ≈ Array(x)
 end
-#
-# if CUBackend.hasnvcc()
-#     @testset "Custom kernel from string function" begin
-#         x = GPUArray(rand(Float32, 100))
-#         y = GPUArray(rand(Float32, 100))
-#         source = """
-#         __global__ void copy(const float *input, float *output)
-#         {
-#             int i = blockIdx.x * blockDim.x + threadIdx.x;
-#             output[i] = input[i];
-#         }
-#         """
-#         f = (source, :copy)
-#         gpu_call(f, x, (x, y))
-#         @test Array(x) == Array(y)
-#     end
-# end
+
+if CUBackend.hasnvcc()
+    @testset "Custom kernel from string function" begin
+        x = GPUArray(rand(Float32, 100))
+        y = GPUArray(rand(Float32, 100))
+        source = """
+        __global__ void copy(const float *input, float *output)
+        {
+            int i = blockIdx.x * blockDim.x + threadIdx.x;
+            output[i] = input[i];
+        }
+        """
+        f = (source, :copy)
+        gpu_call(f, x, (x, y))
+        @test Array(x) == Array(y)
+    end
+end