From 87af609a20ae8ae2c2d355333757c3ebfff0509e Mon Sep 17 00:00:00 2001
From: Valentin Churavy <v.churavy@gmail.com>
Date: Tue, 4 Feb 2025 15:50:59 +0100
Subject: [PATCH] define basic intrinsics

---
 src/KernelAbstractions.jl | 30 +++++++++++++++++-----
 src/intrinsics.jl         | 52 +++++++++++++++++++++++++++++++++++++++
 src/pocl/backend.jl       | 25 ++++++-------------
 3 files changed, 83 insertions(+), 24 deletions(-)
 create mode 100644 src/intrinsics.jl

diff --git a/src/KernelAbstractions.jl b/src/KernelAbstractions.jl
index 1211d2fe..896da9dc 100644
--- a/src/KernelAbstractions.jl
+++ b/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...))
diff --git a/src/intrinsics.jl b/src/intrinsics.jl
new file mode 100644
index 00000000..33ed56fe
--- /dev/null
+++ b/src/intrinsics.jl
@@ -0,0 +1,52 @@
+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
+
+"""
+    get_global_id()::@NamedTuple{x::Int32, y::Int32, z::Int32}
+
+Returns the unique global work-item ID.
+"""
+function get_global_id end
+
+"""
+    get_local_size()::@NamedTuple{x::Int32, y::Int32, z::Int32}
+
+Return the number of local work-items specified.
+"""
+function get_local_size end
+
+"""
+    get_local_id()::@NamedTuple{x::Int32, y::Int32, z::Int32}
+
+Returns the unique local work-item ID.
+"""
+function get_local_id end
+
+"""
+    get_num_groups()::@NamedTuple{x::Int32, y::Int32, z::Int32}
+
+Returns the number of groups.
+"""
+function get_num_groups end
+
+"""
+    get_group_id()::@NamedTuple{x::Int32, y::Int32, z::Int32}
+
+Returns the unique group ID.
+"""
+function get_group_id end
+
+function localmemory end
+function barrier end
+function print end
+
+end
diff --git a/src/pocl/backend.jl b/src/pocl/backend.jl
index 8f8acfcd..16fbeaae 100644
--- a/src/pocl/backend.jl
+++ b/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)