Merge pull request #1 from JuliaGPU/sd/tests

SimonDanisch · web-flow · commit 23fd71420e33 · 2017-09-26T10:41:45.000+02:00
make tests pass on CI
diff --git a/.travis.yml b/.travis.yml
@@ -3,7 +3,7 @@ sudo: false
 
 os:
   - linux
-  - osx
+  #- osx
 
 julia:
   - 0.6
diff --git a/.travis/amd_sdk.sh b/.travis/amd_sdk.sh
@@ -0,0 +1,38 @@
+#!/bin/bash
+
+# Original script from https://github.com/gregvw/amd_sdk/
+
+# Location from which get nonce and file name from
+URL="http://developer.amd.com/tools-and-sdks/opencl-zone/opencl-tools-sdks/amd-accelerated-parallel-processing-app-sdk/"
+URLDOWN="http://developer.amd.com/amd-license-agreement-appsdk/"
+
+NONCE1_STRING='name="amd_developer_central_downloads_page_nonce"'
+FILE_STRING='name="f"'
+POSTID_STRING='name="post_id"'
+NONCE2_STRING='name="amd_developer_central_nonce"'
+
+#For newest FORM=`wget -qO - $URL | sed -n '/download-2/,/64-bit/p'`
+FORM=`wget -qO - $URL | sed -n '/download-5/,/64-bit/p'`
+
+# Get nonce from form
+NONCE1=`echo $FORM | awk -F ${NONCE1_STRING} '{print $2}'`
+NONCE1=`echo $NONCE1 | awk -F'"' '{print $2}'`
+echo $NONCE1
+
+# get the postid
+POSTID=`echo $FORM | awk -F ${POSTID_STRING} '{print $2}'`
+POSTID=`echo $POSTID | awk -F'"' '{print $2}'`
+echo $POSTID
+
+# get file name
+FILE=`echo $FORM | awk -F ${FILE_STRING} '{print $2}'`
+FILE=`echo $FILE | awk -F'"' '{print $2}'`
+echo $FILE
+
+FORM=`wget -qO - $URLDOWN --post-data "amd_developer_central_downloads_page_nonce=${NONCE1}&f=${FILE}&post_id=${POSTID}"`
+
+NONCE2=`echo $FORM | awk -F ${NONCE2_STRING} '{print $2}'`
+NONCE2=`echo $NONCE2 | awk -F'"' '{print $2}'`
+echo $NONCE2
+
+wget --content-disposition --trust-server-names $URLDOWN --post-data "amd_developer_central_nonce=${NONCE2}&f=${FILE}" -O AMD-SDK.tar.bz2;
diff --git a/REQUIRE b/REQUIRE
@@ -1,14 +1,14 @@
 julia 0.6
 
-GPUArrays
+GPUArrays 0.2.0
 
 StaticArrays
 ColorTypes
 
-Transpiler 0.3
-Sugar 0.3
-Matcha 0.0.2
+Transpiler 0.4.1
+Sugar 0.4.1
+Matcha 0.1.1
 
-OpenCL 0.6.0 #proper packed conversion
+OpenCL 0.6.1 #proper packed pointer free
 CLBLAS 1.1.0
 CLFFT 0.4.0 # 0.5.0
diff --git a/src/CLArrays.jl b/src/CLArrays.jl
@@ -14,9 +14,7 @@ include("array.jl")
 include("ondevice.jl")
 include("device.jl")
 include("context.jl")
-include("intrinsics.jl")
 include("compilation.jl")
-include("mapreduce.jl")
 include("3rdparty.jl")
 
 export CLArray, gpu_call
diff --git a/src/array.jl b/src/array.jl
@@ -42,11 +42,14 @@ similar(::Type{<: CLArray}, ::Type{T}, size::Base.Dims{N}) where {T, N} = CLArra
 function unsafe_free!(a::CLArray)
     ptr = pointer(a)
     ctxid = context(ptr).id
-    if cl.is_ctx_id_alive(ctxid) && ctxid != C_NULL
-        Mem.free(ptr)
+    err = Mem.free(ptr)
+    if err == cl.CL_SUCCESS
         Mem.current_allocated_mem[] -= sizeof(eltype(a)) * length(a)
+    elseif err == cl.CL_INVALID_CONTEXT
+        # ignore (log it!)
+    else
+        cl.@check err
     end
-    #TODO logging that we don't free since context is not alive
 end
 
 function unsafe_reinterpret(::Type{T}, A::CLArray{ET}, size::NTuple{N, Integer}) where {T, ET, N}
diff --git a/src/compilation.jl b/src/compilation.jl
@@ -1,24 +1,23 @@
 using OpenCL: cl
 using Transpiler: CLMethod, EmptyStruct
 using Sugar
