llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-mlir-gpu Author: Mehdi Amini (joker-eph) <details> <summary>Changes</summary> Enable strict property assembly format mode for the NVVM dialect and update custom assembly formats to expose property dictionaries explicitly. Refresh NVVM tests so inherent operation properties are printed and parsed through the property dictionary while non-property attributes remain in the attribute dictionary. Assisted-by: Codex --- Patch is 1.10 MiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/196289.diff 80 Files Affected: - (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+93-91) - (modified) mlir/test/Conversion/GPUToNVVM/gpu-to-nvvm.mlir (+4-4) - (modified) mlir/test/Conversion/GPUToNVVM/wmma-ops-to-nvvm.mlir (+14-14) - (modified) mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir (+2-2) - (modified) mlir/test/Conversion/NVVMToLLVM/invalid.mlir (+2-2) - (modified) mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir (+11-11) - (modified) mlir/test/Dialect/LLVMIR/invalid.mlir (+13-13) - (modified) mlir/test/Dialect/LLVMIR/nvvm-transcendentals.mlir (+8-8) - (modified) mlir/test/Dialect/LLVMIR/nvvm.mlir (+27-27) - (modified) mlir/test/Dialect/LLVMIR/nvvm/invalid-convert-stochastic-rounding.mlir (+3-3) - (modified) mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir (+11-11) - (modified) mlir/test/Target/LLVMIR/nvvm/addf/addf.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/addf/addf_invalid.mlir (+8-8) - (modified) mlir/test/Target/LLVMIR/nvvm/addf/addf_vector.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp4x2.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir (+8-8) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir (+26-26) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp8x2_invalid.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_s2f6x2.mlir (+10-10) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_stochastic_rounding.mlir (+7-7) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_tf32.mlir (+10-10) - (modified) mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir (+11-11) - (modified) mlir/test/Target/LLVMIR/nvvm/fence.mlir (+9-9) - (modified) mlir/test/Target/LLVMIR/nvvm/fma/fma.mlir (+32-32) - (modified) mlir/test/Target/LLVMIR/nvvm/fma/fma_invalid.mlir (+11-11) - (modified) mlir/test/Target/LLVMIR/nvvm/fma/fma_vector.mlir (+32-32) - (modified) mlir/test/Target/LLVMIR/nvvm/invalid_convert_fp16x2.mlir (+6-6) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_arr_drop_expect_tx.mlir (+15-15) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_arr_expect_tx.mlir (+15-15) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_arrive.mlir (+18-18) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_arrive_drop.mlir (+18-18) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_complete_tx.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_expect_tx.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_init.mlir (+2-2) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_invalid.mlir (+13-13) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_test_wait.mlir (+12-12) - (modified) mlir/test/Target/LLVMIR/nvvm/mbar_try_wait.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/redux-sync-invalid.mlir (+2-2) - (modified) mlir/test/Target/LLVMIR/nvvm/shfl-sync-invalid.mlir (+1-1) - (modified) mlir/test/Target/LLVMIR/nvvm/subf/subf.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/subf/subf_invalid.mlir (+8-8) - (modified) mlir/test/Target/LLVMIR/nvvm/subf/subf_vector.mlir (+24-24) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir (+34-34) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-ld-invalid.mlir (+1-1) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-ld-red-invalid.mlir (+6-6) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-ld-red.mlir (+84-84) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-ld.mlir (+74-74) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-block-scale-shared.mlir (+48-48) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-block-scale-tensor.mlir (+48-48) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-invalid.mlir (+18-18) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir (+96-96) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-sp-block-scale-shared.mlir (+48-48) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-sp-block-scale-tensor.mlir (+48-48) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-sp-shared.mlir (+96-96) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-sp-tensor.mlir (+144-144) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir (+144-144) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-ws-shared.mlir (+40-40) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-ws-sp-shared.mlir (+40-40) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-ws-sp-tensor.mlir (+40-40) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-ws-tensor.mlir (+40-40) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir (+1-1) - (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-st.mlir (+74-74) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_load_cluster_im2col.mlir (+108-108) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_load_cluster_tile.mlir (+70-70) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_load_cta_im2col.mlir (+18-18) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_load_cta_tile.mlir (+12-12) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_load_invalid.mlir (+10-10) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_prefetch.mlir (+21-21) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_prefetch_invalid.mlir (+6-6) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_store.mlir (+8-8) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_store_invalid.mlir (+3-3) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir (+128-128) - (modified) mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_invalid.mlir (+3-3) - (modified) mlir/test/Target/LLVMIR/nvvm/transcendentals.mlir (+4-4) - (modified) mlir/test/Target/LLVMIR/nvvmir-invalid.mlir (+40-40) - (modified) mlir/test/Target/LLVMIR/nvvmir.mlir (+65-65) - (modified) mlir/test/python/dialects/nvvm.py (+9-9) ``````````diff diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 51396947fad4e..fa510bbf499fd 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -110,6 +110,7 @@ def NVVM_Dialect : Dialect { }]; let useDefaultAttributePrinterParser = 1; + let useStrictPropertiesInAssemblyFormat = 1; } //===----------------------------------------------------------------------===// @@ -308,7 +309,7 @@ class NVVM_F32UnaryApproxOp<string mnemonic, list<Trait> traits = []> : let arguments = (ins F32:$src, DefaultValuedAttr<BoolAttr, "false">:$ftz); let results = (outs F32:$res); - let assemblyFormat = "$src attr-dict `:` type($src)"; + let assemblyFormat = "$src (`,` `ftz` `=` $ftz^)? attr-dict `:` type($src)"; } @@ -610,7 +611,8 @@ def NVVM_ReduxOp : $res = createIntrinsicCall(builder, intId, {$val, $mask_and_clamp}); }]; let assemblyFormat = [{ - $kind $val `,` $mask_and_clamp attr-dict `:` type($val) `->` type($res) + $kind $val `,` $mask_and_clamp (`abs` `=` $abs^)? + (`nan` `=` $nan^)? attr-dict `:` type($val) `->` type($res) }]; } @@ -755,7 +757,7 @@ def NVVM_MBarrierExpectTxOp : NVVM_VoidIntrinsicOp<"mbarrier.expect_tx"> { I32:$txcount, DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope); - let assemblyFormat = "$addr `,` $txcount attr-dict `:` type(operands)"; + let assemblyFormat = "$addr `,` $txcount prop-dict attr-dict `:` type(operands)"; let hasVerifier = 1; } @@ -778,7 +780,7 @@ def NVVM_MBarrierCompleteTxOp : NVVM_VoidIntrinsicOp<"mbarrier.complete_tx"> { I32:$txcount, DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope); - let assemblyFormat = "$addr `,` $txcount attr-dict `:` type(operands)"; + let assemblyFormat = "$addr `,` $txcount prop-dict attr-dict `:` type(operands)"; let hasVerifier = 1; } @@ -827,7 +829,7 @@ def NVVM_MBarrierArriveOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive", DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope, DefaultValuedAttr<BoolAttr, "false">:$relaxed); - let assemblyFormat = "$addr (`,` $count^)? attr-dict `:` type($addr) (`->` type($res)^)?"; + let assemblyFormat = "$addr (`,` $count^)? prop-dict attr-dict `:` type($addr) (`->` type($res)^)?"; let hasVerifier = 1; } @@ -853,7 +855,7 @@ def NVVM_MBarrierArriveDropOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_dr DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope, DefaultValuedAttr<BoolAttr, "false">:$relaxed); - let assemblyFormat = "$addr (`,` $count^)? attr-dict `:` type($addr) (`->` type($res)^)?"; + let assemblyFormat = "$addr (`,` $count^)? prop-dict attr-dict `:` type($addr) (`->` type($res)^)?"; let hasVerifier = 1; } @@ -951,7 +953,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t DefaultValuedAttr<BoolAttr, "false">:$relaxed, PtxPredicate:$predicate); - let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands) (`->` type($res)^)?"; + let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? prop-dict attr-dict `:` type(operands) (`->` type($res)^)?"; let hasVerifier = 1; let extraClassDeclaration = [{ @@ -994,7 +996,7 @@ def NVVM_MBarrierArriveDropExpectTxOp : NVVM_SingleResultIntrinsicOp<"mbarrier.a DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope, DefaultValuedAttr<BoolAttr, "false">:$relaxed); - let assemblyFormat = "$addr `,` $txcount attr-dict `:` type(operands) (`->` type($res)^)?"; + let assemblyFormat = "$addr `,` $txcount prop-dict attr-dict `:` type(operands) (`->` type($res)^)?"; let hasVerifier = 1; } @@ -1118,7 +1120,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.test.wait"> DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope, DefaultValuedAttr<BoolAttr, "false">:$relaxed); - let assemblyFormat = "$addr `,` $stateOrPhase attr-dict `:` type(operands) `->` type($res)"; + let assemblyFormat = "$addr `,` $stateOrPhase prop-dict attr-dict `:` type(operands) `->` type($res)"; let hasVerifier = 1; } @@ -1147,7 +1149,7 @@ def NVVM_MBarrierTryWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.try_wait"> { DefaultValuedAttr<MemScopeKindAttr, "MemScopeKind::CTA">:$scope, DefaultValuedAttr<BoolAttr, "false">:$relaxed); - let assemblyFormat = "$addr `,` $stateOrPhase (`,` $ticks^)? attr-dict `:` type(operands) `->` type($res)"; + let assemblyFormat = "$addr `,` $stateOrPhase (`,` $ticks^)? prop-dict attr-dict `:` type(operands) `->` type($res)"; let hasVerifier = 1; } @@ -1280,7 +1282,7 @@ def NVVM_ClusterArriveOp : NVVM_Op<"cluster.arrive"> { else createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier_cluster_arrive); }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "(`aligned` $aligned^)? attr-dict"; } def NVVM_ClusterArriveRelaxedOp : NVVM_Op<"cluster.arrive.relaxed", [NVVMRequiresSM<90>]> { @@ -1306,7 +1308,7 @@ def NVVM_ClusterArriveRelaxedOp : NVVM_Op<"cluster.arrive.relaxed", [NVVMRequire else createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier_cluster_arrive_relaxed); }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "(`aligned` $aligned^)? attr-dict"; } def NVVM_ClusterWaitOp : NVVM_Op<"cluster.wait", [NVVMRequiresSM<90>]> { @@ -1327,7 +1329,7 @@ def NVVM_ClusterWaitOp : NVVM_Op<"cluster.wait", [NVVMRequiresSM<90>]> { else createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier_cluster_wait); }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "(`aligned` $aligned^)? attr-dict"; } //===----------------------------------------------------------------------===// @@ -1374,7 +1376,7 @@ def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">, [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar) }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "$order attr-dict"; let llvmBuilder = [{ createIntrinsicCall(builder, getFenceSyncRestrictID($order)); }]; @@ -1428,7 +1430,7 @@ def NVVM_FenceProxyOp : NVVM_Op<"fence.proxy">, [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar) }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "$kind (`,` `space` `=` $space^)? attr-dict"; let llvmBuilder = [{ createIntrinsicCall(builder, getFenceProxyID($kind, $space)); @@ -1510,7 +1512,7 @@ def NVVM_FenceProxySyncRestrictOp : NVVM_Op<"fence.proxy.sync_restrict">, [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar) }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "$order (`from_proxy` `=` $fromProxy^)? (`to_proxy` `=` $toProxy^)? attr-dict"; let llvmBuilder = [{ createIntrinsicCall(builder, getFenceProxySyncRestrictID($order)); }]; @@ -1587,7 +1589,7 @@ def NVVM_ShflOp : intId, {$thread_mask, $val, $offset, $mask_and_clamp}); }]; let assemblyFormat = [{ - $kind $thread_mask `,` $val `,` $offset `,` $mask_and_clamp attr-dict + $kind $thread_mask `,` $val `,` $offset `,` $mask_and_clamp prop-dict attr-dict `:` type($val) `->` type($res) }]; let hasVerifier = 1; @@ -1634,7 +1636,7 @@ def NVVM_VoteSyncOp auto intId = getVoteSyncIntrinsicId($kind); $res = createIntrinsicCall(builder, intId, {$mask, $pred}); }]; - let assemblyFormat = "$kind $mask `,` $pred attr-dict `->` type($res)"; + let assemblyFormat = "$kind $mask `,` $pred prop-dict attr-dict `->` type($res)"; let hasVerifier = 1; } @@ -1668,7 +1670,7 @@ def NVVM_SyncWarpOp : [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar-warp-sync) }]; - let assemblyFormat = "$mask attr-dict `:` type($mask)"; + let assemblyFormat = "$mask prop-dict attr-dict `:` type($mask)"; } def NVVM_ElectSyncOp : NVVM_Op<"elect.sync"> @@ -1686,7 +1688,7 @@ def NVVM_ElectSyncOp : NVVM_Op<"elect.sync"> let arguments = (ins Optional<I32>:$membermask); let results = (outs I1:$pred); - let assemblyFormat = "($membermask^)? attr-dict `->` type(results)"; + let assemblyFormat = "($membermask^)? prop-dict attr-dict `->` type(results)"; string llvmBuilder = [{ auto *resultTuple = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_elect_sync, @@ -1805,7 +1807,7 @@ def NVVM_PermuteOp : NVVM_SingleResultIntrinsicOp<"prmt", [Pure]>, }]; let assemblyFormat = [{ - $mode $selector `,` $lo (`,` $hi^)? attr-dict `:` type($res) + $mode $selector `,` $lo (`,` $hi^)? prop-dict attr-dict `:` type($res) }]; let hasVerifier = 1; @@ -1839,7 +1841,7 @@ def NVVM_CpAsyncOp : NVVM_Op<"cp.async.shared.global">, I32Attr:$size, LoadCacheModifierAttr:$modifier, Optional<I32>:$cpSize)> { - let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` $modifier (`,` $cpSize^)? attr-dict `:` type(operands)"; + let assemblyFormat = "$dst `,` $src `,` $size `,` `cache` `=` $modifier (`,` $cpSize^)? prop-dict attr-dict `:` type(operands)"; let hasVerifier = 1; let extraClassDeclaration = [{ static llvm::Intrinsic::ID @@ -1858,7 +1860,7 @@ def NVVM_CpAsyncCommitGroupOp : NVVM_Op<"cp.async.commit.group"> { string llvmBuilder = [{ createIntrinsicCall(builder, llvm::Intrinsic::nvvm_cp_async_commit_group); }]; - let assemblyFormat = "attr-dict"; + let assemblyFormat = "prop-dict attr-dict"; } def NVVM_CpAsyncWaitGroupOp : NVVM_Op<"cp.async.wait.group">, @@ -1871,7 +1873,7 @@ def NVVM_CpAsyncWaitGroupOp : NVVM_Op<"cp.async.wait.group">, llvm::Type::getInt32Ty(moduleTranslation.getLLVMContext()), $n)); }]; - let assemblyFormat = "$n attr-dict"; + let assemblyFormat = "$n prop-dict attr-dict"; } def NVVM_CpAsyncMBarrierArriveOp : NVVM_VoidIntrinsicOp<"cp.async.mbarrier.arrive"> { @@ -1892,7 +1894,7 @@ def NVVM_CpAsyncMBarrierArriveOp : NVVM_VoidIntrinsicOp<"cp.async.mbarrier.arriv AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr, DefaultValuedAttr<I1Attr, "0">:$noinc); - let assemblyFormat = "$addr attr-dict `:` type(operands)"; + let assemblyFormat = "$addr prop-dict attr-dict `:` type(operands)"; } //===----------------------------------------------------------------------===// @@ -1971,7 +1973,7 @@ def NVVM_ConvertFloatToTF32Op : NVVM_Op<"convert.float.to.tf32"> { DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat, DefaultValuedAttr<BoolAttr, "false">:$relu); - let assemblyFormat = "$src attr-dict"; + let assemblyFormat = "$src prop-dict attr-dict"; let extraClassDeclaration = [{ static llvm::Intrinsic::ID getIntrinsicID(NVVM::FPRoundingMode, @@ -2003,7 +2005,7 @@ def NVVM_ConvertF32x2ToF4x2Op : NVVM_Op<"convert.f32x2.to.f4x2"> { let arguments = (ins F32:$a, F32:$b, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttr:$dstTy); - let assemblyFormat = "$a `,` $b attr-dict `:` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$a `,` $b prop-dict attr-dict `:` type($dst) `(` $dstTy `)`"; let hasVerifier = 1; let extraClassDeclaration = [{ @@ -2039,7 +2041,7 @@ class NVVM_ConvertFPx2ToF4x2Op<string srcType> TypeAttrOf<F4E2M1FN>:$dstTy); let assemblyFormat = - "$src attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; + "$src prop-dict attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; let extraClassDeclaration = [{ static NVVM::IDArgPair @@ -2081,7 +2083,7 @@ def NVVM_ConvertF32x2ToF6x2Op : NVVM_Op<"convert.f32x2.to.f6x2"> { F32:$b, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttr:$dstTy); - let assemblyFormat = "$a `,` $b attr-dict `:` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$a `,` $b prop-dict attr-dict `:` type($dst) `(` $dstTy `)`"; let hasVerifier = 1; let extraClassDeclaration = [{ @@ -2123,7 +2125,7 @@ class NVVM_ConvertFPx2ToF6x2Op<string srcType> VectorOfLengthAndType<[2], [!cast<Type>(srcType)]>:$src, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttrOf<AnyTypeOf<[F6E2M3FN, F6E3M2FN]>>:$dstTy); - let assemblyFormat = "$src attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$src prop-dict attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; let extraClassDeclaration = [{ static llvm::Intrinsic::ID getIntrinsicID(mlir::Type dstTy, bool hasRelu); }]; @@ -2169,7 +2171,7 @@ def NVVM_ConvertF32x2ToF8x2Op : NVVM_Op<"convert.f32x2.to.f8x2"> { DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttr:$dstTy); - let assemblyFormat = "$a `,` $b attr-dict `:` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$a `,` $b prop-dict attr-dict `:` type($dst) `(` $dstTy `)`"; let extraClassDeclaration = [{ static llvm::Intrinsic::ID getIntrinsicID(mlir::Type dstTy, @@ -2214,7 +2216,7 @@ def NVVM_ConvertF16x2ToF8x2Op : NVVM_Op<"convert.f16x2.to.f8x2"> { VectorOfLengthAndType<[2], [F16]>:$a, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttr:$dstTy); - let assemblyFormat = "$a attr-dict `:` type($a) `->` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$a prop-dict attr-dict `:` type($a) `->` type($dst) `(` $dstTy `)`"; let extraClassDeclaration = [{ static llvm::Intrinsic::ID getIntrinsicID(mlir::Type dstTy, @@ -2258,7 +2260,7 @@ def NVVM_ConvertBF16x2ToF8x2Op : NVVM_Op<"convert.bf16x2.to.f8x2"> { DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat, DefaultValuedAttr<BoolAttr, "false">:$relu, TypeAttrOf<AnyTypeOf<[F8E8M0FNU, F8E4M3FN, F8E5M2]>>:$dstTy); - let assemblyFormat = "$src attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$src prop-dict attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; let extraClassDeclaration = [{ static llvm::Intrinsic::ID getIntrinsicID(mlir::Type dstTy, @@ -2300,7 +2302,7 @@ class NVVM_ConvertToFP16x2Op_Base <string srcType, Type srcArgType, string dstTy TypeAttr:$srcType), (ins srcArgType:$src, TypeAttr:$srcType)); - let assemblyFormat = "$src attr-dict `:` type($src) `(` $srcType `)` `->` type($dst)"; + let assemblyFormat = "$src prop-dict attr-dict `:` type($src) `(` $srcType `)` `->` type($dst)"; let hasVerifier = 1; } @@ -2338,7 +2340,7 @@ def NVVM_ConvertF32x2ToS2F6x2Op : NVVM_Op<"convert.f32x2.to.s2f6x2"> { Optional<I16>:$scaleFactor, DefaultValuedAttr<BoolAttr, "false">:$relu); let assemblyFormat = - "$a `,` $b (`,` $scaleFactor^)? attr-dict `:` type($dst)"; + "$a `,` $b (`,` $scaleFactor^)? prop-dict attr-dict `:` type($dst)"; let extraClassDeclaration = [{ static IDArgPair getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder); @@ -2382,7 +2384,7 @@ def NVVM_ConvertBF16x2ToS2F6x2Op : NVVM_Op<"convert.bf16x2.to.s2f6x2"> { Optional<I16>:$scaleFactor, DefaultValuedAttr<BoolAttr, "false">:$relu); let assemblyFormat = - "$src (`,` $scaleFactor^)? attr-dict `:` type($src) `->` type($dst)"; + "$src (`,` $scaleFactor^)? prop-dict attr-dict `:` type($src) `->` type($dst)"; let extraClassDeclaration = [{ static IDArgPair getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder); @@ -2421,7 +2423,7 @@ def NVVM_ConvertS2F6x2ToBF16x2Op : NVVM_SingleResultIntrinsicOp<"convert.s2f6x2. DefaultValuedAttr<SaturationModeAttr, "SaturationMode::NONE">:$sat, DefaultValuedAttr<BoolAttr, "false">:$relu); let assemblyFormat = - "$src (`,` $scaleFactor^)? attr-dict `:` type($src) `->` type($dst)"; + "$src (`,` $scaleFactor^)? prop-dict attr-dict `:` type($src) `->` type($dst)"; } //===----------------------------------------------------------------------===// @@ -2458,7 +2460,7 @@ class NVVM_ConvertF32x2ToFPx2OpBase<string dstFormat, string mnemonic, Type dstT [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt) }]; - let assemblyFormat = "$src_hi `,` $src_lo (`,` $random_bits^)? attr-dict `:` type($dst)"; + let assemblyFormat = "$src_hi `,` $src_lo (`,` $random_bits^)? prop-dict attr-dict `:` type($dst)"; let hasVerifier = 1; @@ -2503,7 +2505,7 @@ class NVVM_ConvertF32x4ToFPx4OpBase<string dstFormat, string mnemonic, Type dstT [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt) }]; - let assemblyFormat = "$src `,` $rbits attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; + let assemblyFormat = "$src `,` $rbits prop-dict attr-dict `:` type($src) `->` type($dst) `(` $dstTy `)`"; let hasVerifier = 1; @@ -3111,7 +3113,7 @@ def NVVM_WMMALoadOp: NVVM_Op<"wmma.load">, ``` }]; - let assemblyFormat = "$ptr `,` $stride attr-dict `:` functional-type($ptr, $res)"; + let assemblyFormat = "$ptr `,` $stride prop-dict attr-dict `:` functional-type($ptr, $res)"; let hasVerifier = 1; } @@ -3167,7 +3169,7 @@ def NVVM_WMMAStoreOp : NVVM_Op<"wmma.store">, }]; let assemblyFormat = [{ - $ptr `,` $stride `,` $args attr-dict `:` qualified(type($ptr)) `,` + $ptr `,` $stride `,` $args prop-dict attr-dict `:` qualified(type($ptr)) `,` type($args) }]; let hasVerifier = 1; @@ -3223,7 +3225,7 @@ def NVVM_WMMAMmaOp : NVVM_Op<"wmma.mma">, ``` }]; - let assemblyFormat = "$args attr-dict `:` functional-type($args, $res)"; + let assemblyFormat = "$args prop-dict attr-dict `:` functional-type($args, $res)"; let hasVerifier = 1; } @@ -3263,7 +3265,7 @@ def NVVM_StMatrixOp: NVVM_Op<"stmatrix">, auto intId = getStMatrixIntrinsicId($layout, $sources.size(), $shape, $eltType); createIntrinsicCall(builder, intId, operands, operands[0]->getType()); }]; - let assemblyFormat = "$ptr `,` $sources attr-dict `:` type(operands)"; + let assemblyFormat = "$ptr `,` $sources prop-dict attr-dict `:` type(operands)"; let hasVerifier = 1; } @@ -3310,7 +3312,7 @@ def NVVM_LdMatrixOp: NVVM_Op<"ldmatrix", [InferTypeOpAdaptor]>, ``` }]; - let assemblyFormat = "$ptr attr-dict `:` functional-type($ptr, $res)"; + let assemblyFormat = "$ptr prop-dict attr-dict `:` functional-type($ptr, $res)"; let hasVerifier = 1; } @@ -3340,7 +3342,7 @@ def NVVM_MovMatrixOp DefaultValuedAttr<MMALayoutAttr, "MMALayout::col">:$layout, LdStMatrixEltTypeAttr:$eltType); - let assemblyFormat = "$src attr-dict `:` type($src)"; + let assemblyFormat = "$src prop-dict attr-dict `:` type($src)"; let hasVerifier = 1; } @@ -4209,7 +4211,7 @@ def CTAGroupKindAttr : def NVVM_CpAsyncBulkCommitGroupOp : NVVM_Op<"cp.async.bulk.commit.group">, Arguments<(ins )> { - let assemblyFormat = "attr-dict"; + let assemblyFormat = "prop-dict attr-dict"; let description = [{ This Op commits all prior initiated but uncommitted cp.async.bulk instructions into a cp.async.bulk-group. @@ -4226,7 +4228,7 @@ def NVVM_CpAsyncBulkWaitGroupOp : NVVM_Op<"cp.async.bulk.wait_group", [NVVMRequi Arguments<(ins ConfinedAttr<I32Attr, [IntMinValue<0>]>:$group, OptionalAttr<UnitAttr>:$read)> { - let assemblyFormat = "$group attr-dict"; + let assemblyFormat = "$group prop-dict attr-dict"; let description = [{ Op waits for completion of the most recent bulk async-groups. @@ -4294,7 +4296,7 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp : (`multicast_mask` `=` $multicastMask^ )? (`l2_cache_hint` `=` $l2CacheHint^ )? (`predicate` `=` $predicate^)? ... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/196289 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
