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

[Tom Stellard via Libclc-dev]

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.

-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