summaryrefslogtreecommitdiffstats
path: root/libclc
diff options
context:
space:
mode:
authorJan Vesely <jan.vesely@rutgers.edu>2016-07-22 17:24:24 +0000
committerJan Vesely <jan.vesely@rutgers.edu>2016-07-22 17:24:24 +0000
commita82e080b57a47e9ef8f89c56bdd6b9593db596d2 (patch)
tree392d06b76bd59817ac4897128a1961dc84c39208 /libclc
parent74f02db922b4609095da4218fd3016c2c51d056b (diff)
downloadbcm5719-llvm-a82e080b57a47e9ef8f89c56bdd6b9593db596d2.tar.gz
bcm5719-llvm-a82e080b57a47e9ef8f89c56bdd6b9593db596d2.zip
AMDGPU: Implement get_global_offset builtin
Also fix get_global_id to consider offset No idea how to add this for ptx, so they are stuck with the old get_global_id implementation. v2: split to a separate patch v3: Switch R600 to use implictarg.ptr Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> llvm-svn: 276443
Diffstat (limited to 'libclc')
-rw-r--r--libclc/amdgcn/lib/SOURCES1
-rw-r--r--libclc/amdgcn/lib/workitem/get_global_offset.cl11
-rw-r--r--libclc/generic/include/clc/clc.h1
-rw-r--r--libclc/generic/include/clc/workitem/get_global_offset.h1
-rw-r--r--libclc/generic/lib/workitem/get_global_id.cl2
-rw-r--r--libclc/ptx-nvidiacl/lib/SOURCES1
-rw-r--r--libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl5
-rw-r--r--libclc/r600/lib/SOURCES1
-rw-r--r--libclc/r600/lib/workitem/get_global_offset.cl11
9 files changed, 33 insertions, 1 deletions
diff --git a/libclc/amdgcn/lib/SOURCES b/libclc/amdgcn/lib/SOURCES
index 49d9b531408..ca88b934628 100644
--- a/libclc/amdgcn/lib/SOURCES
+++ b/libclc/amdgcn/lib/SOURCES
@@ -1,5 +1,6 @@
math/ldexp.cl
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/libclc/amdgcn/lib/workitem/get_global_offset.cl b/libclc/amdgcn/lib/workitem/get_global_offset.cl
new file mode 100644
index 00000000000..32aaa4cfa15
--- /dev/null
+++ b/libclc/amdgcn/lib/workitem/get_global_offset.cl
@@ -0,0 +1,11 @@
+#include <clc/clc.h>
+
+_CLC_DEF uint get_global_offset(uint dim)
+{
+ __attribute__((address_space(2))) uint * ptr =
+ (__attribute__((address_space(2))) uint *)
+ __builtin_amdgcn_implicitarg_ptr();
+ if (dim < 3)
+ return ptr[dim + 1];
+ return 0;
+}
diff --git a/libclc/generic/include/clc/clc.h b/libclc/generic/include/clc/clc.h
index 61f9b3571e3..d7931f1ce2e 100644
--- a/libclc/generic/include/clc/clc.h
+++ b/libclc/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/libclc/generic/include/clc/workitem/get_global_offset.h b/libclc/generic/include/clc/workitem/get_global_offset.h
new file mode 100644
index 00000000000..7f4f6032abe
--- /dev/null
+++ b/libclc/generic/include/clc/workitem/get_global_offset.h
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_offset(uint dim);
diff --git a/libclc/generic/lib/workitem/get_global_id.cl b/libclc/generic/lib/workitem/get_global_id.cl
index fdd83d2953d..b6c2ea1d2cc 100644
--- a/libclc/generic/lib/workitem/get_global_id.cl
+++ b/libclc/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/libclc/ptx-nvidiacl/lib/SOURCES b/libclc/ptx-nvidiacl/lib/SOURCES
index 7cdbd850769..ce26bcb8557 100644
--- a/libclc/ptx-nvidiacl/lib/SOURCES
+++ b/libclc/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/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl b/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl
new file mode 100644
index 00000000000..19bc195312c
--- /dev/null
+++ b/libclc/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/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES
index 4178d70b84b..33038f2800f 100644
--- a/libclc/r600/lib/SOURCES
+++ b/libclc/r600/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/libclc/r600/lib/workitem/get_global_offset.cl b/libclc/r600/lib/workitem/get_global_offset.cl
new file mode 100644
index 00000000000..b38ae337757
--- /dev/null
+++ b/libclc/r600/lib/workitem/get_global_offset.cl
@@ -0,0 +1,11 @@
+#include <clc/clc.h>
+
+_CLC_DEF uint get_global_offset(uint dim)
+{
+ __attribute__((address_space(7))) uint * ptr =
+ (__attribute__((address_space(7))) uint *)
+ __builtin_r600_implicitarg_ptr();
+ if (dim < 3)
+ return ptr[dim + 1];
+ return 0;
+}
OpenPOWER on IntegriCloud