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 | 45e6eaaa0545b017d03f71373c983e0e7d9eac4f (patch) | |
| tree | 1e26cea8e7032f417982237b8a22b3146da215d2 /libclc/amdgcn | |
| parent | e5944d97d87b6a71022e247379b11caebcc8e521 (diff) | |
| download | llvm-45e6eaaa0545b017d03f71373c983e0e7d9eac4f.zip llvm-45e6eaaa0545b017d03f71373c983e0e7d9eac4f.tar.gz llvm-45e6eaaa0545b017d03f71373c983e0e7d9eac4f.tar.bz2 | |
amdgcn: Use new workitem intrinsics
llvm-svn: 261042
Diffstat (limited to 'libclc/amdgcn')
| -rw-r--r-- | libclc/amdgcn/lib/SOURCES | 2 | ||||
| -rw-r--r-- | libclc/amdgcn/lib/workitem/get_group_id.ll | 29 | ||||
| -rw-r--r-- | libclc/amdgcn/lib/workitem/get_local_id.ll | 31 |
3 files changed, 62 insertions, 0 deletions
diff --git a/libclc/amdgcn/lib/SOURCES b/libclc/amdgcn/lib/SOURCES index c99f3fc..49c8dd5 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 0000000..9d820e0 --- /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 0000000..c54291c --- /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 } |