-import GPUArrays: gpu_call, linear_index
+import GPUArrays: _gpu_call, linear_index
 using Transpiler: CLMethod
 using Sugar: method_nargs, getslots!, isintrinsic, getcodeinfo!, sugared
 using Sugar: returntype, type_ast, rewrite_ast, newslot!, to_tuple
 using Sugar: isfunction
 
 using Base: tail
 
-function gpu_call(f, A::CLArray, args::Tuple, blocks = nothing, thread = C_NULL)
+
+
+function _gpu_call(f, A::CLArray, args::Tuple, blocks_threads::Tuple{T, T}) where T <: NTuple{N, Integer} where N
     ctx = context(A)
     _args = (KernelState(), args...) # CLArrays "state"
     clfunc = CLFunction(f, _args, ctx)
-    if blocks == nothing
-        blocks, thread = thread_blocks_heuristic(length(A))
-    elseif isa(blocks, Integer)
-        blocks = (blocks,)
-    end
-    clfunc(_args, blocks, thread)
+    blocks, threads = blocks_threads
+    global_size = blocks .* threads
+    clfunc(_args, global_size, threads)
 end
 
 
diff --git a/src/context.jl b/src/context.jl
@@ -11,7 +11,7 @@ end
 function show(io::IO, ctx::CLContext)
     println(io, "OpenCL context with:")
     println(io, "CL version: ", cl.info(ctx.device, :version))
-    device_summary(io, ctx.device)
+    GPUArrays.device_summary(io, ctx.device)
 end
 
 function CLContext(device::cl.Device)
@@ -47,7 +47,7 @@ function device(x::cl.Context)
 end
 
 
-
+global init
 let contexts = Dict{cl.Device, CLContext}(), active_device = cl.Device[], clcontext2context = Dict{cl.Context, CLContext}()
     function getcontext!(ctx::cl.Context)
         if haskey(clcontext2context, ctx)
diff --git a/src/device.jl b/src/device.jl
@@ -4,10 +4,7 @@ import GPUArrays: supports_double, local_memory
 function devices()
     filter(cl.devices()) do dev
         # These drivers have some problems -  TODO figure out if we are the problem
-        !(
-            contains(cl.info(dev, :version), "AMD-APP (2348.3)") || #AMD for intel. Why though?
-            contains(cl.info(dev, :version), "(Build 10)") # intels experimental 2.1 driver
-        )
+        !contains(cl.info(dev, :version), "(Build 10)") # intels experimental 2.1 driver
     end
 end
 
diff --git a/src/intrinsics.jl b/src/intrinsics.jl
diff --git a/src/mapreduce.jl b/src/mapreduce.jl
diff --git a/src/memory.jl b/src/memory.jl
@@ -24,17 +24,16 @@ function convert{T}(::Type{OwnedPtr{T}}, p::OwnedPtr)
 end
 
 function retain(p::OwnedPtr)
-    cl.@check cl.api.clRetainMemObject(p)
-    return
+    err = cl.api.clRetainMemObject(p)
+    return err
 end
 
 
 
 
 const current_allocated_mem = Ref(0)
 function free(p::OwnedPtr)
-    cl.@check_release cl.api.clReleaseMemObject(p)
-    return
+    cl.api.clReleaseMemObject(p)
 end
 
 function pressure_gc!(device, bytes)
diff --git a/src/ondevice.jl b/src/ondevice.jl
@@ -1,9 +1,19 @@
-import Base: setindex!, getindex, size, IndexStyle, next, done, start, sum, eltype
+import Base: setindex!, getindex, size, IndexStyle, sum, eltype
 using Base: IndexLinear
+
+using Transpiler.cli: LocalPointer
+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, get_num_groups
+import GPUArrays: synchronize, synchronize_threads, device, global_size, linear_index
+
+import GPUArrays: LocalMemory
+using GPUArrays: AbstractDeviceArray
+
+
 """
 Array type on the device
 """
-struct DeviceArray{T, N, Ptr} <: AbstractArray{T, N}
+struct DeviceArray{T, N, Ptr} <: AbstractDeviceArray{T, N}
     ptr::Ptr
     size::NTuple{N, Cuint}
 end
@@ -18,13 +28,12 @@ struct HostPtr{T}
 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
+const GlobalArray{T, N} = DeviceArray{T, N, GlobalPointer{T}}
+const LocalArray{T, N} = DeviceArray{T, N, LocalPointer{T}}
+
+const OnDeviceArray{T, N} = Union{GlobalArray{T, N}, LocalArray{T, N}} # 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)
+size(x::DeviceArray) = x.size
 
 getindex(x::OnDeviceArray, ilin::Integer) =  x.ptr[ilin]
 function getindex(x::OnDeviceArray{T, N}, i::Vararg{Integer, N}) where {T, N}
