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

[Jan Vesely via llvm-commits]

This revision was automatically updated to reflect the committed changes.
jvesely marked an inline comment as done.
Closed by commit rL273318: AMDGPU: Remove gcc builtin names from workitem intrinsics (authored by jvesely).

Changed prior to commit:
  http://reviews.llvm.org/D20691?vs=61431&id=61440#toc

Repository:
  rL LLVM

http://reviews.llvm.org/D20691

Files:
  llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

Index: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
+++ llvm/trunk/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 @@
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 >;
 
-def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_r600_read_workdim"
->;
 
 } // End TargetPrefix = "r600"
 
@@ -60,10 +66,9 @@
 
 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 @@
    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.61440.patch
Type: text/x-patch
Size: 4023 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160621/eb420a8a/attachment.bin>