llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Changpeng Fang (changpeng) <details> <summary>Changes</summary> --- Full diff: https://github.com/llvm/llvm-project/pull/225123.diff 3 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+2-2) - (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+9-3) - (modified) llvm/lib/Target/AMDGPU/MIMGInstructions.td (+3-3) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 533ebf4ae69aa..4fd604390d3ce 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -1014,11 +1014,11 @@ def __builtin_amdgcn_global_store_async_from_lds_b128 : AMDGPUBuiltin<"void(_Ext def __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64 : AMDGPUBuiltin<"void(long int address_space<3> *)", [Const], "gfx1250-insts">; def __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64 : AMDGPUBuiltin<"long int(long int address_space<3> *, long int)", [Const], "gfx1250-insts">; -def __builtin_amdgcn_tensor_load_to_lds : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, _ExtVector<8, int>, _ExtVector<4, int>, _ExtVector<4, int>, _ExtVector<8, int>, _Constant int)", [Const], "gfx1250-insts"> { +def __builtin_amdgcn_tensor_load_to_lds : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, _ExtVector<8, int>, _ExtVector<4, int>, _ExtVector<4, int>, _ExtVector<8, int>, _Constant int)", [Const], "tdm-insts"> { let Documentation = [DocTensorLoadToLDS_GFX1250]; let ArgNames = ["D0", "D1", "D2", "D3", "D4", "cpol"]; } -def __builtin_amdgcn_tensor_store_from_lds : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, _ExtVector<8, int>, _ExtVector<4, int>, _ExtVector<4, int>, _ExtVector<8, int>, _Constant int)", [Const], "gfx1250-insts"> { +def __builtin_amdgcn_tensor_store_from_lds : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int>, _ExtVector<8, int>, _ExtVector<4, int>, _ExtVector<4, int>, _ExtVector<8, int>, _Constant int)", [Const], "tdm-insts"> { let Documentation = [DocTensorStoreFromLDS_GFX1250]; let ArgNames = ["D0", "D1", "D2", "D3", "D4", "cpol"]; } diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index fb4d242dfd7a7..99b75c0681b2a 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -1319,6 +1319,10 @@ defm TensorCvtLutInsts : AMDGPUSubtargetFeature<"tensor-cvt-lut-insts", "Has v_perm_pk16* instructions" >; +defm TDMInsts : AMDGPUSubtargetFeature<"tdm-insts", + "Has tensor_load_to_lds/tensor_store_from_lds instructions" +>; + defm TransposeLoadF4F6Insts : AMDGPUSubtargetFeature<"transpose-load-f4f6-insts", "Has ds_load_tr4/tr6 and global_load_tr4/tr6 instructions" >; @@ -2439,6 +2443,7 @@ def FeatureISAVersion12_50_Common : FeatureSet< FeatureBitOp3Insts, FeatureTanhInsts, FeatureTensorCvtLutInsts, + FeatureTDMInsts, FeatureTransposeLoadF4F6Insts, FeatureBF16TransInsts, FeatureBF16ConversionInsts, @@ -3290,9 +3295,10 @@ def AMDGPUFrontendVisibleFeatures { FeatureQsadInsts, FeatureSMemRealTime, FeatureSMemTimeInst, FeatureSWMMACGfx1200Insts, FeatureSWMMACGfx1250Insts, FeatureSWakeupBarrier, FeatureSadInsts, FeatureSetPrioIncWgInst, FeatureSmemPrefetchInsts, - FeatureTanhInsts, FeatureTensorCvtLutInsts, FeatureTransposeLoadF4F6Insts, - FeatureVMemToLDSLoad, FeatureVmemPrefInsts, FeatureWMMA128bInsts, - FeatureWMMA256bInsts, FeatureWMMAN16Insts, FeatureWMMAF4Insts, FeatureXF32Insts, + FeatureTDMInsts, FeatureTanhInsts, FeatureTensorCvtLutInsts, + FeatureTransposeLoadF4F6Insts, FeatureVMemToLDSLoad, FeatureVmemPrefInsts, + FeatureWMMA128bInsts, FeatureWMMA256bInsts, FeatureWMMAN16Insts, + FeatureWMMAF4Insts, FeatureXF32Insts, FeatureWavefrontSize32, FeatureWavefrontSize64, FeatureSupportsWGP, FeatureSupportsWave32, FeatureFastFMAF32, FeatureFastDenormalF32, FeatureSupportsXNACK, FeatureSupportsSRAMECC, FeatureXNACKOnOffModes, diff --git a/llvm/lib/Target/AMDGPU/MIMGInstructions.td b/llvm/lib/Target/AMDGPU/MIMGInstructions.td index 8126135e7a353..a4a4945e89290 100644 --- a/llvm/lib/Target/AMDGPU/MIMGInstructions.td +++ b/llvm/lib/Target/AMDGPU/MIMGInstructions.td @@ -2309,12 +2309,12 @@ class VIMAGE_TENSOR_Pseudo<string opName, bit _UpTo2D = 0> : string AsmOperands = " $vaddr0, $vaddr1"#!if(UpTo2D, "", ", $vaddr2, $vaddr3")#"$r128$cpol"; } -let SubtargetPredicate = isGFX125xOnly in { +let SubtargetPredicate = HasTDMInsts in { def TENSOR_LOAD_TO_LDS_d4 : VIMAGE_TENSOR_Pseudo<"tensor_load_to_lds">; def TENSOR_STORE_FROM_LDS_d4 : VIMAGE_TENSOR_Pseudo<"tensor_store_from_lds">; def TENSOR_LOAD_TO_LDS_d2 : VIMAGE_TENSOR_Pseudo<"tensor_load_to_lds", 1>; def TENSOR_STORE_FROM_LDS_d2 : VIMAGE_TENSOR_Pseudo<"tensor_store_from_lds", 1>; -} // End SubtargetPredicate = isGFX125xOnly. +} // End SubtargetPredicate = HasTDMInsts. class VIMAGE_TENSOR_Real <bits<8> op, VIMAGE_TENSOR_Pseudo ps, string opName = ps.Mnemonic> : InstSI <ps.OutOperandList, ps.InOperandList, opName # ps.AsmOperands, []>, @@ -2346,7 +2346,7 @@ class VIMAGE_TENSOR_Real <bits<8> op, VIMAGE_TENSOR_Pseudo ps, string opName = p } multiclass VIMAGE_TENSOR_Real_gfx1250<bits<8> op> { - let AssemblerPredicate = isGFX125xOnly, DecoderNamespace = "GFX1250" in { + let AssemblerPredicate = HasTDMInsts, DecoderNamespace = "GFX1250" in { foreach DSuffix = ["_d2", "_d4"] in { defvar ps = !cast<VIMAGE_TENSOR_Pseudo>(NAME # DSuffix); def DSuffix # _gfx1250 : VIMAGE_TENSOR_Real<op, ps, ps.Mnemonic>, `````````` </details> https://github.com/llvm/llvm-project/pull/225123 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
