[Libclc-dev] [PATCH v2 2/2] AMDGPU: Implement get_global_offset builtin
Tom Stellard via Libclc-dev
libclc-dev at lists.llvm.org
Wed Jun 1 19:01:24 PDT 2016
LGTM.
On Wed, Jun 01, 2016 at 12:12:35PM -0400, Jan Vesely via Libclc-dev wrote:
> 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.
>
> Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> ---
>
> Depends on http://reviews.llvm.org/D20299
>
> amdgcn/lib/SOURCES | 1 +
> amdgcn/lib/workitem/get_global_offset.cl | 9 +++++++++
> amdgcn/lib/workitem/get_work_dim.cl | 2 +-
> generic/include/clc/clc.h | 1 +
> generic/include/clc/workitem/get_global_offset.h | 2 ++
> generic/lib/workitem/get_global_id.cl | 2 +-
> ptx-nvidiacl/lib/SOURCES | 1 +
> ptx-nvidiacl/lib/workitem/get_global_id.cl | 5 +++++
> r600/lib/SOURCES | 1 +
> r600/lib/workitem/get_global_offset.cl | 11 +++++++++++
> 10 files changed, 33 insertions(+), 2 deletions(-)
> create mode 100644 amdgcn/lib/workitem/get_global_offset.cl
> create mode 100644 generic/include/clc/workitem/get_global_offset.h
> create mode 100644 ptx-nvidiacl/lib/workitem/get_global_id.cl
> create mode 100644 r600/lib/workitem/get_global_offset.cl
>
> diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES
> index b7e1d98..3a90aa5 100644
> --- a/amdgcn/lib/SOURCES
> +++ b/amdgcn/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/amdgcn/lib/workitem/get_global_offset.cl b/amdgcn/lib/workitem/get_global_offset.cl
> new file mode 100644
> index 0000000..4ebdd77
> --- /dev/null
> +++ b/amdgcn/lib/workitem/get_global_offset.cl
> @@ -0,0 +1,9 @@
> +#include <clc/clc.h>
> +
> +_CLC_DEF uint get_global_offset(uint dim)
> +{
> + __constant uint * ptr = __builtin_amdgcn_implicitarg_ptr();
> + if (dim < 3)
> + return ptr[dim +1];
> + return 1;
> +}
> diff --git a/amdgcn/lib/workitem/get_work_dim.cl b/amdgcn/lib/workitem/get_work_dim.cl
> index ed352b2..07052ac 100644
> --- a/amdgcn/lib/workitem/get_work_dim.cl
> +++ b/amdgcn/lib/workitem/get_work_dim.cl
> @@ -2,6 +2,6 @@
>
> _CLC_DEF uint get_work_dim()
> {
> - __global uint * ptr = __builtin_amdgcn_implicitarg_ptr();
> + __constant uint * ptr = __builtin_amdgcn_implicitarg_ptr();
> return *ptr;
> }
> diff --git a/generic/include/clc/clc.h b/generic/include/clc/clc.h
> index 6694f03..f77e495 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..630156e
> --- /dev/null
> +++ b/generic/include/clc/workitem/get_global_offset.h
> @@ -0,0 +1,2 @@
> +_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 1859b98..5b68ce1 100644
> --- a/r600/lib/SOURCES
> +++ b/r600/lib/SOURCES
> @@ -1,4 +1,5 @@
> synchronization/barrier_impl.ll
> +workitem/get_global_offset.cl
> workitem/get_global_size.cl
> workitem/get_group_id.cl
> workitem/get_local_id.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..74adcc3
> --- /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)
> +{
> + switch(dim) {
> + case 0: return __builtin_r600_read_global_offset_x();
> + case 1: return __builtin_r600_read_global_offset_y();
> + case 2: return __builtin_r600_read_global_offset_z();
> + default: return 1;
> + }
> +}
> --
> 2.5.5
>
> _______________________________________________
> Libclc-dev mailing list
> Libclc-dev at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/libclc-dev
More information about the Libclc-dev
mailing list