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/amdgpu | |
| 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/amdgpu')
| -rw-r--r-- | libclc/amdgpu/lib/SOURCES | 2 | ||||
| -rw-r--r-- | libclc/amdgpu/lib/workitem/get_group_id.ll | 18 | ||||
| -rw-r--r-- | libclc/amdgpu/lib/workitem/get_local_id.ll | 18 |
3 files changed, 0 insertions, 38 deletions
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 -} |