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 | e3cd5bd8e6ebffadd222726322f595f28c47f275 (patch) | |
tree | 7098d2b8a5198f55c72807c031b9d1dbfd781b91 | |
parent | ff45572d5e10a96309d58593b967a71228322f8b (diff) |
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>
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@276443 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r-- | amdgcn/lib/SOURCES | 1 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_global_offset.cl | 11 | ||||
-rw-r--r-- | generic/include/clc/clc.h | 1 | ||||
-rw-r--r-- | generic/include/clc/workitem/get_global_offset.h | 1 | ||||
-rw-r--r-- | generic/lib/workitem/get_global_id.cl | 2 | ||||
-rw-r--r-- | ptx-nvidiacl/lib/SOURCES | 1 | ||||
-rw-r--r-- | ptx-nvidiacl/lib/workitem/get_global_id.cl | 5 | ||||
-rw-r--r-- | r600/lib/SOURCES | 1 | ||||
-rw-r--r-- | r600/lib/workitem/get_global_offset.cl | 11 |
9 files changed, 33 insertions, 1 deletions
diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES index 49d9b53..ca88b93 100644 --- a/amdgcn/lib/SOURCES +++ b/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/amdgcn/lib/workitem/get_global_offset.cl b/amdgcn/lib/workitem/get_global_offset.cl new file mode 100644 index 0000000..32aaa4c --- /dev/null +++ b/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/generic/include/clc/clc.h b/generic/include/clc/clc.h index 61f9b35..d7931f1 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..7f4f603 --- /dev/null +++ b/generic/include/clc/workitem/get_global_offset.h @@ -0,0 +1 @@ +_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 4178d70..33038f2 100644 --- a/r600/lib/SOURCES +++ b/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/r600/lib/workitem/get_global_offset.cl b/r600/lib/workitem/get_global_offset.cl new file mode 100644 index 0000000..b38ae33 --- /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) +{ + __attribute__((address_space(7))) uint * ptr = + (__attribute__((address_space(7))) uint *) + __builtin_r600_implicitarg_ptr(); + if (dim < 3) + return ptr[dim + 1]; + return 0; +} |