Author: jvesely
Date: Fri Jul 22 12:24:24 2016
New Revision: 276443
URL: http://llvm.org/viewvc/llvm-project?rev=276443&view=rev
Log:
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.ves...@rutgers.edu>
Added:
libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl
libclc/trunk/generic/include/clc/workitem/get_global_offset.h
libclc/trunk/ptx-nvidiacl/lib/workitem/get_global_id.cl
libclc/trunk/r600/lib/workitem/get_global_offset.cl
Modified:
libclc/trunk/amdgcn/lib/SOURCES
libclc/trunk/generic/include/clc/clc.h
libclc/trunk/generic/lib/workitem/get_global_id.cl
libclc/trunk/ptx-nvidiacl/lib/SOURCES
libclc/trunk/r600/lib/SOURCES
Modified: libclc/trunk/amdgcn/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES?rev=276443&r1=276442&r2=276443&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES (original)
+++ libclc/trunk/amdgcn/lib/SOURCES Fri Jul 22 12:24:24 2016
@@ -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
Added: libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl?rev=276443&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl (added)
+++ libclc/trunk/amdgcn/lib/workitem/get_global_offset.cl Fri Jul 22 12:24:24
2016
@@ -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;
+}
Modified: libclc/trunk/generic/include/clc/clc.h
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/clc.h?rev=276443&r1=276442&r2=276443&view=diff
==============================================================================
--- libclc/trunk/generic/include/clc/clc.h (original)
+++ libclc/trunk/generic/include/clc/clc.h Fri Jul 22 12:24:24 2016
@@ -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>
Added: libclc/trunk/generic/include/clc/workitem/get_global_offset.h
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/workitem/get_global_offset.h?rev=276443&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/workitem/get_global_offset.h (added)
+++ libclc/trunk/generic/include/clc/workitem/get_global_offset.h Fri Jul 22
12:24:24 2016
@@ -0,0 +1 @@
+_CLC_DECL size_t get_global_offset(uint dim);
Modified: libclc/trunk/generic/lib/workitem/get_global_id.cl
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/lib/workitem/get_global_id.cl?rev=276443&r1=276442&r2=276443&view=diff
==============================================================================
--- libclc/trunk/generic/lib/workitem/get_global_id.cl (original)
+++ libclc/trunk/generic/lib/workitem/get_global_id.cl Fri Jul 22 12:24:24 2016
@@ -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);
}
Modified: libclc/trunk/ptx-nvidiacl/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/SOURCES?rev=276443&r1=276442&r2=276443&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/SOURCES (original)
+++ libclc/trunk/ptx-nvidiacl/lib/SOURCES Fri Jul 22 12:24:24 2016
@@ -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
Added: libclc/trunk/ptx-nvidiacl/lib/workitem/get_global_id.cl
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_global_id.cl?rev=276443&view=auto
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/workitem/get_global_id.cl (added)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_global_id.cl Fri Jul 22 12:24:24
2016
@@ -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);
+}
Modified: libclc/trunk/r600/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/SOURCES?rev=276443&r1=276442&r2=276443&view=diff
==============================================================================
--- libclc/trunk/r600/lib/SOURCES (original)
+++ libclc/trunk/r600/lib/SOURCES Fri Jul 22 12:24:24 2016
@@ -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
Added: libclc/trunk/r600/lib/workitem/get_global_offset.cl
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_global_offset.cl?rev=276443&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_global_offset.cl (added)
+++ libclc/trunk/r600/lib/workitem/get_global_offset.cl Fri Jul 22 12:24:24 2016
@@ -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;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
