AMDGPU: Remove gcc builtin names from workitem intrinsics

We'll need to emit these manually in clang to add range metadata

Reviewers: arsenm

Differential Revision: http://reviews.llvm.org/D20691

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@273318 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/llvm/IR/IntrinsicsAMDGPU.td b/include/llvm/IR/IntrinsicsAMDGPU.td
index 18008c6787d99b5030df45db159be6c381af1578..48e215e232f1a55c762e2a1fb73d56c345fe4606 100644 (file)
--- a/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -11,28 +11,37 @@
 //
 //===----------------------------------------------------------------------===//
 
-class AMDGPUReadPreloadRegisterIntrinsic<string name>
-  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
-    GCCBuiltin<name>;
+class AMDGPUReadPreloadRegisterIntrinsic
+  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
+
+class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
+  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<name>;
 
 let TargetPrefix = "r600" in {
 
-multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz<string prefix> {
-  def _x : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>;
-  def _y : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>;
-  def _z : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>;
+multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
+  def _x : AMDGPUReadPreloadRegisterIntrinsic;
+  def _y : AMDGPUReadPreloadRegisterIntrinsic;
+  def _z : AMDGPUReadPreloadRegisterIntrinsic;
+}
+
+multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {
+  def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;
+  def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;
+  def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;
 }
 
-defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_global_size">;
-defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_local_size">;
-defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_ngroups">;
-defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_tgid">;
-defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_tidig">;
+defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
+                                 <"__builtin_r600_read_global_size">;
+defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
+                             <"__builtin_r600_read_ngroups">;
+defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
+                          <"__builtin_r600_read_tgid">;
+
+defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+
+def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
 
 def int_r600_rat_store_typed :
   // 1st parameter: Data
@@ -45,9 +54,6 @@ def int_r600_rsq : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 >;
 
-def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_r600_read_workdim"
->;
 
 } // End TargetPrefix = "r600"
 
@@ -60,10 +66,9 @@ def int_AMDGPU_ldexp : Intrinsic<
 
 let TargetPrefix = "amdgcn" in {
 
-defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-  "__builtin_amdgcn_workitem_id">;
-defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-  "__builtin_amdgcn_workgroup_id">;
+defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
+                               <"__builtin_amdgcn_workgroup_id">;
 
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
   Intrinsic<[], [], [IntrConvergent]>;
@@ -289,8 +294,7 @@ def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
    llvm_i1_ty],       // slc(imm)
   []>;
 
-def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_amdgcn_read_workdim">;
+def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
 
 
 def int_amdgcn_buffer_wbinvl1_sc :