From 45e6eaaa0545b017d03f71373c983e0e7d9eac4f Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 17 Feb 2016 00:27:27 +0000 Subject: amdgcn: Use new workitem intrinsics llvm-svn: 261042 --- libclc/r600/lib/SOURCES | 2 ++ libclc/r600/lib/workitem/get_group_id.ll | 29 +++++++++++++++++++++++++++++ libclc/r600/lib/workitem/get_local_id.ll | 31 +++++++++++++++++++++++++++++++ 3 files changed, 62 insertions(+) create mode 100644 libclc/r600/lib/workitem/get_group_id.ll create mode 100644 libclc/r600/lib/workitem/get_local_id.ll (limited to 'libclc/r600/lib') diff --git a/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES index c99f3fc7052..49c8dd53a56 100644 --- a/libclc/r600/lib/SOURCES +++ b/libclc/r600/lib/SOURCES @@ -1 +1,3 @@ synchronization/barrier_impl.ll +workitem/get_group_id.ll +workitem/get_local_id.ll diff --git a/libclc/r600/lib/workitem/get_group_id.ll b/libclc/r600/lib/workitem/get_group_id.ll new file mode 100644 index 00000000000..837c799395d --- /dev/null +++ b/libclc/r600/lib/workitem/get_group_id.ll @@ -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 } diff --git a/libclc/r600/lib/workitem/get_local_id.ll b/libclc/r600/lib/workitem/get_local_id.ll new file mode 100644 index 00000000000..da37ca0c7b4 --- /dev/null +++ b/libclc/r600/lib/workitem/get_local_id.ll @@ -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 } -- cgit v1.2.1