libclc: Make all built-ins overloadable

[daniels]

Summary

The SPIR spec states that all OpenCL built-in functions should be
overloadable and mangled, to ensure consistency.

Add the overload attribute to functions which were missing them:
work dimensions, memory barriers and fences, and events.

[daniels]

This fixes things for us running when emitting SPIR-V, as the LLVM -> SPIR-V translator assumes (per the SPIR spec) that all builtins are overloaded and mangled.

The unrelated formatting changes here are from clang-format automatically running under arc.

[jenatali]

Looks good to me. Verified that the async functions now work in SPIR-V with this patch.

[daniels]

CI failures are due to not ignoring CL source files when running clang-tidy, per previously-submitted PR, as well as an unrelated failure in a clang-go test.

[daniels]

@jvesely @tstellar Hi, any thoughts on this change please? We need this in order to have a fully working SPIR-V backend.

[tstellar]

LGTM.

libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl

 
 #if __clang_major__ >= 6
 #define __dispatch_ptr __builtin_amdgcn_dispatch_ptr
 #else
 #define __dispatch_ptr __clc_amdgcn_dispatch_ptr
 CONST_AS uchar * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
 #endif
 
-_CLC_DEF size_t get_global_size(uint dim)
-{
-  CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
-  if (dim < 3)
-    return ptr[3 + dim];
-  return 1;
+_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
+  CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr();
+  if (dim < 3)
+    return ptr[3 + dim];
+  return 1;
 }

libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl

 
 #if __clang_major__ >= 6
 #define __dispatch_ptr __builtin_amdgcn_dispatch_ptr
 #else
 #define __dispatch_ptr __clc_amdgcn_dispatch_ptr
 CONST_AS char * __clc_amdgcn_dispatch_ptr(void) __asm("llvm.amdgcn.dispatch.ptr");
 #endif
 
-_CLC_DEF size_t get_local_size(uint dim)
-{
-  CONST_AS uint * ptr = (CONST_AS uint *) __dispatch_ptr();
-  switch (dim) {
-  case 0:
-    return ptr[1] & 0xffffu;
-  case 1:
-    return ptr[1] >> 16;
-  case 2:
-    return ptr[2] & 0xffffu;
-  }
-  return 1;
+_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
+  CONST_AS uint *ptr = (CONST_AS uint *)__dispatch_ptr();
+  switch (dim) {
+  case 0:
+    return ptr[1] & 0xffffu;
+  case 1:
+    return ptr[1] >> 16;
+  case 2:
+    return ptr[2] & 0xffffu;
+  }
+  return 1;
 }

libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl

 
 #include <clc/clc.h>
 
-_CLC_DEF size_t get_num_groups(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
   size_t global_size = get_global_size(dim);
   size_t local_size = get_local_size(dim);
   size_t num_groups = global_size / local_size;
   if (global_size % local_size != 0) {
     num_groups++;
   }
   return num_groups;
 }

libclc/amdgcn/lib/mem_fence/fence.cl

 // Newer clang supports __builtin_amdgcn_s_waitcnt
 #if __clang_major__ >= 5
 #  define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
 #else
 #  define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
 _CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned)  __asm("llvm.amdgcn.s.waitcnt");
 #endif
 
-_CLC_DEF void mem_fence(cl_mem_fence_flags flags)
-{
-  if (flags & CLK_GLOBAL_MEM_FENCE) {
-    // scalar loads are counted with LGKM but we don't know whether
-    // the compiler turned any loads to scalar
-    __waitcnt(0);
-  } else if (flags & CLK_LOCAL_MEM_FENCE)
-    __waitcnt(0xff); // LGKM is [12:8]
+_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
+  if (flags & CLK_GLOBAL_MEM_FENCE) {
+    // scalar loads are counted with LGKM but we don't know whether
+    // the compiler turned any loads to scalar
+    __waitcnt(0);
+  } else if (flags & CLK_LOCAL_MEM_FENCE)
+    __waitcnt(0xff); // LGKM is [12:8]
 }
 #undef __waitcnt
 
 // We don't have separate mechanism for read and write fences
-_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags)
-{
-  mem_fence(flags);
+_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
+  mem_fence(flags);
 }
 
-_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags)
-{
-  mem_fence(flags);
+_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
+  mem_fence(flags);
 }

libclc/amdgcn/lib/synchronization/barrier.cl

 #include <clc/clc.h>
 
-_CLC_DEF void barrier(cl_mem_fence_flags flags)
-{
-  mem_fence(flags);
-  __builtin_amdgcn_s_barrier();
+_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
+  mem_fence(flags);
+  __builtin_amdgcn_s_barrier();
 }

