Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -11,30 +11,35 @@
 //
 //===----------------------------------------------------------------------===//
 
-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;
+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_named
+                                      <"__builtin_r600_read_global_offset">;
+
+def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
 
 def int_r600_rat_store_typed :
   // 1st parameter: Data
@@ -47,9 +52,6 @@
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 >;
 
-def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
-  "__builtin_r600_read_workdim"
->;
 
 } // End TargetPrefix = "r600"
 
@@ -62,10 +64,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 +279,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 :