[Libclc-dev] [PATCH 2/2] amdgcn--mesa3d: Implement get_num_groups()

[Jan Vesely via Libclc-dev]

On Tue, 2016-07-26 at 14:55 +0000, Tom Stellard wrote:
> On Tue, Jul 26, 2016 at 10:42:32AM -0400, Jan Vesely via Libclc-dev
> wrote:
> > 
> > On Tue, 2016-07-26 at 00:00 +0000, Tom Stellard via Libclc-dev
> > wrote:
> > > 
> > > ---
> > >  amdgcn-mesa3d/lib/OVERRIDES                  |  1 +
> > >  amdgcn-mesa3d/lib/SOURCES                    |  1 +
> > >  amdgcn-mesa3d/lib/workitem/get_num_groups.cl | 10 ++++++++++
> > >  3 files changed, 12 insertions(+)
> > >  create mode 100644 amdgcn-mesa3d/lib/OVERRIDES
> > >  create mode 100644 amdgcn-mesa3d/lib/SOURCES
> > >  create mode 100644 amdgcn-mesa3d/lib/workitem/get_num_groups.cl
> > > 
> > > diff --git a/amdgcn-mesa3d/lib/OVERRIDES b/amdgcn-
> > > mesa3d/lib/OVERRIDES
> > > new file mode 100644
> > > index 0000000..c9bd69b
> > > --- /dev/null
> > > +++ b/amdgcn-mesa3d/lib/OVERRIDES
> > > @@ -0,0 +1 @@
> > > +workitem/get_num_groups.ll
> > > diff --git a/amdgcn-mesa3d/lib/SOURCES b/amdgcn-
> > > mesa3d/lib/SOURCES
> > > new file mode 100644
> > > index 0000000..77eda67
> > > --- /dev/null
> > > +++ b/amdgcn-mesa3d/lib/SOURCES
> > > @@ -0,0 +1 @@
> > > +workitem/get_num_groups.cl
> > > diff --git a/amdgcn-mesa3d/lib/workitem/get_num_groups.cl
> > > b/amdgcn-
> > > mesa3d/lib/workitem/get_num_groups.cl
> > > new file mode 100644
> > > index 0000000..3a4a6f4
> > > --- /dev/null
> > > +++ b/amdgcn-mesa3d/lib/workitem/get_num_groups.cl
> > > @@ -0,0 +1,10 @@
> > > +#include <clc/clc.h>
> > > +
> > > +_CLC_DEF size_t get_num_groups(uint dim) {
> > > +  switch (dim) {
> > > +  case 0: return __builtin_amdgcn_workgroup_count_x();
> > > +  case 1: return __builtin_amdgcn_workgroup_count_y();
> > > +  case 2: return __builtin_amdgcn_workgroup_count_z();
> > 
> > I planned to switch r600 to reading these (and global size and
> > local
> > size) from implicitarg ptr. Is there an advantage to gcn reading
> > this
> > from dispatch info, rather than implicitarg ptr?
> > 
> 
> These intrinsics read these values from the user SGPRs and not the
> dispatch info.  This is really the only difference between what Mesa
> does and what HSA does.

The questions are mostly out of my curiosity, I can't really comment on
HSA patches.
I take it SGPRs are faster?
If these intrinsics follow HSA ABI shouldn't they be in amdgcn-amdhsa
directory?
The first patch introduced amdgcn-mesa3d directory, but if mesa/clover
is switched to use HSA ABI shouldn't libclc produced for amdgcn-amdhsa
work for clover as well?

thanks,
Jan

> 
> -Tom
> 
> > 
> > Jan
> > 
> > > 
> > > +  default: return 0;
> > > +  }
> > > +}
> > -- 
> > Jan Vesely <jan.vesely at rutgers.edu>
> 
> 
> 
> > 
> > _______________________________________________
> > Libclc-dev mailing list
> > Libclc-dev at lists.llvm.org
> > http://lists.llvm.org/cgi-bin/mailman/listinfo/libclc-dev
> 
-- 
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/20160726/86755674/attachment.sig>