libclc/amdgcn/lib/workitem/get_global_offset.cl

 #include <clc/clc.h>
 
 #if __clang_major__ >= 8
 #define CONST_AS __constant
 #elif __clang_major__ >= 7
 #define CONST_AS __attribute__((address_space(4)))
 #else
 #define CONST_AS __attribute__((address_space(2)))
 #endif
 
-_CLC_DEF size_t get_global_offset(uint dim)
-{
-  CONST_AS uint * ptr =
-    (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
-  if (dim < 3)
-    return ptr[dim + 1];
-  return 0;
+_CLC_DEF _CLC_OVERLOAD size_t get_global_offset(uint dim) {
+  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
+  if (dim < 3)
+    return ptr[dim + 1];
+  return 0;
 }

libclc/amdgcn/lib/workitem/get_global_size.cl

 #include <clc/clc.h>
 
 uint __clc_amdgcn_get_global_size_x(void) __asm("llvm.r600.read.global.size.x");
 uint __clc_amdgcn_get_global_size_y(void) __asm("llvm.r600.read.global.size.y");
 uint __clc_amdgcn_get_global_size_z(void) __asm("llvm.r600.read.global.size.z");
 
-_CLC_DEF size_t get_global_size(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_amdgcn_get_global_size_x();
-  case 1: return __clc_amdgcn_get_global_size_y();
-  case 2: return __clc_amdgcn_get_global_size_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_amdgcn_get_global_size_x();
+  case 1:
+    return __clc_amdgcn_get_global_size_y();
+  case 2:
+    return __clc_amdgcn_get_global_size_z();
+  default:
+    return 1;
+  }
 }

libclc/amdgcn/lib/workitem/get_group_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_group_id(uint dim)
-{
-  switch(dim) {
-  case 0: return __builtin_amdgcn_workgroup_id_x();
-  case 1: return __builtin_amdgcn_workgroup_id_y();
-  case 2: return __builtin_amdgcn_workgroup_id_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_amdgcn_workgroup_id_x();
+  case 1:
+    return __builtin_amdgcn_workgroup_id_y();
+  case 2:
+    return __builtin_amdgcn_workgroup_id_z();
+  default:
+    return 1;
+  }
 }

libclc/amdgcn/lib/workitem/get_local_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_local_id(uint dim)
-{
-  switch(dim) {
-  case 0: return __builtin_amdgcn_workitem_id_x();
-  case 1: return __builtin_amdgcn_workitem_id_y();
-  case 2: return __builtin_amdgcn_workitem_id_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_amdgcn_workitem_id_x();
+  case 1:
+    return __builtin_amdgcn_workitem_id_y();
+  case 2:
+    return __builtin_amdgcn_workitem_id_z();
+  default:
+    return 1;
+  }
 }

libclc/amdgcn/lib/workitem/get_local_size.cl

 #include <clc/clc.h>
 
 uint __clc_amdgcn_get_local_size_x(void) __asm("llvm.r600.read.local.size.x");
 uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y");
 uint __clc_amdgcn_get_local_size_z(void) __asm("llvm.r600.read.local.size.z");
 
-_CLC_DEF size_t get_local_size(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_amdgcn_get_local_size_x();
-  case 1: return __clc_amdgcn_get_local_size_y();
-  case 2: return __clc_amdgcn_get_local_size_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_amdgcn_get_local_size_x();
+  case 1:
+    return __clc_amdgcn_get_local_size_y();
+  case 2:
+    return __clc_amdgcn_get_local_size_z();
+  default:
+    return 1;
+  }
 }

libclc/amdgcn/lib/workitem/get_num_groups.cl

 #include <clc/clc.h>
 
 uint __clc_amdgcn_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x");
 uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y");
 uint __clc_amdgcn_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z");
 
-_CLC_DEF size_t get_num_groups(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_amdgcn_get_num_groups_x();
-  case 1: return __clc_amdgcn_get_num_groups_y();
-  case 2: return __clc_amdgcn_get_num_groups_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_amdgcn_get_num_groups_x();
+  case 1:
+    return __clc_amdgcn_get_num_groups_y();
+  case 2:
+    return __clc_amdgcn_get_num_groups_z();
+  default:
+    return 1;
+  }
 }

libclc/amdgcn/lib/workitem/get_work_dim.cl

 #include <clc/clc.h>
 
 #if __clang_major__ >= 8
 #define CONST_AS __constant
 #elif __clang_major__ >= 7
 #define CONST_AS __attribute__((address_space(4)))
 #else
 #define CONST_AS __attribute__((address_space(2)))
 #endif
 
