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

[Tom Stellard via Libclc-dev]

On Fri, Jul 22, 2016 at 12:01:02PM -0400, Jan Vesely wrote:
> 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.
> 

Ok, that makes sense. LGTM.

-Tom

> 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>