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 | |
| 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')
| -rw-r--r-- | amdgcn/lib/SOURCES | 2 | ||||
| -rw-r--r-- | amdgcn/lib/workitem/get_group_id.ll | 29 | ||||
| -rw-r--r-- | amdgcn/lib/workitem/get_local_id.ll | 31 |
3 files changed, 62 insertions, 0 deletions
diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES index c99f3fc..49c8dd5 100644 --- a/amdgcn/lib/SOURCES +++ b/amdgcn/lib/SOURCES @@ -1 +1,3 @@ synchronization/barrier_impl.ll +workitem/get_group_id.ll +workitem/get_local_id.ll diff --git a/amdgcn/lib/workitem/get_group_id.ll b/amdgcn/lib/workitem/get_group_id.ll new file mode 100644 index 0000000..9d820e0 --- /dev/null +++ b/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/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 } |