diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_global_size.cl
@@ -15,10 +15,9 @@
 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;
 }
diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_local_size.cl
@@ -15,16 +15,15 @@
 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;
 }
diff --git a/libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl b/libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl
--- a/libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl
+++ b/libclc/amdgcn-amdhsa/lib/workitem/get_num_groups.cl
@@ -1,7 +1,7 @@
 
 #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;
diff --git a/libclc/amdgcn/lib/mem_fence/fence.cl b/libclc/amdgcn/lib/mem_fence/fence.cl
--- a/libclc/amdgcn/lib/mem_fence/fence.cl
+++ b/libclc/amdgcn/lib/mem_fence/fence.cl
@@ -17,24 +17,21 @@
 _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);
 }
diff --git a/libclc/amdgcn/lib/synchronization/barrier.cl b/libclc/amdgcn/lib/synchronization/barrier.cl
--- a/libclc/amdgcn/lib/synchronization/barrier.cl
+++ b/libclc/amdgcn/lib/synchronization/barrier.cl
@@ -1,7 +1,6 @@
 #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();
 }
diff --git a/libclc/amdgcn/lib/workitem/get_global_offset.cl b/libclc/amdgcn/lib/workitem/get_global_offset.cl
--- a/libclc/amdgcn/lib/workitem/get_global_offset.cl
+++ b/libclc/amdgcn/lib/workitem/get_global_offset.cl
@@ -8,11 +8,9 @@
 #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;
 }
diff --git a/libclc/amdgcn/lib/workitem/get_global_size.cl b/libclc/amdgcn/lib/workitem/get_global_size.cl
--- a/libclc/amdgcn/lib/workitem/get_global_size.cl
+++ b/libclc/amdgcn/lib/workitem/get_global_size.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/amdgcn/lib/workitem/get_group_id.cl b/libclc/amdgcn/lib/workitem/get_group_id.cl
--- a/libclc/amdgcn/lib/workitem/get_group_id.cl
+++ b/libclc/amdgcn/lib/workitem/get_group_id.cl
@@ -1,11 +1,14 @@
 #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;
+  }
 }
diff --git a/libclc/amdgcn/lib/workitem/get_local_id.cl b/libclc/amdgcn/lib/workitem/get_local_id.cl
--- a/libclc/amdgcn/lib/workitem/get_local_id.cl
+++ b/libclc/amdgcn/lib/workitem/get_local_id.cl
@@ -1,11 +1,14 @@
 #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;
+  }
 }
diff --git a/libclc/amdgcn/lib/workitem/get_local_size.cl b/libclc/amdgcn/lib/workitem/get_local_size.cl
--- a/libclc/amdgcn/lib/workitem/get_local_size.cl
+++ b/libclc/amdgcn/lib/workitem/get_local_size.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/amdgcn/lib/workitem/get_num_groups.cl b/libclc/amdgcn/lib/workitem/get_num_groups.cl
--- a/libclc/amdgcn/lib/workitem/get_num_groups.cl
+++ b/libclc/amdgcn/lib/workitem/get_num_groups.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/amdgcn/lib/workitem/get_work_dim.cl b/libclc/amdgcn/lib/workitem/get_work_dim.cl
--- a/libclc/amdgcn/lib/workitem/get_work_dim.cl
+++ b/libclc/amdgcn/lib/workitem/get_work_dim.cl
@@ -8,9 +8,7 @@
 #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];
 }
