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,7 +451,10 @@ end
 # Internal kernel functions
 ###
 
-function __index_Local_Linear end
+function __index_Local_Linear(ctx)
+    KernelIntrinsics.get_local_id().x
+end
+
 function __index_Group_Linear end
 function __index_Global_Linear end
 

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