amdgcn: Use new workitem intrinsics

llvm-svn: 261042

9 files changed, 124 insertions, 38 deletions

diff --git a/libclc/amdgcn/lib/SOURCES b/libclc/amdgcn/lib/SOURCES
index c99f3fc7052..49c8dd53a56 100644
--- a/libclc/amdgcn/lib/SOURCES
+++ b/libclc/amdgcn/lib/SOURCES
@@ -1 +1,3 @@
 synchronization/barrier_impl.ll
+workitem/get_group_id.ll
+workitem/get_local_id.ll
diff --git a/libclc/amdgcn/lib/workitem/get_group_id.ll b/libclc/amdgcn/lib/workitem/get_group_id.ll
new file mode 100644
index 00000000000..9d820e0b3e4
--- /dev/null
+++ b/libclc/amdgcn/lib/workitem/get_group_id.ll
@@ -0,0 +1,29 @@
+declare i32 @llvm.amdgcn.workgroup.id.x() #0
+declare i32 @llvm.amdgcn.workgroup.id.y() #0
+declare i32 @llvm.amdgcn.workgroup.id.z() #0
+
+define i32 @get_group_id(i32 %dim) #1 {
+  switch i32 %dim, label %default [
+    i32 0, label %x_dim
+    i32 1, label %y_dim
+    i32 2, label %z_dim
+  ]
+
+x_dim:
+  %x = tail call i32 @llvm.amdgcn.workgroup.id.x()
+  ret i32 %x
+
+y_dim:
+  %y = tail call i32 @llvm.amdgcn.workgroup.id.y()
+  ret i32 %y
+
+z_dim:
+  %z = tail call i32 @llvm.amdgcn.workgroup.id.z()
+  ret i32 %z
+
+default:
+  ret i32 0
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { alwaysinline norecurse nounwind readnone }
diff --git a/libclc/amdgcn/lib/workitem/get_local_id.ll b/libclc/amdgcn/lib/workitem/get_local_id.ll
new file mode 100644
index 00000000000..c54291c0b8f
--- /dev/null
+++ b/libclc/amdgcn/lib/workitem/get_local_id.ll
@@ -0,0 +1,31 @@
+declare i32 @llvm.amdgcn.workitem.id.x() #0
+declare i32 @llvm.amdgcn.workitem.id.y() #0
+declare i32 @llvm.amdgcn.workitem.id.z() #0
+
+define i32 @get_local_id(i32 %dim) #1 {
+  switch i32 %dim, label %default [
+    i32 0, label %x_dim
+    i32 1, label %y_dim
+    i32 2, label %z_dim
+  ]
+
+x_dim:
+  %x = tail call i32 @llvm.amdgcn.workitem.id.x(), !range !0
+  ret i32 %x
+
+y_dim:
+  %y = tail call i32 @llvm.amdgcn.workitem.id.y(), !range !0
+  ret i32 %y
+
+z_dim:
+  %z = tail call i32 @llvm.amdgcn.workitem.id.z(), !range !0
+  ret i32 %z
+
+default:
+  ret i32 0
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { alwaysinline norecurse nounwind readnone }
+
+!0 = !{ i32 0, i32 2048 }
diff --git a/libclc/amdgpu/lib/SOURCES b/libclc/amdgpu/lib/SOURCES
index 7505f3f9692..0f99fe17fea 100644
--- a/libclc/amdgpu/lib/SOURCES
+++ b/libclc/amdgpu/lib/SOURCES
@@ -3,9 +3,7 @@ math/ldexp.cl
 math/nextafter.cl
 math/sqrt.cl
 workitem/get_num_groups.ll
-workitem/get_group_id.ll
 workitem/get_local_size.ll
-workitem/get_local_id.ll
 workitem/get_global_size.ll
 workitem/get_work_dim.ll
 synchronization/barrier.cl
diff --git a/libclc/amdgpu/lib/workitem/get_group_id.ll b/libclc/amdgpu/lib/workitem/get_group_id.ll
deleted file mode 100644
index 0dc86e5edfe..00000000000
--- a/libclc/amdgpu/lib/workitem/get_group_id.ll
+++ /dev/null
@@ -1,18 +0,0 @@
-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_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.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
-}
diff --git a/libclc/amdgpu/lib/workitem/get_local_id.ll b/libclc/amdgpu/lib/workitem/get_local_id.ll
deleted file mode 100644
index ac5522a7822..00000000000
--- a/libclc/amdgpu/lib/workitem/get_local_id.ll
+++ /dev/null
@@ -1,18 +0,0 @@
-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/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES
index c99f3fc7052..49c8dd53a56 100644
--- a/libclc/r600/lib/SOURCES
+++ b/libclc/r600/lib/SOURCES
@@ -1 +1,3 @@
 synchronization/barrier_impl.ll
+workitem/get_group_id.ll
+workitem/get_local_id.ll
diff --git a/libclc/r600/lib/workitem/get_group_id.ll b/libclc/r600/lib/workitem/get_group_id.ll
new file mode 100644
index 00000000000..837c799395d
--- /dev/null
+++ b/libclc/r600/lib/workitem/get_group_id.ll
@@ -0,0 +1,29 @@
+declare i32 @llvm.r600.read.tgid.x() #0
+declare i32 @llvm.r600.read.tgid.y() #0
+declare i32 @llvm.r600.read.tgid.z() #0
+
+define i32 @get_group_id(i32 %dim) #1 {
+  switch i32 %dim, label %default [
+    i32 0, label %x_dim
+    i32 1, label %y_dim
+    i32 2, label %z_dim
+  ]
+
+x_dim:
+  %x = tail call i32 @llvm.r600.read.tgid.x()
+  ret i32 %x
+
+y_dim:
+  %y = tail call i32 @llvm.r600.read.tgid.y()
+  ret i32 %y
+
+z_dim:
+  %z = tail call i32 @llvm.r600.read.tgid.z()
+  ret i32 %z
+
+default:
+  ret i32 0
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { alwaysinline norecurse nounwind readnone }
diff --git a/libclc/r600/lib/workitem/get_local_id.ll b/libclc/r600/lib/workitem/get_local_id.ll
new file mode 100644
index 00000000000..da37ca0c7b4
--- /dev/null
+++ b/libclc/r600/lib/workitem/get_local_id.ll
@@ -0,0 +1,31 @@
+declare i32 @llvm.r600.read.tidig.x() #0
+declare i32 @llvm.r600.read.tidig.y() #0
+declare i32 @llvm.r600.read.tidig.z() #0
+
+define i32 @get_local_id(i32 %dim) #1 {
+  switch i32 %dim, label %default [
+    i32 0, label %x_dim
+    i32 1, label %y_dim
+    i32 2, label %z_dim
+  ]
+
+x_dim:
+  %x = tail call i32 @llvm.r600.read.tidig.x(), !range !0
+  ret i32 %x
+
+y_dim:
+  %y = tail call i32 @llvm.r600.read.tidig.y(), !range !0
+  ret i32 %y
+z_dim:
+
+  %z = tail call i32 @llvm.r600.read.tidig.z(), !range !0
+  ret i32 %z
+
+default:
+  ret i32 0
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { alwaysinline norecurse nounwind readnone }
+
+!0 = !{ i32 0, i32 2048 }