get many more Knet tests working

Author: SimonDanisch

REQUIRE

@@ -1 +1,29 @@
 julia 0.6
+CLFFT
+CLBLAS
+OpenCL
+Transpiler
+Sugar
+GPUArrays
+
+
+julia 0.6
+StaticArrays
+ColorTypes
+
+Transpiler 0.3
+Sugar 0.3
+Matcha 0.0.2
+
+CUDAnative 0.4.1 # llvm codegen fix
+CUDAdrv 0.5.1
+CUDArt 0.4.0 # for cuda c compiler support
+CUBLAS 0.2.0 # for CUDAdrv support
+CUFFT
+
+OpenCL 0.6.0 #proper packed conversion
+CLBLAS 1.1.0
+CLFFT 0.4.0 # 0.5.0
+
+Interpolations
+IterTools

src/CLArrays.jl

@@ -11,13 +11,13 @@ function context end
 
 include("memory.jl")
 include("array.jl")
+include("ondevice.jl")
 include("device.jl")
 include("context.jl")
 include("intrinsics.jl")
 include("compilation.jl")
 include("mapreduce.jl")
 include("3rdparty.jl")
-include("device_funcs.jl")
 
 export CLArray
 

src/array.jl

@@ -8,79 +8,12 @@ using Sugar: to_tuple
 import Base: pointer, similar, size, copy!, convert
 using Base: RefValue
 
-# pointer MUST be a type parameter, to make it easier to replace it with a non pointer type for host upload
 mutable struct CLArray{T, N} <: GPUArray{T, N}
     ptr::OwnedPtr{T}
     size::NTuple{N, Cuint}
 end
 
 
-struct HostPtr{T}
-    ptr::Int32
-    (::Type{HostPtr{T}})() where T = new{T}(Int32(0))
-end
-Base.eltype(::Type{HostPtr{T}}) where T = T
-
-struct DeviceArray{T, N, Ptr} <: AbstractArray{T, N}
-    ptr::Ptr
-    size::NTuple{N, Cuint}
-end
-const PreDeviceArray{T, N} = DeviceArray{T, N, HostPtr{T}}
-const OnDeviceArray{T, N} = DeviceArray{T, N, GlobalPointer{T}}
-
-kernel_convert(A::CLArray{T, N}) where {T, N} = PreDeviceArray{T, N}(HostPtr{T}(), A.size)
-predevice_type(::Type{OnDeviceArray{T, N}}) where {T, N} = PreDeviceArray{T, N}
-device_type(::CLArray{T, N}) where {T, N} = OnDeviceArray{T, N}
-reconstruct(x::PreDeviceArray{T, N}, ptr::GlobalPointer{T}) where {T, N} = OnDeviceArray{T, N}(ptr, x.size)
-
-kernel_convert(x::RefValue{T}) where T <: CLArray = RefValue(kernel_convert(x[]))
-predevice_type(::Type{RefValue{T}}) where T <: OnDeviceArray = RefValue{predevice_type(T)}
-device_type(x::RefValue{T}) where T <: CLArray = RefValue{device_type(x[])}
-reconstruct(x::RefValue{T}, ptr::GlobalPointer) where T <: PreDeviceArray = RefValue(reconstruct(x[], ptr))
-
-kernel_convert(x::Tuple) = kernel_convert.(x)
-predevice_type(::Type{T}) where T <: Tuple = Tuple{predevice_type.((T.parameters...))...}
-device_type(x::T) where T <: Tuple = Tuple{device_type.(x)...}
-
-@generated function reconstruct(x::Tuple, ptrs::GlobalPointer...)
-    ptrlist = to_tuple(ptrs)
-    tup = Expr(:tuple)
-    ptr_idx = 0
-    for (xi, T) in enumerate(to_tuple(x))
-        hasptr, fields = contains_pointer(T)
-        if hasptr
-            # consume the n pointers that T contains
-            ptr_args = ntuple(i-> :(ptrs[$(i + ptr_idx)]), length(fields))
-            ptr_idx += 1
-            push!(tup.args, :(reconstruct(x[$xi], $(ptr_args...))))
-        else
-            push!(tup.args, :(x[$xi]))
-        end
-    end
-    return tup
-end
-
-
-Base.size(x::OnDeviceArray) = x.size
-
-
-function Base.getindex(x::OnDeviceArray{T, N}, i::Vararg{Integer, N}) where {T, N}
-    ilin = gpu_sub2ind(size(x), Cuint.(i))
-    return x.ptr[ilin]
-end
-function Base.setindex!(x::OnDeviceArray{T, N}, val, i::Vararg{Integer, N}) where {T, N}
-    ilin = gpu_sub2ind(size(x), Cuint.(i))
-    x.ptr[ilin] = T(val)
-    return
-end
-function Base.setindex!(x::OnDeviceArray{T, N}, val, ilin::Integer) where {T, N}
-    x.ptr[ilin] = T(val)
-    return
-end
-
-function Base.getindex(x::OnDeviceArray, ilin::Integer)
-    return x.ptr[ilin]
-end
 # arguments are swapped to not override default constructor
 function (::Type{CLArray{T, N}})(size::NTuple{N, Integer}, ptr::OwnedPtr{T}) where {T, N}
     arr = CLArray{T, N}(ptr, size)

