diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-02-17 00:27:27 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-02-17 00:27:27 +0000 |
| commit | ba9858caa1e927a6fcc601e3466faa693835db5e (patch) | |
| tree | 8207e6f195cc00841c6403e2baaf7d3a526e7348 /amdgcn/lib/workitem/get_local_id.ll | |
| parent | 03d52f1f52d8b7b4f738f62c459926e2796a9aab (diff) | |
amdgcn: Use new workitem intrinsics
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@261042 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'amdgcn/lib/workitem/get_local_id.ll')
| -rw-r--r-- | amdgcn/lib/workitem/get_local_id.ll | 31 |
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 } |