Index: include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- include/llvm/IR/IntrinsicsAMDGPU.td +++ include/llvm/IR/IntrinsicsAMDGPU.td @@ -11,30 +11,35 @@ // //===----------------------------------------------------------------------===// -class AMDGPUReadPreloadRegisterIntrinsic - : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, - GCCBuiltin; +class AMDGPUReadPreloadRegisterIntrinsic + : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; + +class AMDGPUReadPreloadRegisterIntrinsicNamed + : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin; let TargetPrefix = "r600" in { -multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz { - def _x : AMDGPUReadPreloadRegisterIntrinsic; - def _y : AMDGPUReadPreloadRegisterIntrinsic; - def _z : AMDGPUReadPreloadRegisterIntrinsic; +multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz { + def _x : AMDGPUReadPreloadRegisterIntrinsic; + def _y : AMDGPUReadPreloadRegisterIntrinsic; + def _z : AMDGPUReadPreloadRegisterIntrinsic; +} + +multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named { + def _x : AMDGPUReadPreloadRegisterIntrinsicNamed; + def _y : AMDGPUReadPreloadRegisterIntrinsicNamed; + def _z : AMDGPUReadPreloadRegisterIntrinsicNamed; } -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 :