AMDGPU: Remove gcc builtin names from workitem intrinsics

[jvesely]

Summary

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.

[arsenm]

Why do you need this? I know I had a patch to do the same thing at some point, but don't remember why. Since these don't need mangling I would expect GCCBuiltin to work correctly

[jvesely]

In D20691#441248, @arsenm wrote:

Why do you need this? I know I had a patch to do the same thing at some point, but don't remember why. Since these don't need mangling I would expect GCCBuiltin to work correctly

clang in CodeGenFunction::EmitBuiltinExpr() considers embedded intrinsic names before calling EmitTargetBuiltinExpr(). Since we want these intrinsics to be emitted with range metadata we need to do it in EmitAMDGPUBuiltinExpr(called by EmitTargetBUiltinExpr). The embedded names need to be removed, otherwise the Emit* functions won't run.

[jvesely]

global offset is unbounded so we can provide the GCC builtin here

[jvesely]

ngroups, global size, and group id are also unbounded.

sorry for the noise

[arsenm]

LGTM

[jvesely]

rebase to remove dependence on D20298

[arsenm]

LGTM

include/llvm/IR/IntrinsicsAMDGPU.td

40

This comment isn't needed / is misleading since the intrinsic itself has no indication of the range limit

[jvesely]

include/llvm/IR/IntrinsicsAMDGPU.td

40

Removed locally. IMO it would be preferable if the range metadata were tied to the intrinsic rather than added by clang.

include/llvm/IR/IntrinsicsAMDGPU.td

 //===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines all of the R600-specific intrinsics.
 //
 //===----------------------------------------------------------------------===//
 
-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_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz <
-                                       "__builtin_r600_read_global_offset">;
+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_global_offset : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
+                                   <"__builtin_r600_read_global_offset">;
+
+// Range limited
+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
   // 2nd parameter: Index
   // 3rd parameter: Constant RAT ID
   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
   GCCBuiltin<"__builtin_r600_rat_store_typed">;
 
 def int_r600_rsq : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 >;
 
-def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_r600_read_workdim"
->;
 
 } // End TargetPrefix = "r600"
 
 // FIXME: These should be renamed/moved to r600
 let TargetPrefix = "AMDGPU" in {
 def int_AMDGPU_ldexp : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 >;
 }
 
 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]>;
 
 def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>;
 
 def int_amdgcn_div_scale : Intrinsic<
   // 1st parameter: Numerator
def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
   [llvm_i32_ty,       // src(VGPR)
    llvm_i32_ty,       // cmp(VGPR)
    llvm_v4i32_ty,     // rsrc(SGPR)
    llvm_i32_ty,       // vindex(VGPR)
    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
    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 :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_buffer_wbinvl1 :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,