libclc: Make all built-ins overloadable

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.

Reviewed By: tstellar, jenatali

Differential Revision: https://reviews.llvm.org/D82078

8 files changed, 68 insertions, 56 deletions

diff --git a/libclc/r600/lib/synchronization/barrier.cl b/libclc/r600/lib/synchronization/barrier.cl
index 98200e7eda92..6a28ee3201de 100644
--- 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
index b38ae3377570..7c2e403ea6ec 100644
--- 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
index d356929c4948..628136150d84 100644
--- 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_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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_group_id.cl b/libclc/r600/lib/workitem/get_group_id.cl
index e5efc0a85778..1fb993ace72e 100644
--- 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
index a871a5d77f0c..80fdc344193b 100644
--- 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
index 89e2612786e4..6edab7c46c2f 100644
--- 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_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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_num_groups.cl b/libclc/r600/lib/workitem/get_num_groups.cl
index dfe6cef22f8e..ab4f5f629c27 100644
--- 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_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;
+  }
 }
diff --git a/libclc/r600/lib/workitem/get_work_dim.cl b/libclc/r600/lib/workitem/get_work_dim.cl
index fccf716cf7c9..e18a83b8dd95 100644
--- 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];
 }