diff --git a/libclc/generic/include/clc/async/wait_group_events.h b/libclc/generic/include/clc/async/wait_group_events.h
--- a/libclc/generic/include/clc/async/wait_group_events.h
+++ b/libclc/generic/include/clc/async/wait_group_events.h
@@ -1 +1,2 @@
-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);
diff --git a/libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h b/libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h
--- a/libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h
+++ b/libclc/generic/include/clc/explicit_fence/explicit_memory_fence.h
@@ -1,3 +1,3 @@
-_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);
diff --git a/libclc/generic/include/clc/synchronization/barrier.h b/libclc/generic/include/clc/synchronization/barrier.h
--- a/libclc/generic/include/clc/synchronization/barrier.h
+++ b/libclc/generic/include/clc/synchronization/barrier.h
@@ -1 +1 @@
-_CLC_DECL void barrier(cl_mem_fence_flags flags);
+_CLC_DECL _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags);
diff --git a/libclc/generic/include/clc/workitem/get_global_id.h b/libclc/generic/include/clc/workitem/get_global_id.h
--- a/libclc/generic/include/clc/workitem/get_global_id.h
+++ b/libclc/generic/include/clc/workitem/get_global_id.h
@@ -1 +1 @@
-_CLC_DECL size_t get_global_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_id(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_global_offset.h b/libclc/generic/include/clc/workitem/get_global_offset.h
--- a/libclc/generic/include/clc/workitem/get_global_offset.h
+++ b/libclc/generic/include/clc/workitem/get_global_offset.h
@@ -1 +1 @@
-_CLC_DECL size_t get_global_offset(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_offset(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_global_size.h b/libclc/generic/include/clc/workitem/get_global_size.h
--- a/libclc/generic/include/clc/workitem/get_global_size.h
+++ b/libclc/generic/include/clc/workitem/get_global_size.h
@@ -1 +1 @@
-_CLC_DECL size_t get_global_size(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_global_size(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_group_id.h b/libclc/generic/include/clc/workitem/get_group_id.h
--- a/libclc/generic/include/clc/workitem/get_group_id.h
+++ b/libclc/generic/include/clc/workitem/get_group_id.h
@@ -1 +1 @@
-_CLC_DECL size_t get_group_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_group_id(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_local_id.h b/libclc/generic/include/clc/workitem/get_local_id.h
--- a/libclc/generic/include/clc/workitem/get_local_id.h
+++ b/libclc/generic/include/clc/workitem/get_local_id.h
@@ -1 +1 @@
-_CLC_DECL size_t get_local_id(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_local_id(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_local_size.h b/libclc/generic/include/clc/workitem/get_local_size.h
--- a/libclc/generic/include/clc/workitem/get_local_size.h
+++ b/libclc/generic/include/clc/workitem/get_local_size.h
@@ -1 +1 @@
-_CLC_DECL size_t get_local_size(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_local_size(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_num_groups.h b/libclc/generic/include/clc/workitem/get_num_groups.h
--- a/libclc/generic/include/clc/workitem/get_num_groups.h
+++ b/libclc/generic/include/clc/workitem/get_num_groups.h
@@ -1 +1 @@
-_CLC_DECL size_t get_num_groups(uint dim);
+_CLC_DECL _CLC_OVERLOAD size_t get_num_groups(uint dim);
diff --git a/libclc/generic/include/clc/workitem/get_work_dim.h b/libclc/generic/include/clc/workitem/get_work_dim.h
--- a/libclc/generic/include/clc/workitem/get_work_dim.h
+++ b/libclc/generic/include/clc/workitem/get_work_dim.h
@@ -1 +1 @@
-_CLC_DECL uint get_work_dim(void);
+_CLC_DECL _CLC_OVERLOAD uint get_work_dim(void);
diff --git a/libclc/generic/lib/async/wait_group_events.cl b/libclc/generic/lib/async/wait_group_events.cl
--- a/libclc/generic/lib/async/wait_group_events.cl
+++ b/libclc/generic/lib/async/wait_group_events.cl
@@ -1,5 +1,6 @@
 #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);
 }
diff --git a/libclc/generic/lib/workitem/get_global_id.cl b/libclc/generic/lib/workitem/get_global_id.cl
--- a/libclc/generic/lib/workitem/get_global_id.cl
+++ b/libclc/generic/lib/workitem/get_global_id.cl
@@ -1,5 +1,5 @@
 #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);
 }
diff --git a/libclc/generic/lib/workitem/get_global_size.cl b/libclc/generic/lib/workitem/get_global_size.cl
--- a/libclc/generic/lib/workitem/get_global_size.cl
+++ b/libclc/generic/lib/workitem/get_global_size.cl
@@ -1,5 +1,5 @@
 #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);
 }
diff --git a/libclc/ptx-nvidiacl/lib/mem_fence/fence.cl b/libclc/ptx-nvidiacl/lib/mem_fence/fence.cl
--- a/libclc/ptx-nvidiacl/lib/mem_fence/fence.cl
+++ b/libclc/ptx-nvidiacl/lib/mem_fence/fence.cl
@@ -1,15 +1,15 @@
 #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) {
+  if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))
+    __nvvm_membar_cta();
 }
 
 // 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);
 }
diff --git a/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl b/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl
--- a/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl
+++ b/libclc/ptx-nvidiacl/lib/synchronization/barrier.cl
@@ -1,6 +1,5 @@
 #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();
 }
-
diff --git a/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl b/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl
--- a/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl
+++ b/libclc/ptx-nvidiacl/lib/workitem/get_global_id.cl
@@ -1,5 +1,5 @@
 #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);
 }
diff --git a/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl b/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl
--- a/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl
+++ b/libclc/ptx-nvidiacl/lib/workitem/get_group_id.cl
@@ -1,6 +1,6 @@
 #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();
diff --git a/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl b/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl
--- a/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl
+++ b/libclc/ptx-nvidiacl/lib/workitem/get_local_id.cl
@@ -1,6 +1,6 @@
 #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();
diff --git a/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl b/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl
--- a/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl
+++ b/libclc/ptx-nvidiacl/lib/workitem/get_local_size.cl
@@ -1,6 +1,6 @@
 #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();
diff --git a/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl b/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl
--- a/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl
+++ b/libclc/ptx-nvidiacl/lib/workitem/get_num_groups.cl
@@ -1,6 +1,6 @@
 #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();
diff --git a/libclc/r600/lib/synchronization/barrier.cl b/libclc/r600/lib/synchronization/barrier.cl
--- a/libclc/r600/lib/synchronization/barrier.cl
+++ b/libclc/r600/lib/synchronization/barrier.cl
@@ -2,8 +2,7 @@
 
 _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();
 }
