diff options
| -rw-r--r-- | libclc/amdgcn/lib/SOURCES | 1 | ||||
| -rw-r--r-- | libclc/amdgcn/lib/workitem/get_global_offset.cl | 11 | ||||
| -rw-r--r-- | libclc/generic/include/clc/clc.h | 1 | ||||
| -rw-r--r-- | libclc/generic/include/clc/workitem/get_global_offset.h | 1 | ||||
| -rw-r--r-- | libclc/generic/lib/workitem/get_global_id.cl | 2 | ||||
| -rw-r--r-- | libclc/ptx-nvidiacl/lib/SOURCES | 1 | ||||
| -rw-r--r-- | libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl | 5 | ||||
| -rw-r--r-- | libclc/r600/lib/SOURCES | 1 | ||||
| -rw-r--r-- | libclc/r600/lib/workitem/get_global_offset.cl | 11 |
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; +} |
