Author: arsenm
Date: Tue Feb 16 18:27:27 2016
New Revision: 261042
URL: http://llvm.org/viewvc/llvm-project?rev=261042&view=rev
Log:
amdgcn: Use new workitem intrinsics
Added:
libclc/trunk/amdgcn/lib/workitem/
libclc/trunk/amdgcn/lib/workitem/get_group_id.ll
libclc/trunk/amdgcn/lib/workitem/get_local_id.ll
libclc/trunk/r600/lib/workitem/get_group_id.ll
libclc/trunk/r600/lib/workitem/get_local_id.ll
Removed:
libclc/trunk/amdgpu/lib/workitem/get_group_id.ll
libclc/trunk/amdgpu/lib/workitem/get_local_id.ll
Modified:
libclc/trunk/amdgcn/lib/SOURCES
libclc/trunk/amdgpu/lib/SOURCES
libclc/trunk/r600/lib/SOURCES
Modified: libclc/trunk/amdgcn/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES?rev=261042&r1=261041&r2=261042&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES (original)
+++ libclc/trunk/amdgcn/lib/SOURCES Tue Feb 16 18:27:27 2016
@@ -1 +1,3 @@
synchronization/barrier_impl.ll
+workitem/get_group_id.ll
+workitem/get_local_id.ll
Added: libclc/trunk/amdgcn/lib/workitem/get_group_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_group_id.ll?rev=261042&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_group_id.ll (added)
+++ libclc/trunk/amdgcn/lib/workitem/get_group_id.ll Tue Feb 16 18:27:27 2016
@@ -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 }
Added: libclc/trunk/amdgcn/lib/workitem/get_local_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/workitem/get_local_id.ll?rev=261042&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/workitem/get_local_id.ll (added)
+++ libclc/trunk/amdgcn/lib/workitem/get_local_id.ll Tue Feb 16 18:27:27 2016
@@ -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 }
Modified: libclc/trunk/amdgpu/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/SOURCES?rev=261042&r1=261041&r2=261042&view=diff
==============================================================================
--- libclc/trunk/amdgpu/lib/SOURCES (original)
+++ libclc/trunk/amdgpu/lib/SOURCES Tue Feb 16 18:27:27 2016
@@ -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
Removed: libclc/trunk/amdgpu/lib/workitem/get_group_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_group_id.ll?rev=261041&view=auto
==============================================================================
--- libclc/trunk/amdgpu/lib/workitem/get_group_id.ll (original)
+++ libclc/trunk/amdgpu/lib/workitem/get_group_id.ll (removed)
@@ -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
-}
Removed: libclc/trunk/amdgpu/lib/workitem/get_local_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/workitem/get_local_id.ll?rev=261041&view=auto
==============================================================================
--- libclc/trunk/amdgpu/lib/workitem/get_local_id.ll (original)
+++ libclc/trunk/amdgpu/lib/workitem/get_local_id.ll (removed)
@@ -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
-}
Modified: libclc/trunk/r600/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/SOURCES?rev=261042&r1=261041&r2=261042&view=diff
==============================================================================
--- libclc/trunk/r600/lib/SOURCES (original)
+++ libclc/trunk/r600/lib/SOURCES Tue Feb 16 18:27:27 2016
@@ -1 +1,3 @@
synchronization/barrier_impl.ll
+workitem/get_group_id.ll
+workitem/get_local_id.ll
Added: libclc/trunk/r600/lib/workitem/get_group_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_group_id.ll?rev=261042&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_group_id.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_group_id.ll Tue Feb 16 18:27:27 2016
@@ -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 }
Added: libclc/trunk/r600/lib/workitem/get_local_id.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/r600/lib/workitem/get_local_id.ll?rev=261042&view=auto
==============================================================================
--- libclc/trunk/r600/lib/workitem/get_local_id.ll (added)
+++ libclc/trunk/r600/lib/workitem/get_local_id.ll Tue Feb 16 18:27:27 2016
@@ -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 }
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