diff --git a/libclc/r600/lib/workitem/get_global_offset.cl b/libclc/r600/lib/workitem/get_global_offset.cl
--- a/libclc/r600/lib/workitem/get_global_offset.cl
+++ b/libclc/r600/lib/workitem/get_global_offset.cl
@@ -1,11 +1,10 @@
 #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;
 }
diff --git a/libclc/r600/lib/workitem/get_global_size.cl b/libclc/r600/lib/workitem/get_global_size.cl
--- a/libclc/r600/lib/workitem/get_global_size.cl
+++ b/libclc/r600/lib/workitem/get_global_size.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_group_id.cl b/libclc/r600/lib/workitem/get_group_id.cl
--- a/libclc/r600/lib/workitem/get_group_id.cl
+++ b/libclc/r600/lib/workitem/get_group_id.cl
@@ -1,11 +1,14 @@
 #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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_local_id.cl b/libclc/r600/lib/workitem/get_local_id.cl
--- a/libclc/r600/lib/workitem/get_local_id.cl
+++ b/libclc/r600/lib/workitem/get_local_id.cl
@@ -1,11 +1,14 @@
 #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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_local_size.cl b/libclc/r600/lib/workitem/get_local_size.cl
--- a/libclc/r600/lib/workitem/get_local_size.cl
+++ b/libclc/r600/lib/workitem/get_local_size.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_num_groups.cl b/libclc/r600/lib/workitem/get_num_groups.cl
--- a/libclc/r600/lib/workitem/get_num_groups.cl
+++ b/libclc/r600/lib/workitem/get_num_groups.cl
@@ -4,12 +4,15 @@
 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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_work_dim.cl b/libclc/r600/lib/workitem/get_work_dim.cl
--- a/libclc/r600/lib/workitem/get_work_dim.cl
+++ b/libclc/r600/lib/workitem/get_work_dim.cl
@@ -1,9 +1,8 @@
 #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];
 }