diff options
author | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 17:24:20 +0000 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 17:24:20 +0000 |
commit | ff45572d5e10a96309d58593b967a71228322f8b (patch) | |
tree | 3f0479e6a400a4c38b9384be7e9a5dd5cac86179 | |
parent | 3d39fb557be2e349ded2d2055211b2757627f40b (diff) |
AMDGPU: Use clang intrinsics for workitem builtins
v2: split into 2 patches
use clang builtins for other intrinsics as well
v3: Fix warnings
Switch r600 to use implictarg.ptr
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@276442 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r-- | amdgcn/lib/SOURCES | 5 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_group_id.cl | 11 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_group_id.ll | 29 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_local_id.cl | 11 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_local_id.ll | 31 | ||||
-rw-r--r-- | amdgcn/lib/workitem/get_work_dim.cl | 9 | ||||
-rw-r--r-- | amdgpu/lib/SOURCES | 7 | ||||
-rw-r--r-- | amdgpu/lib/workitem/get_work_dim.ll | 8 | ||||
-rw-r--r-- | r600/lib/SOURCES | 5 | ||||
-rw-r--r-- | r600/lib/workitem/get_group_id.cl | 11 | ||||
-rw-r--r-- | r600/lib/workitem/get_group_id.ll | 29 | ||||
-rw-r--r-- | r600/lib/workitem/get_local_id.cl | 11 | ||||
-rw-r--r-- | r600/lib/workitem/get_local_id.ll | 31 | ||||
-rw-r--r-- | r600/lib/workitem/get_work_dim.cl | 9 |
14 files changed, 71 insertions, 136 deletions
diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES index ada06d2..49d9b53 100644 --- a/amdgcn/lib/SOURCES +++ b/amdgcn/lib/SOURCES @@ -1,4 +1,5 @@ math/ldexp.cl synchronization/barrier_impl.ll -workitem/get_group_id.ll -workitem/get_local_id.ll +workitem/get_group_id.cl +workitem/get_local_id.cl +workitem/get_work_dim.cl diff --git a/amdgcn/lib/workitem/get_group_id.cl b/amdgcn/lib/workitem/get_group_id.cl new file mode 100644 index 0000000..4b4e7a7 --- /dev/null +++ b/amdgcn/lib/workitem/get_group_id.cl @@ -0,0 +1,11 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_group_id(uint dim) +{ + switch(dim) { + case 0: return __builtin_amdgcn_workgroup_id_x(); + case 1: return __builtin_amdgcn_workgroup_id_y(); + case 2: return __builtin_amdgcn_workgroup_id_z(); + default: return 1; + } +} diff --git a/amdgcn/lib/workitem/get_group_id.ll b/amdgcn/lib/workitem/get_group_id.ll deleted file mode 100644 index 9d820e0..0000000 --- a/amdgcn/lib/workitem/get_group_id.ll +++ /dev/null @@ -1,29 +0,0 @@ -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.cl b/amdgcn/lib/workitem/get_local_id.cl new file mode 100644 index 0000000..257c30f --- /dev/null +++ b/amdgcn/lib/workitem/get_local_id.cl @@ -0,0 +1,11 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_local_id(uint dim) +{ + switch(dim) { + case 0: return __builtin_amdgcn_workitem_id_x(); + case 1: return __builtin_amdgcn_workitem_id_y(); + case 2: return __builtin_amdgcn_workitem_id_z(); + default: return 1; + } +} diff --git a/amdgcn/lib/workitem/get_local_id.ll b/amdgcn/lib/workitem/get_local_id.ll deleted file mode 100644 index c54291c..0000000 --- a/amdgcn/lib/workitem/get_local_id.ll +++ /dev/null @@ -1,31 +0,0 @@ -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/amdgcn/lib/workitem/get_work_dim.cl b/amdgcn/lib/workitem/get_work_dim.cl new file mode 100644 index 0000000..dd2c64f --- /dev/null +++ b/amdgcn/lib/workitem/get_work_dim.cl @@ -0,0 +1,9 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_work_dim() +{ + __attribute__((address_space(2))) uint * ptr = + (__attribute__((address_space(2))) uint *) + __builtin_amdgcn_implicitarg_ptr(); + return ptr[0]; +} diff --git a/amdgpu/lib/SOURCES b/amdgpu/lib/SOURCES index 39287bf..403e1e7 100644 --- a/amdgpu/lib/SOURCES +++ b/amdgpu/lib/SOURCES @@ -1,10 +1,6 @@ atomic/atomic.cl math/nextafter.cl math/sqrt.cl -workitem/get_num_groups.ll -workitem/get_local_size.ll -workitem/get_global_size.ll -workitem/get_work_dim.ll synchronization/barrier.cl image/get_image_width.cl image/get_image_height.cl @@ -20,3 +16,6 @@ image/write_imagef.cl image/write_imagei.cl image/write_imageui.cl image/write_image_impl.ll +workitem/get_num_groups.ll +workitem/get_local_size.ll +workitem/get_global_size.ll diff --git a/amdgpu/lib/workitem/get_work_dim.ll b/amdgpu/lib/workitem/get_work_dim.ll deleted file mode 100644 index 1f86b5e..0000000 --- a/amdgpu/lib/workitem/get_work_dim.ll +++ /dev/null @@ -1,8 +0,0 @@ -declare i32 @llvm.AMDGPU.read.workdim() nounwind readnone - -define i32 @get_work_dim() nounwind readnone alwaysinline { - %x = call i32 @llvm.AMDGPU.read.workdim() nounwind readnone , !range !0 - ret i32 %x -} - -!0 = !{ i32 1, i32 4 } diff --git a/r600/lib/SOURCES b/r600/lib/SOURCES index 49c8dd5..4178d70 100644 --- a/r600/lib/SOURCES +++ b/r600/lib/SOURCES @@ -1,3 +1,4 @@ synchronization/barrier_impl.ll -workitem/get_group_id.ll -workitem/get_local_id.ll +workitem/get_group_id.cl +workitem/get_local_id.cl +workitem/get_work_dim.cl diff --git a/r600/lib/workitem/get_group_id.cl b/r600/lib/workitem/get_group_id.cl new file mode 100644 index 0000000..e5efc0a --- /dev/null +++ b/r600/lib/workitem/get_group_id.cl @@ -0,0 +1,11 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_group_id(uint dim) +{ + switch(dim) { + case 0: return __builtin_r600_read_tgid_x(); + case 1: return __builtin_r600_read_tgid_y(); + case 2: return __builtin_r600_read_tgid_z(); + default: return 1; + } +} diff --git a/r600/lib/workitem/get_group_id.ll b/r600/lib/workitem/get_group_id.ll deleted file mode 100644 index 837c799..0000000 --- a/r600/lib/workitem/get_group_id.ll +++ /dev/null @@ -1,29 +0,0 @@ -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/r600/lib/workitem/get_local_id.cl b/r600/lib/workitem/get_local_id.cl new file mode 100644 index 0000000..a871a5d --- /dev/null +++ b/r600/lib/workitem/get_local_id.cl @@ -0,0 +1,11 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_local_id(uint dim) +{ + switch(dim) { + case 0: return __builtin_r600_read_tidig_x(); + case 1: return __builtin_r600_read_tidig_y(); + case 2: return __builtin_r600_read_tidig_z(); + default: return 1; + } +} diff --git a/r600/lib/workitem/get_local_id.ll b/r600/lib/workitem/get_local_id.ll deleted file mode 100644 index da37ca0..0000000 --- a/r600/lib/workitem/get_local_id.ll +++ /dev/null @@ -1,31 +0,0 @@ -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 } diff --git a/r600/lib/workitem/get_work_dim.cl b/r600/lib/workitem/get_work_dim.cl new file mode 100644 index 0000000..826a655 --- /dev/null +++ b/r600/lib/workitem/get_work_dim.cl @@ -0,0 +1,9 @@ +#include <clc/clc.h> + +_CLC_DEF uint get_work_dim() +{ + __attribute__((address_space(7))) uint * ptr = + (__attribute__((address_space(7))) uint *) + __builtin_r600_implicitarg_ptr(); + return ptr[0]; +} |