src/device_funcs.jl

@@ -1,8 +1,13 @@
+import GPUArrays: acc_mapreduce
+using Transpiler.cli: get_local_id, get_global_id, barrier, CLK_LOCAL_MEM_FENCE
+using Transpiler.cli: get_local_size, get_global_size, get_group_id
+using GPUArrays: blockdim_x, blockidx_x, threadidx_x, synchronize, synchronize_threads, device
+
 for i = 0:10
     args = ntuple(x-> Symbol("arg_", x), i)
     fargs = ntuple(x-> :(broadcast_index($(args[x]), length, global_index)), i)
     @eval begin
-        function reduce_kernel(f, op, v0, A, tmp_local, length, result, $(args...))
+        function reduce_kernel(state, f, op, v0, A, tmp_local, length, result, $(args...))
             ui1 = Cuint(1)
             global_index = get_global_id(0) + ui1
             local_v0 = v0
@@ -14,21 +19,21 @@ for i = 0:10
             end
 
             # Perform parallel reduction
-            local_index = threadidx_x(A)
+            local_index = threadidx_x(state)
             tmp_local[local_index + ui1] = local_v0
-            synchronize_threads(A)
-            offset = blockdim_x(A) ÷ ui1
+            barrier(CLK_LOCAL_MEM_FENCE)
+            offset = blockdim_x(state) ÷ ui1
             while offset > 0
                 if (local_index < offset)
                     other = tmp_local[local_index + offset + ui1]
                     mine = tmp_local[local_index + ui1]
                     tmp_local[local_index + ui1] = op(mine, other)
                 end
-                synchronize_threads(A)
+                barrier(CLK_LOCAL_MEM_FENCE)
                 offset = offset ÷ Cuint(2)
             end
             if local_index == Cuint(0)
-                result[blockidx_x(A) + ui1] = tmp_local[1]
+                result[blockidx_x(state) + ui1] = tmp_local[1]
             end
             return
         end
@@ -38,16 +43,14 @@ end
 function acc_mapreduce{T, OT, N}(
         f, op, v0::OT, A::CLArray{T, N}, rest::Tuple
     )
-    dev = context(A).device
+    dev = device(A)
     block_size = 16
     group_size = ceil(Int, length(A) / block_size)
     out = similar(A, OT, (group_size,))
     fill!(out, v0)
     lmem = LocalMemory{OT}(block_size)
     args = (f, op, v0, A, lmem, Cuint(length(A)), out, rest...)
-
-    func = CLFunction(A, reduce_kernel, args...)
-    func(args, group_size * block_size, (block_size,))
-    x = reduce(op, Array(out))
-    x
+    gpu_call(reduce_kernel, A, args, (group_size * block_size,), (block_size,))
+    println(Array(out))
+    reduce(op, Array(out))
 end

src/mapreduce.jl

