[Libclc-dev] [PATCH v3 2/2] AMDGPU: Implement get_global_offset builtin

[Tom Stellard via Libclc-dev]

On Thu, Jul 14, 2016 at 04:12:42PM -0400, Jan Vesely 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.
> 
> v2: split to a separate patch
> 
> v3: Switch R600 to use implictarg.ptr
> 
> Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> ---
>  amdgcn/lib/SOURCES                               |  1 +
>  amdgcn/lib/workitem/get_global_offset.cl         | 11 +++++++++++
>  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 +++++++++++
>  9 files changed, 34 insertions(+), 1 deletion(-)
>  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 4178d70..33038f2 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..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();

Why did you use __attribute__((address_space(2))) instead of 'constant'?

> +	if (dim < 3)
> +		return ptr[dim + 1];
> +	return 0;
> +}
> 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 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;
> +}
> -- 
> 2.7.4
>