R600: Replace cl implementations with LLVM IR implementation

This allows libclc to be built for R600 with a patched clang +
unpatch LLVM by dropping the dependency on R600 Intrinsics

13 files changed, 76 insertions, 36 deletions

diff --git a/r600/include/clc/workitem/get_global_id.h b/r600/include/clc/workitem/get_global_id.h
deleted file mode 100644
index b61450f..0000000
--- a/r600/include/clc/workitem/get_global_id.h
+++ /dev/null
@@ -1 +0,0 @@
-size_t get_global_id(uint dim);
diff --git a/r600/include/clc/workitem/get_global_size.h b/r600/include/clc/workitem/get_global_size.h
deleted file mode 100644
index 8d7b9a1..0000000
--- a/r600/include/clc/workitem/get_global_size.h
+++ /dev/null
@@ -1 +0,0 @@
-size_t get_global_size(uint dim);
diff --git a/r600/include/clc/workitem/get_group_id.h b/r600/include/clc/workitem/get_group_id.h
deleted file mode 100644
index 6862dba..0000000
--- a/r600/include/clc/workitem/get_group_id.h
+++ /dev/null
@@ -1,3 +0,0 @@
-_CLC_INLINE size_t get_group_id(uint dim) {
-  return 0;
-}
diff --git a/r600/include/clc/workitem/get_local_id.h b/r600/include/clc/workitem/get_local_id.h
deleted file mode 100644
index 22749cd..0000000
--- a/r600/include/clc/workitem/get_local_id.h
+++ /dev/null
@@ -1,3 +0,0 @@
-_CLC_INLINE size_t get_local_id(uint dim) {
-  return 0;
-}
diff --git a/r600/include/clc/workitem/get_local_size.h b/r600/include/clc/workitem/get_local_size.h
deleted file mode 100644
index 51d9762..0000000
--- a/r600/include/clc/workitem/get_local_size.h
+++ /dev/null
@@ -1,3 +0,0 @@
-_CLC_INLINE size_t get_local_size(uint dim) {
-  return 0;
-}
diff --git a/r600/include/clc/workitem/get_num_groups.h b/r600/include/clc/workitem/get_num_groups.h
deleted file mode 100644
index fe1f343..0000000
--- a/r600/include/clc/workitem/get_num_groups.h
+++ /dev/null
@@ -1,3 +0,0 @@
-_CLC_INLINE size_t get_num_groups(uint dim) {
-  return 0;
-}
diff --git a/r600/lib/SOURCES b/r600/lib/SOURCES
index 644d2f3..af8c8c8 100644
--- a/r600/lib/SOURCES
+++ b/r600/lib/SOURCES
@@ -1,2 +1,4 @@
-workitem/get_global_id.cl
-workitem/get_global_size.cl
+workitem/get_group_id.ll
+workitem/get_local_size.ll
+workitem/get_local_id.ll
+workitem/get_global_size.ll
diff --git a/r600/lib/workitem/get_global_id.cl b/r600/lib/workitem/get_global_id.cl
deleted file mode 100644
index 671f657..0000000
--- a/r600/lib/workitem/get_global_id.cl
+++ /dev/null
@@ -1,10 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_global_id(uint dim) {
-  switch (dim) {
-  case 0:  return __builtin_r600_read_tgid_x()*__builtin_r600_read_local_size_x()+__builtin_r600_read_tidig_x();
-  case 1:  return __builtin_r600_read_tgid_y()*__builtin_r600_read_local_size_y()+__builtin_r600_read_tidig_y();
-  case 2:  return __builtin_r600_read_tgid_z()*__builtin_r600_read_local_size_z()+__builtin_r600_read_tidig_z();
-  default: return 0;
-  }
-}
diff --git a/r600/lib/workitem/get_global_size.cl b/r600/lib/workitem/get_global_size.cl
deleted file mode 100644
index 4e47bbe..0000000
--- a/r600/lib/workitem/get_global_size.cl
+++ /dev/null
@@ -1,10 +0,0 @@
-#include <clc/clc.h>
-
-_CLC_DEF size_t get_global_size(uint dim) {
-  switch (dim) {
-  case 0: return __builtin_r600_read_global_size_x();
-  case 1: return __builtin_r600_read_global_size_y();
-  case 2: return __builtin_r600_read_global_size_z();
-  default: return 1;
-  }
-}
diff --git a/r600/lib/workitem/get_global_size.ll b/r600/lib/workitem/get_global_size.ll
new file mode 100644
index 0000000..ac2d08d
--- /dev/null
+++ b/r600/lib/workitem/get_global_size.ll
@@ -0,0 +1,18 @@
+declare i32 @llvm.r600.read.global.size.x() nounwind readnone
+declare i32 @llvm.r600.read.global.size.y() nounwind readnone
+declare i32 @llvm.r600.read.global.size.z() nounwind readnone
+
+define i32 @get_global_size(i32 %dim) nounwind readnone alwaysinline {
+  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
+x_dim:
+  %x = call i32 @llvm.r600.read.global.size.x() nounwind readnone
+  ret i32 %x
+y_dim:
+  %y = call i32 @llvm.r600.read.global.size.y() nounwind readnone
+  ret i32 %y
+z_dim:
+  %z = call i32 @llvm.r600.read.global.size.z() nounwind readnone
+  ret i32 %z
+default:
+  ret i32 0
+}
diff --git a/r600/lib/workitem/get_group_id.ll b/r600/lib/workitem/get_group_id.ll
new file mode 100644
index 0000000..5131648
--- /dev/null
+++ b/r600/lib/workitem/get_group_id.ll
@@ -0,0 +1,18 @@
+declare i32 @llvm.r600.read.local.size.x() nounwind readnone
+declare i32 @llvm.r600.read.local.size.y() nounwind readnone
+declare i32 @llvm.r600.read.local.size.z() nounwind readnone
+
+define i32 @get_group_id(i32 %dim) nounwind readnone alwaysinline {
+  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
+x_dim:
+  %x = call i32 @llvm.r600.read.local.size.x() nounwind readnone
+  ret i32 %x
+y_dim:
+  %y = call i32 @llvm.r600.read.local.size.y() nounwind readnone
+  ret i32 %y
+z_dim:
+  %z = call i32 @llvm.r600.read.local.size.z() nounwind readnone
+  ret i32 %z
+default:
+  ret i32 0
+}
diff --git a/r600/lib/workitem/get_local_id.ll b/r600/lib/workitem/get_local_id.ll
new file mode 100644
index 0000000..ac5522a
--- /dev/null
+++ b/r600/lib/workitem/get_local_id.ll
@@ -0,0 +1,18 @@
+declare i32 @llvm.r600.read.tidig.x() nounwind readnone
+declare i32 @llvm.r600.read.tidig.y() nounwind readnone
+declare i32 @llvm.r600.read.tidig.z() nounwind readnone
+
+define i32 @get_local_id(i32 %dim) nounwind readnone alwaysinline {
+  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
+x_dim:
+  %x = call i32 @llvm.r600.read.tidig.x() nounwind readnone
+  ret i32 %x
+y_dim:
+  %y = call i32 @llvm.r600.read.tidig.y() nounwind readnone
+  ret i32 %y
+z_dim:
+  %z = call i32 @llvm.r600.read.tidig.z() nounwind readnone
+  ret i32 %z
+default:
+  ret i32 0
+}
diff --git a/r600/lib/workitem/get_local_size.ll b/r600/lib/workitem/get_local_size.ll
new file mode 100644
index 0000000..6a71f75
--- /dev/null
+++ b/r600/lib/workitem/get_local_size.ll
@@ -0,0 +1,18 @@
+declare i32 @llvm.r600.read.tgid.x() nounwind readnone
+declare i32 @llvm.r600.read.tgid.y() nounwind readnone
+declare i32 @llvm.r600.read.tgid.z() nounwind readnone
+
+define i32 @get_local_size(i32 %dim) nounwind readnone alwaysinline {
+  switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim]
+x_dim:
+  %x = call i32 @llvm.r600.read.tgid.x() nounwind readnone
+  ret i32 %x
+y_dim:
+  %y = call i32 @llvm.r600.read.tgid.y() nounwind readnone
+  ret i32 %y
+z_dim:
+  %z = call i32 @llvm.r600.read.tgid.z() nounwind readnone
+  ret i32 %z
+default:
+  ret i32 0
+}