From 45e6eaaa0545b017d03f71373c983e0e7d9eac4f Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 17 Feb 2016 00:27:27 +0000 Subject: [PATCH] amdgcn: Use new workitem intrinsics llvm-svn: 261042 --- libclc/amdgcn/lib/SOURCES | 2 ++ libclc/amdgcn/lib/workitem/get_group_id.ll | 29 ++++++++++++++++++++++++++++ libclc/amdgcn/lib/workitem/get_local_id.ll | 31 ++++++++++++++++++++++++++++++ libclc/amdgpu/lib/SOURCES | 2 -- libclc/amdgpu/lib/workitem/get_group_id.ll | 18 ----------------- libclc/amdgpu/lib/workitem/get_local_id.ll | 18 ----------------- libclc/r600/lib/SOURCES | 2 ++ libclc/r600/lib/workitem/get_group_id.ll | 29 ++++++++++++++++++++++++++++ libclc/r600/lib/workitem/get_local_id.ll | 31 ++++++++++++++++++++++++++++++ 9 files changed, 124 insertions(+), 38 deletions(-) create mode 100644 libclc/amdgcn/lib/workitem/get_group_id.ll create mode 100644 libclc/amdgcn/lib/workitem/get_local_id.ll delete mode 100644 libclc/amdgpu/lib/workitem/get_group_id.ll delete mode 100644 libclc/amdgpu/lib/workitem/get_local_id.ll create mode 100644 libclc/r600/lib/workitem/get_group_id.ll create mode 100644 libclc/r600/lib/workitem/get_local_id.ll 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 } diff --git a/libclc/amdgpu/lib/SOURCES b/libclc/amdgpu/lib/SOURCES index 7505f3f..0f99fe1 100644 --- a/libclc/amdgpu/lib/SOURCES +++ b/libclc/amdgpu/lib/SOURCES @@ -3,9 +3,7 @@ math/ldexp.cl math/nextafter.cl math/sqrt.cl workitem/get_num_groups.ll -workitem/get_group_id.ll workitem/get_local_size.ll -workitem/get_local_id.ll workitem/get_global_size.ll workitem/get_work_dim.ll synchronization/barrier.cl diff --git a/libclc/amdgpu/lib/workitem/get_group_id.ll b/libclc/amdgpu/lib/workitem/get_group_id.ll deleted file mode 100644 index 0dc86e5..0000000 --- a/libclc/amdgpu/lib/workitem/get_group_id.ll +++ /dev/null @@ -1,18 +0,0 @@ -declare i32 @llvm.r600.read.tgid.x() nounwind readnone -declare i32 @llvm.r600.read.tgid.y() nounwind readnone -declare i32 @llvm.r600.read.tgid.z() nounwind readnone - -define i32 @get_group_id(i32 %dim) nounwind readnone alwaysinline { - switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] -x_dim: - %x = call i32 @llvm.r600.read.tgid.x() nounwind readnone - ret i32 %x -y_dim: - %y = call i32 @llvm.r600.read.tgid.y() nounwind readnone - ret i32 %y -z_dim: - %z = call i32 @llvm.r600.read.tgid.z() nounwind readnone - ret i32 %z -default: - ret i32 0 -} diff --git a/libclc/amdgpu/lib/workitem/get_local_id.ll b/libclc/amdgpu/lib/workitem/get_local_id.ll deleted file mode 100644 index ac5522a..0000000 --- a/libclc/amdgpu/lib/workitem/get_local_id.ll +++ /dev/null @@ -1,18 +0,0 @@ -declare i32 @llvm.r600.read.tidig.x() nounwind readnone -declare i32 @llvm.r600.read.tidig.y() nounwind readnone -declare i32 @llvm.r600.read.tidig.z() nounwind readnone - -define i32 @get_local_id(i32 %dim) nounwind readnone alwaysinline { - switch i32 %dim, label %default [i32 0, label %x_dim i32 1, label %y_dim i32 2, label %z_dim] -x_dim: - %x = call i32 @llvm.r600.read.tidig.x() nounwind readnone - ret i32 %x -y_dim: - %y = call i32 @llvm.r600.read.tidig.y() nounwind readnone - ret i32 %y -z_dim: - %z = call i32 @llvm.r600.read.tidig.z() nounwind readnone - ret i32 %z -default: - ret i32 0 -} diff --git a/libclc/r600/lib/SOURCES b/libclc/r600/lib/SOURCES index c99f3fc..49c8dd5 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 0000000..837c799 --- /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 0000000..da37ca0 --- /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 } -- 2.7.4