@@ -0,0 +1,86 @@
+import Base: setindex!, getindex, size, IndexStyle, next, done, start, sum, eltype
+using Base: IndexLinear
+"""
+Array type on the device
+"""
+struct DeviceArray{T, N, Ptr} <: AbstractArray{T, N}
+    ptr::Ptr
+    size::NTuple{N, Cuint}
+end
+# shaninagans for uploading CLArrays to OpenCL as a DeviceArray
+# (spoiler alert: they can't contain pointers while uploading, but can on the device)
+"""
+Dummy pointer type for inlining into structs that get uploaded to the GPU
+"""
+struct HostPtr{T}
+    ptr::Int32
+    (::Type{HostPtr{T}})() where T = new{T}(Int32(0))
+end
+eltype(::Type{HostPtr{T}}) where T = T
+const PreDeviceArray{T, N} = DeviceArray{T, N, HostPtr{T}} # Pointer free variant for kernel upload
+const OnDeviceArray{T, N} = DeviceArray{T, N, GlobalPointer{T}} # Variant on the device containing the correct pointer
+
+size(x::OnDeviceArray) = x.size
+IndexStyle(::OnDeviceArray) = IndexLinear()
+start(x::OnDeviceArray) = Cuint(1)
+next(x::OnDeviceArray, state::Cuint) = x[state], state + Cuint(1)
+done(x::OnDeviceArray, state::Cuint) = state > length(x)
+
+function getindex(x::OnDeviceArray{T, N}, i::Vararg{Integer, N}) where {T, N}
+    ilin = gpu_sub2ind(size(x), Cuint.(i))
+    return x.ptr[ilin]
+end
+function setindex!(x::OnDeviceArray{T, N}, val, i::Vararg{Integer, N}) where {T, N}
+    ilin = gpu_sub2ind(size(x), Cuint.(i))
+    x.ptr[ilin] = T(val)
+    return
+end
+function setindex!(x::OnDeviceArray{T, N}, val, ilin::Integer) where {T, N}
+    x.ptr[ilin] = T(val)
+    return
+end
+getindex(x::OnDeviceArray, ilin::Integer) =  x.ptr[ilin]
+
+
+kernel_convert(A::CLArray{T, N}) where {T, N} = PreDeviceArray{T, N}(HostPtr{T}(), A.size)
+predevice_type(::Type{OnDeviceArray{T, N}}) where {T, N} = PreDeviceArray{T, N}
+device_type(::CLArray{T, N}) where {T, N} = OnDeviceArray{T, N}
+reconstruct(x::PreDeviceArray{T, N}, ptr::GlobalPointer{T}) where {T, N} = OnDeviceArray{T, N}(ptr, x.size)
+
+# some converts to inline CLArrays into tuples and refs
+kernel_convert(x::RefValue{T}) where T <: CLArray = RefValue(kernel_convert(x[]))
+predevice_type(::Type{RefValue{T}}) where T <: OnDeviceArray = RefValue{predevice_type(T)}
+device_type(x::RefValue{T}) where T <: CLArray = RefValue{device_type(x[])}
+reconstruct(x::RefValue{T}, ptr::GlobalPointer) where T <: PreDeviceArray = RefValue(reconstruct(x[], ptr))
+
+kernel_convert(x::Tuple) = kernel_convert.(x)
+predevice_type(::Type{T}) where T <: Tuple = Tuple{predevice_type.((T.parameters...))...}
+device_type(x::T) where T <: Tuple = Tuple{device_type.(x)...}
+
+@generated function reconstruct(x::Tuple, ptrs::GlobalPointer...)
+    ptrlist = to_tuple(ptrs)
+    tup = Expr(:tuple)
+    ptr_idx = 0
+    for (xi, T) in enumerate(to_tuple(x))
+        hasptr, fields = contains_pointer(T)
+        if hasptr
+            # consume the n pointers that T contains
+            ptr_args = ntuple(i-> :(ptrs[$(i + ptr_idx)]), length(fields))
+            ptr_idx += 1
+            push!(tup.args, :(reconstruct(x[$xi], $(ptr_args...))))
+        else
+            push!(tup.args, :(x[$xi]))
+        end
+    end
+    return tup
+end
+
+
+
+function sum(A::CLArrays.DeviceArray{T}) where T
+    acc = zero(T)
+    for elem in A
+        acc += elem
+    end
+    acc
+end

src/ondevice.jl

@@ -1,3 +1,4 @@
 using CLArrays
 using GPUArrays.TestSuite, Base.Test
+
 TestSuite.run_tests(CLArray)