[PATCH] D20691: AMDGPU: Remove gcc builtin names from workitem intrinsics

[Jan Vesely via llvm-commits]

jvesely created this revision.
jvesely added reviewers: tstellarAMD, arsenm.
jvesely added a subscriber: llvm-commits.
jvesely set the repository for this revision to rL LLVM.
Herald added a subscriber: kzhuravl.

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

Names in llvm/IR/Intrinsics*.td are considered before custom emit so they need to go.

Repository:
  rL LLVM

http://reviews.llvm.org/D20691

Files:
  include/llvm/IR/IntrinsicsAMDGPU.td

Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -11,30 +11,25 @@
 //
 //===----------------------------------------------------------------------===//
 
-class AMDGPUReadPreloadRegisterIntrinsic<string name>
-  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
-    GCCBuiltin<name>;
+class AMDGPUReadPreloadRegisterIntrinsic
+  : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
 
 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;
 }
 
-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_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_global_offset">;
+defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+defm int_r600_read_global_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz;
+
+def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
 
 def int_r600_rat_store_typed :
   // 1st parameter: Data
@@ -47,9 +42,6 @@
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 >;
 
-def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_r600_read_workdim"
->;
 
 } // End TargetPrefix = "r600"
 
@@ -62,10 +54,8 @@
 
 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;
 
 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
   Intrinsic<[], [], [IntrConvergent]>;
@@ -279,8 +269,7 @@
    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 :


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D20691.58652.patch
Type: text/x-patch
Size: 3497 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160526/66cb58b2/attachment.bin>