From 45e6eaaa0545b017d03f71373c983e0e7d9eac4f Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Wed, 17 Feb 2016 00:27:27 +0000
Subject: amdgcn: Use new workitem intrinsics

llvm-svn: 261042
---
 libclc/amdgpu/lib/SOURCES                  |  2 --
 libclc/amdgpu/lib/workitem/get_group_id.ll | 18 ------------------
 libclc/amdgpu/lib/workitem/get_local_id.ll | 18 ------------------
 3 files changed, 38 deletions(-)
 delete mode 100644 libclc/amdgpu/lib/workitem/get_group_id.ll
 delete mode 100644 libclc/amdgpu/lib/workitem/get_local_id.ll

(limited to 'libclc/amdgpu')

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
-}
-- 
cgit v1.1