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

[Jan Vesely via Libclc-dev]

On Fri, 2016-07-22 at 09:36 -0400, Tom Stellard wrote:
> 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 I use __constant uint *, clang complains that the cast changes AS
(guess it does not know that AS 2 is constant AS at this point).
It's also more consistent with R600 that uses AS 7.

Jan

> 
> > 
> > +	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
> > 
-- 
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <http://lists.llvm.org/pipermail/libclc-dev/attachments/20160722/0c051a2f/attachment.sig>