add backend

Author: vchuravy

src/abstract_gpu_interface.jl

@@ -106,8 +106,8 @@ end
 #     CUDAnative.__syncthreads()
 # end
 
-
-
+abstract type GPUBackend end
+backend(::Type{T}) where T = error("Can't choose GPU backend for $T")
 
 """
     gpu_call(kernel::Function, A::GPUArray, args::Tuple, configuration = length(A))
@@ -124,7 +124,7 @@ Optionally, a launch configuration can be supplied in the following way:
     2) Pass a tuple of integer tuples to define blocks and threads per blocks!
 
 """
-function gpu_call(kernel, A::GPUArray, args::Tuple, configuration = length(A))
+function gpu_call(kernel, A::AbstractArray, args::Tuple, configuration = length(A))
     ITuple = NTuple{N, Integer} where N
     # If is a single integer, we assume it to be the global size / total number of threads one wants to launch
     thread_blocks = if isa(configuration, Integer)
@@ -148,8 +148,8 @@ function gpu_call(kernel, A::GPUArray, args::Tuple, configuration = length(A))
                 `linear_index` will be inbetween 1:prod((blocks..., threads...))
         """)
     end
-    _gpu_call(kernel, A, args, thread_blocks)
+    _gpu_call(backend(typeof(A)), kernel, A, args, thread_blocks)
 end
 
 # Internal GPU call function, that needs to be overloaded by the backends.
-_gpu_call(f, A, args, thread_blocks) = error("Not implemented")
+_gpu_call(::Any, f, A, args, thread_blocks) = error("Not implemented")

src/array.jl

@@ -21,6 +21,8 @@ function JLArray{T, N}(size::NTuple{N, Integer}) where {T, N}
     JLArray{T, N}(Array{T, N}(undef, size), size)
 end
 
+struct JLBackend <: GPUBackend end
+backend(::Type{<:JLArray}) = JLBackend()
 
 ## getters
 
@@ -120,7 +122,7 @@ function AbstractDeviceArray(ptr::Array, shape::Vararg{Integer, N}) where N
     reshape(ptr, shape)
 end
 
-function _gpu_call(f, A::JLArray, args::Tuple, blocks_threads::Tuple{T, T}) where T <: NTuple{N, Integer} where N
+function _gpu_call(::JLBackend, f, A, args::Tuple, blocks_threads::Tuple{T, T}) where T <: NTuple{N, Integer} where N
     blocks, threads = blocks_threads
     idx = ntuple(i-> 1, length(blocks))
     blockdim = blocks

src/broadcast.jl

@@ -16,6 +16,9 @@ BroadcastStyle(::Type{T}) where {T<:GPUArray} = ArrayStyle{T}()
 BroadcastStyle(::Type{<:LinearAlgebra.Transpose{<:Any,T}}) where {T<:GPUArray} = BroadcastStyle(T)
 BroadcastStyle(::Type{<:LinearAlgebra.Adjoint{<:Any,T}}) where {T<:GPUArray} = BroadcastStyle(T)
 
+backend(::Type{<:LinearAlgebra.Transpose{<:Any,T}}) where {T<:GPUArray} = backend(T)
+backend(::Type{<:LinearAlgebra.Adjoint{<:Any,T}}) where {T<:GPUArray} = backend(T)
+
 # This Union is a hack. Ideally Base would have a Transpose <: WrappedArray <: AbstractArray
 # and we could define our methods in terms of Union{GPUArray, WrappedArray{<:Any, <:GPUArray}}
 const GPUDestArray = Union{GPUArray,