diff options
| -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; +} |