KernelIntrinsics API

src/KernelAbstractions.jl

@@ -194,6 +194,15 @@ function unsafe_free! end
 
 unsafe_free!(::AbstractArray) = return
 
+"""
+Abstract type for all KernelAbstractions backends.
+"""
+abstract type Backend end
+
+include("intrinsics.jl")
+import .KernelIntrinsics
+export KernelIntrinsics
+
 ###
 # Kernel language
 # - @localmem
@@ -239,11 +248,8 @@ end
 Declare storage that is local to a workgroup.
 """
 macro localmem(T, dims)
-    # Stay in sync with CUDAnative
-    id = gensym("static_shmem")
-
     return quote
-        $SharedMemory($(esc(T)), Val($(esc(dims))), Val($(QuoteNode(id))))
+        $SharedMemory($(esc(T)), Val($(esc(dims))))
     end
 end
 
@@ -460,13 +466,27 @@ end
 # Internal kernel functions
 ###
 
-function __index_Local_Linear end
-function __index_Group_Linear end
-function __index_Global_Linear end
+@inline function __index_Local_Linear(ctx)
+    return KernelIntrinsics.get_local_id().x
+end
+
+@inline function __index_Group_Linear(ctx)
+    return KernelIntrinsics.get_group_id().x
+end
 
-function __index_Local_Cartesian end
-function __index_Group_Cartesian end
-function __index_Global_Cartesian end
+@inline function __index_Global_Linear(ctx)
+    return KernelIntrinsics.get_global_id().x
+end
+
+@inline function __index_Local_Cartesian(ctx)
+    return @inbounds workitems(__iterspace(ctx))[KernelIntrinsics.get_local_id().x]
+end
+@inline function __index_Group_Cartesian(ctx)
+    return @inbounds blocks(__iterspace(ctx))[KernelIntrinsics.get_group_id().x]
+end
+@inline function __index_Global_Cartesian(ctx)
+    return @inbounds expand(__iterspace(ctx), KernelIntrinsics.get_group_id().x, KernelIntrinsics.get_local_id().x)
+end
 
 @inline __index_Local_NTuple(ctx, I...) = Tuple(__index_Local_Cartesian(ctx, I...))
 @inline __index_Group_NTuple(ctx, I...) = Tuple(__index_Group_Cartesian(ctx, I...))
@@ -482,11 +502,6 @@ constify(arg) = adapt(ConstAdaptor(), arg)
 # Backend hierarchy
 ###
 
-"""
-
-Abstract type for all KernelAbstractions backends.
-"""
-abstract type Backend end
 
 """
 Abstract type for all GPU based KernelAbstractions backends.
@@ -796,29 +811,12 @@ include("macros.jl")
 ###
 
 function Scratchpad end
-function SharedMemory end
-
-function __synchronize()
-    error("@synchronize used outside kernel or not captured")
-end
-
-@generated function __print(items...)
-    str = ""
-    args = []
+# SharedMemory(t::Type{T}, dims::Val{Dims}, id::Val{Id}) where {T, Dims, Id} = KernelIntrinsics.localmemory(t, dims, id)
+SharedMemory(t::Type{T}, dims::Val{Dims}) where {T, Dims} = KernelIntrinsics.localmemory(t, dims)
 
-    for i in 1:length(items)
-        item = :(items[$i])
-        T = items[i]
-        if T <: Val
-            item = QuoteNode(T.parameters[1])
-        end
-        push!(args, item)
-    end
+__synchronize() = KernelIntrinsics.barrier()
 
-    return quote
-        print($(args...))
-    end
-end
+__print(args...) = KernelIntrinsics._print(args...)
 
 # Utils
 __size(args::Tuple) = Tuple{args...}