summaryrefslogtreecommitdiff
path: root/amdgcn/lib/workitem/get_local_id.ll
diff options
context:
space:
mode:
Diffstat (limited to 'amdgcn/lib/workitem/get_local_id.ll')
-rw-r--r--amdgcn/lib/workitem/get_local_id.ll31
1 files changed, 31 insertions, 0 deletions
diff --git a/amdgcn/lib/workitem/get_local_id.ll b/amdgcn/lib/workitem/get_local_id.ll
new file mode 100644
index 0000000..c54291c
--- /dev/null
+++ b/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 }