-_CLC_DEF uint get_work_dim(void)
-{
-  CONST_AS uint * ptr =
-    (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
-  return ptr[0];
+_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) {
+  CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr();
+  return ptr[0];
 }

libclc/generic/include/clc/async/wait_group_events.h

-void wait_group_events(int num_events, event_t *event_list);
+_CLC_DECL _CLC_OVERLOAD void wait_group_events(int num_events,
+                                               event_t *event_list);

libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h

-_CLC_DECL void mem_fence(cl_mem_fence_flags flags);
-_CLC_DECL void read_mem_fence(cl_mem_fence_flags flags);
-_CLC_DECL void write_mem_fence(cl_mem_fence_flags flags);
+_CLC_DECL _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags);
+_CLC_DECL _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags);
+_CLC_DECL _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags);

libclc/generic/include/clc/synchronization/barrier.h

-_CLC_DECL void barrier(cl_mem_fence_flags flags);
+_CLC_DECL _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags);

libclc/generic/include/clc/workitem/get_global_id.h

-_CLC_DECL size_t get_global_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);

libclc/generic/include/clc/workitem/get_global_offset.h

-_CLC_DECL size_t get_global_offset(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim);

libclc/generic/include/clc/workitem/get_global_size.h

-_CLC_DECL size_t get_global_size(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim);

libclc/generic/include/clc/workitem/get_group_id.h

-_CLC_DECL size_t get_group_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim);

libclc/generic/include/clc/workitem/get_local_id.h

-_CLC_DECL size_t get_local_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim);

libclc/generic/include/clc/workitem/get_local_size.h

-_CLC_DECL size_t get_local_size(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim);

libclc/generic/include/clc/workitem/get_num_groups.h

-_CLC_DECL size_t get_num_groups(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim);

libclc/generic/include/clc/workitem/get_work_dim.h

-_CLC_DECL uint get_work_dim(void);
+_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void);

libclc/generic/lib/async/wait_group_events.cl

 #include <clc/clc.h>
 
-_CLC_DEF void wait_group_events(int num_events, event_t *event_list) {
+_CLC_DEF _CLC_OVERLOAD void wait_group_events(int num_events,
+                                              event_t *event_list) {
   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
 }

libclc/generic/lib/workitem/get_global_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_global_id(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
   return get_group_id(dim) * get_local_size(dim) + get_local_id(dim) + get_global_offset(dim);
 }

libclc/generic/lib/workitem/get_global_size.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_global_size(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
   return get_num_groups(dim)*get_local_size(dim);
 }

libclc/ptx-nvidiacl/lib/mem_fence/fence.cl

 #include <clc/clc.h>
 
-_CLC_DEF void mem_fence(cl_mem_fence_flags flags) {
-   if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))
-     __nvvm_membar_cta();
+_CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
 }
 
 // We do not have separate mechanism for read and write fences.
-_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags) {
+_CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
   mem_fence(flags);
 }
 
-_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags) {
+_CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
   mem_fence(flags);
 }

libclc/ptx-nvidiacl/lib/synchronization/barrier.cl

 #include <clc/clc.h>
 
-_CLC_DEF void barrier(cl_mem_fence_flags flags) {
+_CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
   __syncthreads();
 }
-

libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_global_id(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_global_id(uint dim) {
   return get_group_id(dim) * get_local_size(dim) + get_local_id(dim);
 }

libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_group_id(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
   switch (dim) {
   case 0:  return __nvvm_read_ptx_sreg_ctaid_x();
   case 1:  return __nvvm_read_ptx_sreg_ctaid_y();
   case 2:  return __nvvm_read_ptx_sreg_ctaid_z();
   default: return 0;
   }
 }

libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_local_id(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
   switch (dim) {
   case 0:  return __nvvm_read_ptx_sreg_tid_x();
   case 1:  return __nvvm_read_ptx_sreg_tid_y();
   case 2:  return __nvvm_read_ptx_sreg_tid_z();
   default: return 0;
   }
 }

libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_local_size(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
   switch (dim) {
   case 0:  return __nvvm_read_ptx_sreg_ntid_x();
   case 1:  return __nvvm_read_ptx_sreg_ntid_y();
   case 2:  return __nvvm_read_ptx_sreg_ntid_z();
   default: return 0;
   }
 }

libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl

 #include <clc/clc.h>
 
-_CLC_DEF size_t get_num_groups(uint dim) {
+_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
   switch (dim) {
   case 0:  return __nvvm_read_ptx_sreg_nctaid_x();
   case 1:  return __nvvm_read_ptx_sreg_nctaid_y();
   case 2:  return __nvvm_read_ptx_sreg_nctaid_z();
   default: return 0;
   }
 }

libclc/r600/lib/synchronization/barrier.cl

 #include <clc/clc.h>
 
 _CLC_DEF void __clc_r600_barrier(void) __asm("llvm.r600.group.barrier");
 
-_CLC_DEF void barrier(uint flags)
-{
+_CLC_DEF _CLC_OVERLOAD void barrier(uint flags) {
   // We should call mem_fence here, but that is not implemented for r600 yet
   __clc_r600_barrier();
 }

libclc/r600/lib/workitem/get_global_offset.cl

 #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;
+_CLC_DEF _CLC_OVERLOAD 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;
 }

libclc/r600/lib/workitem/get_global_size.cl

 #include <clc/clc.h>
 
 uint __clc_r600_get_global_size_x(void) __asm("llvm.r600.read.global.size.x");
 uint __clc_r600_get_global_size_y(void) __asm("llvm.r600.read.global.size.y");
 uint __clc_r600_get_global_size_z(void) __asm("llvm.r600.read.global.size.z");
 
-_CLC_DEF size_t get_global_size(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_r600_get_global_size_x();
-  case 1: return __clc_r600_get_global_size_y();
-  case 2: return __clc_r600_get_global_size_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_global_size(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_r600_get_global_size_x();
+  case 1:
+    return __clc_r600_get_global_size_y();
+  case 2:
+    return __clc_r600_get_global_size_z();
+  default:
+    return 1;
+  }
 }

libclc/r600/lib/workitem/get_group_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF uint get_group_id(uint dim)
-{
-  switch(dim) {
-  case 0: return __builtin_r600_read_tgid_x();
-  case 1: return __builtin_r600_read_tgid_y();
-  case 2: return __builtin_r600_read_tgid_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD uint get_group_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_r600_read_tgid_x();
+  case 1:
+    return __builtin_r600_read_tgid_y();
+  case 2:
+    return __builtin_r600_read_tgid_z();
+  default:
+    return 1;
+  }
 }

libclc/r600/lib/workitem/get_local_id.cl

 #include <clc/clc.h>
 
-_CLC_DEF uint get_local_id(uint dim)
-{
-  switch(dim) {
-  case 0: return __builtin_r600_read_tidig_x();
-  case 1: return __builtin_r600_read_tidig_y();
-  case 2: return __builtin_r600_read_tidig_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD uint get_local_id(uint dim) {
+  switch (dim) {
+  case 0:
+    return __builtin_r600_read_tidig_x();
+  case 1:
+    return __builtin_r600_read_tidig_y();
+  case 2:
+    return __builtin_r600_read_tidig_z();
+  default:
+    return 1;
+  }
 }

libclc/r600/lib/workitem/get_local_size.cl

 #include <clc/clc.h>
 
 uint __clc_r600_get_local_size_x(void) __asm("llvm.r600.read.local.size.x");
 uint __clc_r600_get_local_size_y(void) __asm("llvm.r600.read.local.size.y");
 uint __clc_r600_get_local_size_z(void) __asm("llvm.r600.read.local.size.z");
 
-_CLC_DEF size_t get_local_size(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_r600_get_local_size_x();
-  case 1: return __clc_r600_get_local_size_y();
-  case 2: return __clc_r600_get_local_size_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_local_size(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_r600_get_local_size_x();
+  case 1:
+    return __clc_r600_get_local_size_y();
+  case 2:
+    return __clc_r600_get_local_size_z();
+  default:
+    return 1;
+  }
 }

libclc/r600/lib/workitem/get_num_groups.cl

 #include <clc/clc.h>
 
 uint __clc_r600_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x");
 uint __clc_r600_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y");
 uint __clc_r600_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z");
 
-_CLC_DEF size_t get_num_groups(uint dim)
-{
-  switch (dim) {
-  case 0: return __clc_r600_get_num_groups_x();
-  case 1: return __clc_r600_get_num_groups_y();
-  case 2: return __clc_r600_get_num_groups_z();
-  default: return 1;
-  }
+_CLC_DEF _CLC_OVERLOAD size_t get_num_groups(uint dim) {
+  switch (dim) {
+  case 0:
+    return __clc_r600_get_num_groups_x();
+  case 1:
+    return __clc_r600_get_num_groups_y();
+  case 2:
+    return __clc_r600_get_num_groups_z();
+  default:
+    return 1;
+  }
 }

libclc/r600/lib/workitem/get_work_dim.cl

 #include <clc/clc.h>
 
-_CLC_DEF uint get_work_dim(void)
-{
-  __attribute__((address_space(7))) uint * ptr =
-    (__attribute__((address_space(7))) uint *)
-    __builtin_r600_implicitarg_ptr();
-  return ptr[0];
+_CLC_DEF _CLC_OVERLOAD uint get_work_dim(void) {
+  __attribute__((address_space(7))) uint *ptr =
+      (__attribute__((address_space(7)))
+       uint *)__builtin_r600_implicitarg_ptr();
+  return ptr[0];
 }