@@ -45,13 +54,13 @@ end
 
 
 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)
+predevice_type(::Type{GlobalArray{T, N}}) where {T, N} = PreDeviceArray{T, N}
+device_type(::CLArray{T, N}) where {T, N} = GlobalArray{T, N}
+reconstruct(x::PreDeviceArray{T, N}, ptr::GlobalPointer{T}) where {T, N} = GlobalArray{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)}
+predevice_type(::Type{RefValue{T}}) where T <: GlobalArray = 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))
 
@@ -78,11 +87,44 @@ device_type(x::T) where T <: Tuple = Tuple{device_type.(x)...}
 end
 
 
+#synchronize
+function synchronize(x::CLArray)
+    cl.finish(global_queue(x)) # TODO figure out the diverse ways of synchronization
+end
+
+
+immutable KernelState
+    empty::Int32
+    KernelState() = new(Int32(0))
+end
+
+for (i, sym) in enumerate((:x, :y, :z))
+    for (f, fcl, isidx) in (
+            (:blockidx, get_group_id, true),
+            (:blockdim, get_local_size, false),
+            (:threadidx, get_local_id, true),
+            (:griddim, get_num_groups, false)
+        )
 
-function sum(A::CLArrays.DeviceArray{T}) where T
-    acc = zero(T)
-    for elem in A
-        acc += elem
+        fname = Symbol(string(f, '_', sym))
+        if isidx
+            @eval GPUArrays.$fname(::KernelState)::Cuint = $fcl($(i-1)) + Cuint(1)
+        else
+            @eval GPUArrays.$fname(::KernelState)::Cuint = $fcl($(i-1))
+        end
     end
-    acc
+end
+
+global_size(state::KernelState) = get_global_size(0)
+linear_index(state::KernelState) = get_global_id(0) + Cuint(1)
+
+synchronize_threads(::KernelState) = cli.barrier(CLK_LOCAL_MEM_FENCE)
+
+LocalMemory(state::KernelState, ::Type{T}, ::Val{N}, ::Val{C}) where {T, N, C} = Transpiler.cli.LocalPointer{T}()
+
+function (::Type{AbstractDeviceArray})(ptr::PtrT, shape::Vararg{Integer, N}) where PtrT <: Transpiler.cli.LocalPointer{T} where {T, N}
+    DeviceArray{T, N, PtrT}(ptr, shape)
+end
+function (::Type{AbstractDeviceArray})(ptr::PtrT, shape::NTuple{N, Integer}) where PtrT <: Transpiler.cli.LocalPointer{T} where {T, N}
+    DeviceArray{T, N, PtrT}(ptr, shape)
 end
diff --git a/test/runtests.jl b/test/runtests.jl
@@ -1,15 +1,35 @@
-if (
-        get(ENV, "TRAVIS", "") == "true" ||
-        get(ENV, "APPVEYOR", "") == "true" ||
-        get(ENV, "CI", "") == "true"
-    )
-
-    Pkg.checkout("Transpiler", "sd/struct_buffer")
-    Pkg.checkout("Sugar", "sd/struct_buffer")
-    Pkg.checkout("GPUArrays", "sd/abstractgpu")
-end
-
 using CLArrays
 using GPUArrays.TestSuite, Base.Test
+for dev in CLArrays.devices()
+    CLArrays.init(dev)
+    @testset "Device: $dev" begin
+        TestSuite.run_tests(CLArray)
+
+        @testset "muladd & abs" begin
+            a = rand(Float32, 32) - 0.5f0
+            A = CLArray(a)
+            x = abs.(A)
+            @test Array(x) == abs.(a)
+            y = muladd.(A, 2f0, x)
+            @test Array(y) == muladd(a, 2f0, abs.(a))
+            ###########
+            # issue #20
+
+            against_base(a-> abs.(a), CLArray{Float32}, (10, 10))
+        end
+    end
+end
+
 
-TestSuite.run_tests(CLArray)
+# The above is equal to:
+# Typ = CuArray
+# GPUArrays.allowslow(false)
+# TestSuite.run_gpuinterface(Typ)
+# TestSuite.run_base(Typ)
+# TestSuite.run_blas(Typ)
+# TestSuite.run_broadcasting(Typ)
+# TestSuite.run_construction(Typ)
+# TestSuite.run_fft(Typ)
+# TestSuite.run_linalg(Typ)
+# TestSuite.run_mapreduce(Typ)
+# TestSuite.run_indexing(Typ)

-Original file line number
+Diff line change
 os:
   - linux
 -  - osx
 +  #- osx
 julia:
   - 0.6