fixup! define basic intrinsics

Author: vchuravy

src/KernelAbstractions.jl

@@ -189,6 +189,10 @@ function unsafe_free! end
 
 unsafe_free!(::AbstractArray) = return
 
+include("intrinsics.jl")
+import .KernelIntrinsics
+export KernelIntrinsics
+
 ###
 # Kernel language
 # - @localmem
@@ -447,13 +451,27 @@ end
 # Internal kernel functions
 ###
 
-function __index_Local_Linear end
-function __index_Group_Linear end
-function __index_Global_Linear end
+function __index_Local_Linear(ctx)
+    return KernelIntrinsics.get_local_id().x
+end
 
-function __index_Local_Cartesian end
-function __index_Group_Cartesian end
-function __index_Global_Cartesian end
+function __index_Group_Linear(ctx)
+    return KernelIntrinsics.get_group_id().x
+end
+
+function __index_Global_Linear(ctx)
+    return KernelIntrinsics.get_global_id().x
+end
+
+function __index_Local_Cartesian(ctx)
+    return @inbounds workitems(__iterspace(ctx))[KernelIntrinsics.get_local_id().x]
+end
+function __index_Group_Cartesian(ctx)
+    return @inbounds blocks(__iterspace(ctx))[KernelIntrinsics.get_group_id().x]
+end
+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...))

src/intrinsics.jl

@@ -4,6 +4,9 @@ module KernelIntrinsics
     get_global_size()::@NamedTuple{x::Int32, y::Int32, z::Int32}
 
 Return the number of global work-items specified.
+
+!!! note
+    1-based.
 """
 function get_global_size end
 
@@ -12,7 +15,7 @@ function get_global_size end
 
 Returns the unique global work-item ID.
 """
-function get_global_id end 
+function get_global_id end
 
 """
     get_local_size()::@NamedTuple{x::Int32, y::Int32, z::Int32}
@@ -40,10 +43,10 @@ function get_num_groups end
 
 Returns the unique group ID.
 """
-function get_group_id end 
+function get_group_id end
 
 function localmemory end
 function barrier end
 function print end
 
-end
+end

src/pocl/backend.jl

@@ -139,29 +139,18 @@ end
 
 
 ## Indexing Functions
+const KI = KA.KernelIntrinsics
 
-@device_override @inline function KA.__index_Local_Linear(ctx)
-    return get_local_id(1)
+@device_override @inline function KI.get_local_id()
+    return (; x = get_local_id(1), y = get_local_id(2), z = get_local_id(3))
 end
 
-@device_override @inline function KA.__index_Group_Linear(ctx)
-    return get_group_id(1)
+@device_override @inline function KI.get_group_id()
+    return (; x = get_group_id(1), y = get_group_id(2), z = get_group_id(3))
 end
 
-@device_override @inline function KA.__index_Global_Linear(ctx)
-    return get_global_id(1)
-end
-
-@device_override @inline function KA.__index_Local_Cartesian(ctx)
-    @inbounds KA.workitems(KA.__iterspace(ctx))[get_local_id(1)]
-end
-
-@device_override @inline function KA.__index_Group_Cartesian(ctx)
-    @inbounds KA.blocks(KA.__iterspace(ctx))[get_group_id(1)]
-end
-
-@device_override @inline function KA.__index_Global_Cartesian(ctx)
-    return @inbounds KA.expand(KA.__iterspace(ctx), get_group_id(1), get_local_id(1))
+@device_override @inline function KI.get_global_id()
+    return (; x = get_global_id(1), y = get_global_id(2), z = get_global_id(3))
 end
 
 @device_override @inline function KA.__validindex(ctx)