From 26c6bb103f3878b8d5e809717c77d566be7b35dd Mon Sep 17 00:00:00 2001 From: Jan Vesely Date: Tue, 21 Jun 2016 20:46:22 +0000 Subject: [PATCH] AMDGPU: Remove gcc builtin names from workitem intrinsics We'll need to emit these manually in clang to add range metadata Reviewers: arsenm Differential Revision: http://reviews.llvm.org/D20691 llvm-svn: 273318 --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 56 +++++++++++++----------- 1 file changed, 30 insertions(+), 26 deletions(-) diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 18008c6787d9..48e215e232f1 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -11,28 +11,37 @@ // //===----------------------------------------------------------------------===// -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; } -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">; +multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named { + def _x : AMDGPUReadPreloadRegisterIntrinsicNamed; + def _y : AMDGPUReadPreloadRegisterIntrinsicNamed; + def _z : AMDGPUReadPreloadRegisterIntrinsicNamed; +} + +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 @@ def int_r600_rsq : Intrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] >; -def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic < - "__builtin_r600_read_workdim" ->; } // End TargetPrefix = "r600" @@ -60,10 +66,9 @@ def int_AMDGPU_ldexp : Intrinsic< 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 @@ def int_amdgcn_buffer_atomic_cmpswap : Intrinsic< 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 :