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

[Jan Vesely via Libclc-dev]

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?

Jan

> +  default: return 0;
> +  }
> +}
-- 
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/d1139ba6/attachment.sig>