Reimplement ROCArray using GPUArrays

Manifest.toml

@@ -49,7 +49,9 @@ version = "5.1.0"
 
 [[GPUCompiler]]
 deps = ["DataStructures", "InteractiveUtils", "LLVM", "Libdl", "TimerOutputs", "UUIDs"]
-git-tree-sha1 = "10b1a3aa52de30e9219f3ed147cb09e72cf6d2e8"
+git-tree-sha1 = "e0137fdb7c1d0fe217c39a5a3586a4e10a94ddda"
+repo-rev = "master"
+repo-url = "https://github.com/JuliaGPU/GPUCompiler.jl.git"
 uuid = "61eb1bfa-7361-4325-ad38-22787b887f55"
 version = "0.7.0"
 

Project.toml

@@ -1,7 +1,7 @@
 name = "AMDGPU"
 uuid = "21141c5a-9bdb-4563-92ae-f87d6854732e"
 authors = ["Julian P Samaroo <[email protected]>"]
-version = "0.1.2"
+version = "0.2.0"
 
 [deps]
 AbstractFFTs = "621f4979-c628-5d54-868e-fcf4e3e8185c"

deps/deps.jl

@@ -1,14 +1,13 @@
 # HSA runtime
 ## copied from CUDAdrv/src/CUDAdrv.jl
 const hsa_ext = joinpath(@__DIR__, "hsa", "ext.jl")
-if isfile(hsa_ext)
-    include(hsa_ext)
-end
-if !isdefined(@__MODULE__, :hsa_configured)
+if !isfile(hsa_ext)
+    @warn "Didn't find $hsa_ext, please build AMDGPU.jl"
     const hsa_configured = false
+else
+    include(hsa_ext)
 end
 if !hsa_configured
-    const hsa_configured = false
     const libhsaruntime_version = v"0.0"
     const libhsaruntime_vendor = "none"
     const libhsaruntime_path = nothing
@@ -19,11 +18,11 @@ const device_libs_path = joinpath(@__DIR__, "device-libs", "usr", "lib")
 
 # ROCm External Libraries
 const libs_ext = joinpath(@__DIR__, "rocm-external", "ext.jl")
-if isfile(libs_ext)
-    include(libs_ext)
-end
-if !isdefined(@__MODULE__, :ext_libs_configured)
+if !isfile(libs_ext)
+    @warn "Didn't find $libs_ext, please build AMDGPU.jl"
     const ext_libs_configured = false
+else
+    include(libs_ext)
 end
 if !ext_libs_configured
     # default (non-functional) values for critical variables,

src/AMDGPU.jl

@@ -22,10 +22,6 @@ export ROCArray, ROCVector, ROCMatrix, ROCVecOrMat
 export roc, roczeros, rocones, rocfill
 export HSAArray
 
-### Binary Dependencies ###
-
-include(joinpath(dirname(@__DIR__), "deps", "deps.jl"))
-
 ### HSA Runtime ###
 
 include(joinpath(@__DIR__, "hsa", "HSA.jl"))
@@ -84,40 +80,41 @@ include("array.jl")
 roc(xs) = adapt(ROCArray{Float32}, xs)
 allowscalar(x::Bool) = nothing
 
-### External Libraries ###
-
-# TODO: add check
-include("hip/HIP.jl")
-librocblas !== nothing     && include("blas/rocBLAS.jl")
-librocfft !== nothing      && include("fft/rocFFT.jl")
-#librocsparse !== nothing  && include("sparse/rocSPARSE.jl")
-#librocalution !== nothing && include("solver/rocALUTION.jl")
-#librocrand !== nothing    && include("rand/rocRAND.jl")
-#libmiopen !== nothing     && include("dnn/MIOpen.jl")
-
 ### Initialization and Shutdown ###
 
-atexit() do
-    configured && HSA.shut_down()
-end
 function __init__()
     deps_failed() = @warn """
         AMDGPU dependencies have not been built, some functionality may be missing.
         Please run Pkg.build("AMDGPU") and reload AMDGPU.
         """
 
+    # Load binary dependencies
+    include(joinpath(dirname(@__DIR__), "deps", "deps.jl"))
+
     # We want to always be able to load the package
     if !configured
         deps_failed()
         return
     end
 
+    # TODO: add check
+    include(joinpath(@__DIR__, "hip", "HIP.jl"))
+    librocblas !== nothing     && include(joinpath(@__DIR__, "blas", "rocBLAS.jl"))
+    librocfft !== nothing      && include(joinpath(@__DIR__, "fft", "rocFFT.jl"))
+    #librocsparse !== nothing  && include("sparse/rocSPARSE.jl")
+    #librocalution !== nothing && include("solver/rocALUTION.jl")
+    #librocrand !== nothing    && include("rand/rocRAND.jl")
+    #libmiopen !== nothing     && include("dnn/MIOpen.jl")
+
     # Make sure we load the library found by the last `] build`
     push!(Libdl.DL_LOAD_PATH, dirname(libhsaruntime_path))
     # TODO: Do the same (if possible) for the debug library
 
     # Initialize the HSA runtime
     HSA.init() |> check
+    atexit() do
+        configured && HSA.shut_down()
+    end
 
     # Populate the default agent
     agents = get_agents(:gpu)