diff options
| author | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 17:24:24 +0000 |
|---|---|---|
| committer | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 17:24:24 +0000 |
| commit | a82e080b57a47e9ef8f89c56bdd6b9593db596d2 (patch) | |
| tree | 392d06b76bd59817ac4897128a1961dc84c39208 /libclc/amdgcn/lib | |
| parent | 74f02db922b4609095da4218fd3016c2c51d056b (diff) | |
| download | bcm5719-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/amdgcn/lib')
| -rw-r--r-- | libclc/amdgcn/lib/SOURCES | 1 | ||||
| -rw-r--r-- | libclc/amdgcn/lib/workitem/get_global_offset.cl | 11 |
2 files changed, 12 insertions, 0 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; +} |
