[Libclc-dev] [PATCH RFC 1/1] AMDGPU: Implement get_global_offset builtin

[Jan Vesely via Libclc-dev]

Also fix get_global_id to consider offset
Fixes global-offset piglit and GEGL video-degradation tests on both r600 (Turks) and GCN(Kaveri)

Depends on: http://reviews.llvm.org/D20299

Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
---
TODO:
for some reason GCN needs the pointer to be global to work.
No idea how to add this for ptx, so they are stuck with the old get_global_id
implementation for now
Also no idea how this is used for HSA, so it might need different offsets.

 amdgcn/lib/SOURCES                               |  1 +
 amdgcn/lib/workitem/get_global_offset.cl         |  9 +++++++++
 generic/include/clc/clc.h                        |  1 +
 generic/include/clc/workitem/get_global_offset.h |  2 ++
 generic/lib/workitem/get_global_id.cl            |  2 +-
 ptx-nvidiacl/lib/SOURCES                         |  1 +
 ptx-nvidiacl/lib/workitem/get_global_id.cl       |  5 +++++
 r600/lib/SOURCES                                 |  1 +
 r600/lib/workitem/get_global_offset.cl           | 11 +++++++++++
 9 files changed, 32 insertions(+), 1 deletion(-)
 create mode 100644 amdgcn/lib/workitem/get_global_offset.cl
 create mode 100644 generic/include/clc/workitem/get_global_offset.h
 create mode 100644 ptx-nvidiacl/lib/workitem/get_global_id.cl
 create mode 100644 r600/lib/workitem/get_global_offset.cl

diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES
index b7e1d98..3a90aa5 100644
--- a/amdgcn/lib/SOURCES
+++ b/amdgcn/lib/SOURCES
@@ -1,4 +1,5 @@
 synchronization/barrier_impl.ll
+workitem/get_global_offset.cl
 workitem/get_group_id.cl
 workitem/get_local_id.cl
 workitem/get_work_dim.cl
diff --git a/amdgcn/lib/workitem/get_global_offset.cl b/amdgcn/lib/workitem/get_global_offset.cl
new file mode 100644
index 0000000..4691409
--- /dev/null
+++ b/amdgcn/lib/workitem/get_global_offset.cl
@@ -0,0 +1,9 @@
+#include <clc/clc.h>
+
+_CLC_DEF uint get_global_offset(uint dim)
+{
+	__global uint * ptr = __builtin_amdgcn_kernarg_segment_ptr();
+	if (dim < 3)
+		return ptr[dim +1];
+	return 1;
+}
diff --git a/generic/include/clc/clc.h b/generic/include/clc/clc.h
index 6694f03..f77e495 100644
--- a/generic/include/clc/clc.h
+++ b/generic/include/clc/clc.h
@@ -30,6 +30,7 @@
 #include <clc/workitem/get_local_id.h>
 #include <clc/workitem/get_num_groups.h>
 #include <clc/workitem/get_group_id.h>
+#include <clc/workitem/get_global_offset.h>
 
 /* 6.11.2 Math Functions */
 #include <clc/math/acos.h>
diff --git a/generic/include/clc/workitem/get_global_offset.h b/generic/include/clc/workitem/get_global_offset.h
new file mode 100644
index 0000000..630156e
--- /dev/null
+++ b/generic/include/clc/workitem/get_global_offset.h
@@ -0,0 +1,2 @@
+_CLC_DECL size_t get_global_offset(uint dim);
+
diff --git a/generic/lib/workitem/get_global_id.cl b/generic/lib/workitem/get_global_id.cl
index fdd83d2..b6c2ea1 100644
--- a/generic/lib/workitem/get_global_id.cl
+++ b/generic/lib/workitem/get_global_id.cl
@@ -1,5 +1,5 @@
 #include <clc/clc.h>
 
 _CLC_DEF size_t get_global_id(uint dim) {
-  return get_group_id(dim)*get_local_size(dim) + get_local_id(dim);
+  return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + get_global_offset(dim);
 }
diff --git a/ptx-nvidiacl/lib/SOURCES b/ptx-nvidiacl/lib/SOURCES
index 7cdbd85..ce26bcb 100644
--- a/ptx-nvidiacl/lib/SOURCES
+++ b/ptx-nvidiacl/lib/SOURCES
@@ -1,4 +1,5 @@
 synchronization/barrier.cl
+workitem/get_global_id.cl
 workitem/get_group_id.cl
 workitem/get_local_id.cl
 workitem/get_local_size.cl
diff --git a/ptx-nvidiacl/lib/workitem/get_global_id.cl b/ptx-nvidiacl/lib/workitem/get_global_id.cl
new file mode 100644
index 0000000..19bc195
--- /dev/null
+++ b/ptx-nvidiacl/lib/workitem/get_global_id.cl
@@ -0,0 +1,5 @@
+#include <clc/clc.h>
+
+_CLC_DEF size_t get_global_id(uint dim) {
+  return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
+}
diff --git a/r600/lib/SOURCES b/r600/lib/SOURCES
index 1859b98..5b68ce1 100644
--- a/r600/lib/SOURCES
+++ b/r600/lib/SOURCES
@@ -1,4 +1,5 @@
 synchronization/barrier_impl.ll
+workitem/get_global_offset.cl
 workitem/get_global_size.cl
 workitem/get_group_id.cl
 workitem/get_local_id.cl
diff --git a/r600/lib/workitem/get_global_offset.cl b/r600/lib/workitem/get_global_offset.cl
new file mode 100644
index 0000000..74adcc3
--- /dev/null
+++ b/r600/lib/workitem/get_global_offset.cl
@@ -0,0 +1,11 @@
+#include <clc/clc.h>
+
+_CLC_DEF uint get_global_offset(uint dim)
+{
+	switch(dim) {
+	case 0: return __builtin_r600_read_global_offset_x();
+	case 1: return __builtin_r600_read_global_offset_y();
+	case 2: return __builtin_r600_read_global_offset_z();
+	default: return 1;
+	}
+}
-- 
2.5.5