34#include "llvm/ADT/STLExtras.h"
35#include "llvm/ADT/StringExtras.h"
36#include "llvm/ADT/TypeSwitch.h"
37#include "llvm/IR/IRBuilder.h"
38#include "llvm/IR/NVVMIntrinsicUtils.h"
39#include "llvm/Support/Casting.h"
40#include "llvm/Support/FormatVariadic.h"
41#include "llvm/Support/MathExtras.h"
42#include "llvm/Support/NVPTXAddrSpace.h"
43#include "llvm/Support/raw_ostream.h"
54#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.cpp.inc"
56static constexpr unsigned notIntrinsic = llvm::Intrinsic::not_intrinsic;
63 auto ptrTy = llvm::cast<LLVM::LLVMPointerType>(
ptr.getType());
64 return ptrTy.getAddressSpace() ==
static_cast<unsigned>(targetAS);
81 NVVMMemorySpace targetAS) {
82 unsigned AS =
static_cast<unsigned>(targetAS);
83 return builder.CreateAddrSpaceCast(
84 ptr, llvm::PointerType::get(builder.getContext(), AS));
88static llvm::nvvm::CTAGroupKind
91 case NVVM::CTAGroupKind::CTA_1:
92 return llvm::nvvm::CTAGroupKind::CG_1;
93 case NVVM::CTAGroupKind::CTA_2:
94 return llvm::nvvm::CTAGroupKind::CG_2;
96 llvm_unreachable(
"unsupported cta_group value");
100 NVVM::CTAGroupKindAttr &groupAttr) {
104 std::optional<NVVM::CTAGroupKind> group =
105 NVVM::symbolizeCTAGroupKind(keyword);
108 groupAttr = NVVM::CTAGroupKindAttr::get(parser.
getContext(), *group);
113 NVVM::CTAGroupKindAttr groupAttr) {
114 printer << NVVM::stringifyCTAGroupKind(groupAttr.getValue());
117template <
typename AttrTy>
124 using EnumTy =
decltype(attr.getValue());
125 std::optional<EnumTy> value = NVVM::symbolizeEnum<EnumTy>(keyword);
127 return parser.
emitError(loc) <<
"unknown enum value '" << keyword <<
"'";
129 attr = AttrTy::get(parser.
getContext(), *value);
142 size_t numIm2ColOffsets,
144 if (tensorDims < 1 || tensorDims > 5)
145 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
153 "to use im2col mode, the tensor has to be at least 3-dimensional");
155 if (numIm2ColOffsets && (tensorDims != (numIm2ColOffsets + 2)))
157 loc,
"im2col offsets must be 2 less than number of coordinates");
166 if (!tensorSize.empty() && coordinates.size() != tensorSize.size()) {
168 emitError(loc,
"Expected coordinates size to be equal to tensor size");
171 if (!lowerStride.empty() && tensorSize.empty()) {
174 "Expected tensor_size to be present when lower_stride is provided");
175 }
else if (!lowerStride.empty() &&
176 lowerStride.size() != tensorSize.size() - 1) {
179 "Expected lower_stride size to be equal to one less than tensor size");
182 if (!lowerStride.empty() !=
static_cast<bool>(upperStride)) {
184 "Expected lower_stride and upper_stride to be either both "
185 "present or both absent");
188 bool isDimStride = tensorSize.size() > 0;
189 if (!isTile && isDimStride) {
191 loc,
"Only tile mode supports override address with dim and stride");
197LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
198 TMAStoreMode mode = getMode();
202 if (getPredicate()) {
203 if (mode != TMAStoreMode::TILE)
204 return emitError(
"Inline-ptx lowering supported only for Tile mode.");
205 if (getL2CacheHint())
206 return emitError(
"Inline-ptx lowering unsupported with L2 cache-hint.");
211 case TMAStoreMode::TILE:
213 case TMAStoreMode::IM2COL:
214 case TMAStoreMode::IM2COL_W:
216 case TMAStoreMode::TILE_SCATTER4:
218 return emitError(
"Scatter4 mode expects 5 coordinates");
223LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::verify() {
224 TMAStoreMode mode = getMode();
226 mode == TMAStoreMode::IM2COL || mode == TMAStoreMode::IM2COL_W;
227 bool isTile = mode == TMAStoreMode::TILE;
233 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
236 if (mode == TMAStoreMode::TILE_SCATTER4 &&
getCoordinates().size() != 5)
237 overrideAddrRes =
emitError(
"Mode tile scatter4 expects 5 coordinates");
242LogicalResult CpAsyncOp::verify() {
243 if (getModifier() != LoadCacheModifierKind::CG &&
244 getModifier() != LoadCacheModifierKind::CA)
245 return emitError(
"Only CG and CA cache modifiers are supported.");
246 if (getSize() != 4 && getSize() != 8 && getSize() != 16)
247 return emitError(
"expected byte size to be either 4, 8 or 16.");
248 if (getModifier() == LoadCacheModifierKind::CG && getSize() != 16)
249 return emitError(
"CG cache modifier is only support for 16 bytes copy.");
256 if (tensorDims < 1 || tensorDims > 5)
257 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
259 auto checkTMALoadParams = [&](TMALoadMode mode,
bool isIm2col,
260 size_t expectedIm2colOff) -> LogicalResult {
261 if (isIm2col && (tensorDims < 3))
264 <<
" mode, the tensor has to be at least 3-dimensional";
266 if (numIm2colOff != expectedIm2colOff)
267 return emitError(loc) <<
" im2col offsets expected " << expectedIm2colOff
268 <<
" (provided " << numIm2colOff <<
")";
274 case TMALoadMode::TILE:
275 return checkTMALoadParams(mode,
false, 0);
276 case TMALoadMode::IM2COL:
277 return checkTMALoadParams(mode,
true, tensorDims - 2);
278 case TMALoadMode::IM2COL_W:
279 case TMALoadMode::IM2COL_W_128:
280 return checkTMALoadParams(mode,
true, 2);
281 case TMALoadMode::TILE_GATHER4:
282 return (tensorDims == 5)
283 ? checkTMALoadParams(mode,
false, 0)
284 :
emitError(loc,
"Gather4 mode expects 5 coordinates");
289LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
291 getMode(), getLoc());
294LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
295 TMALoadMode mode = getMode();
296 bool isCTAOnly = getIsCTAOnly();
297 if (getPredicate()) {
299 return emitError(
"Predicate is supported only for shared::cluster mode.");
300 if (mode != TMALoadMode::TILE && mode != TMALoadMode::IM2COL)
302 "Predicate is supported only for Tile and Im2col modes.");
304 NVVMMemorySpace expectedAS =
305 isCTAOnly ? NVVMMemorySpace::Shared : NVVMMemorySpace::SharedCluster;
306 unsigned AS = llvm::cast<LLVM::LLVMPointerType>(getDstMem().
getType())
308 if (AS != expectedAS)
311 ?
"Shared::cta destination requires address-space 3."
312 :
"Shared::cluster destination requires address-space 7.");
315 if (getMulticastMask())
316 return emitError(
"Multicast is not supported with shared::cta mode.");
318 return emitError(
"CTAGroup is not supported with shared::cta mode.");
323 getMode(), getLoc());
326LogicalResult CpAsyncBulkTensorReduceOp::verify() {
327 TMAStoreMode mode = getMode();
330 case TMAStoreMode::TILE:
332 case TMAStoreMode::IM2COL:
333 case TMAStoreMode::IM2COL_W:
335 case TMAStoreMode::TILE_SCATTER4:
336 return emitError(
"Scatter mode unsupported for CpAsyncBulkTensorReduceOp");
341LogicalResult CpAsyncBulkTensorReduceOverrideAddrOp::verify() {
343 getMode() == TMAStoreMode::IM2COL || getMode() == TMAStoreMode::IM2COL_W;
344 bool isTile = getMode() == TMAStoreMode::TILE;
350 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
353 if (getMode() == TMAStoreMode::TILE_SCATTER4)
355 "Scatter mode unsupported for CpAsyncBulkTensorReduceOverrideAddrOp");
360LogicalResult CpAsyncBulkGlobalToSharedClusterOp::verify() {
362 if (isSharedCTA && getMulticastMask())
363 return emitError(
"Multicast is not supported with shared::cta mode.");
369 NVVM::MemScopeKind scope,
370 Value retVal =
nullptr) {
371 if (scope != NVVM::MemScopeKind::CTA && scope != NVVM::MemScopeKind::CLUSTER)
372 return op->
emitError(
"mbarrier scope must be either CTA or Cluster");
375 bool hasRetValue =
static_cast<bool>(retVal);
376 if (isSharedCluster && hasRetValue)
378 "mbarrier in shared_cluster space cannot return any value");
383LogicalResult MBarrierArriveOp::verify() {
388LogicalResult MBarrierArriveDropOp::verify() {
393LogicalResult MBarrierArriveExpectTxOp::verify() {
397 if (getPredicate()) {
398 if (getScope() != NVVM::MemScopeKind::CTA)
399 return emitError(
"mbarrier scope must be CTA when using predicate");
402 return emitError(
"mbarrier in shared_cluster space is not supported when "
406 return emitError(
"return-value is not supported when using predicate");
408 if (getRelaxed() ==
true)
409 return emitError(
"mbarrier with relaxed semantics is not supported when "
416LogicalResult MBarrierArriveDropExpectTxOp::verify() {
431 inferredReturnTypes.push_back(IntegerType::get(context, 64));
436MBarrierArriveOp::inferReturnTypes(
MLIRContext *context,
437 std::optional<Location> location,
438 MBarrierArriveOp::Adaptor adaptor,
441 inferredReturnTypes);
444LogicalResult MBarrierArriveDropOp::inferReturnTypes(
445 MLIRContext *context, std::optional<Location> location,
446 MBarrierArriveDropOp::Adaptor adaptor,
449 inferredReturnTypes);
452LogicalResult MBarrierArriveExpectTxOp::inferReturnTypes(
453 MLIRContext *context, std::optional<Location> location,
454 MBarrierArriveExpectTxOp::Adaptor adaptor,
458 if (adaptor.getPredicate())
461 inferredReturnTypes);
464LogicalResult MBarrierArriveDropExpectTxOp::inferReturnTypes(
465 MLIRContext *context, std::optional<Location> location,
466 MBarrierArriveDropExpectTxOp::Adaptor adaptor,
469 inferredReturnTypes);
479 return inferred == actual;
488bool MBarrierArriveExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
492bool MBarrierArriveDropExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
497LogicalResult MBarrierExpectTxOp::verify() {
501LogicalResult MBarrierCompleteTxOp::verify() {
505LogicalResult MBarrierTestWaitOp::verify() {
509LogicalResult MBarrierTryWaitOp::verify() {
513LogicalResult ConvertFloatToTF32Op::verify() {
514 using RndMode = NVVM::FPRoundingMode;
518 return emitError(
"Relu not supported with rna rounding mode.");
525 "Only {rn,rz,rna} rounding modes supported for ConvertFloatToTF32Op.");
530LogicalResult ConvertF32x2ToF6x2Op::verify() {
533 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy())) {
534 return emitOpError(
"Only ")
535 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
536 << mlir::Float6E3M2FNType::get(ctx)
537 <<
" types are supported for conversions from f32x2 to f6x2.";
542LogicalResult ConvertF32x2ToF8x2Op::verify() {
543 using RndMode = NVVM::FPRoundingMode;
544 using SatMode = NVVM::SaturationMode;
546 bool isRoundingModeRN = getRnd() == RndMode::RN;
547 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
548 bool isRoundingModeRP = getRnd() == RndMode::RP;
549 bool isSatFinite = getSat() == SatMode::SATFINITE;
551 bool hasRelu = getRelu();
552 bool hasPzo = getPzo();
557 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
559 if (!(isRoundingModeRN || isRoundingModeRZ)) {
560 return emitOpError(
"Only RN and RZ rounding modes are supported "
561 "for conversions from f32x2 to ")
562 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
563 << mlir::Float8E5M2Type::get(ctx) <<
" types";
566 return emitOpError(
"Only SATFINITE saturation mode is supported "
569 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
570 << mlir::Float8E5M2Type::get(ctx) <<
" types";
574 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
575 if (!(isRoundingModeRZ || isRoundingModeRP)) {
576 return emitOpError(
"Only RZ and RP rounding modes are supported for "
577 "conversions from f32x2 to ")
578 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
580 if (hasRelu || hasPzo) {
581 return emitOpError() << (hasRelu ?
"relu" :
"pzo")
582 <<
" not supported for conversions to "
583 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
588 return emitOpError(
"Only ")
589 << mlir::Float8E4M3FNType::get(ctx) <<
", "
590 << mlir::Float8E5M2Type::get(ctx) <<
", and "
591 << mlir::Float8E8M0FNUType::get(ctx)
593 "supported for conversions from f32x2 to f8x2";
597LogicalResult ConvertF16x2ToF8x2Op::verify() {
600 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy())) {
601 return emitOpError(
"Only ")
602 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
603 << mlir::Float8E5M2Type::get(ctx)
604 <<
" types are supported for conversions from f16x2 to f8x2.";
609LogicalResult ConvertBF16x2ToF8x2Op::verify() {
610 using RndMode = NVVM::FPRoundingMode;
611 using SatMode = NVVM::SaturationMode;
613 bool isRoundingModeRN = getRnd() == RndMode::RN;
614 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
615 bool isRoundingModeRP = getRnd() == RndMode::RP;
616 bool isSatFinite = getSat() == SatMode::SATFINITE;
617 bool hasRelu = getRelu();
618 bool hasPzo = getPzo();
623 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
625 if (!(isRoundingModeRN || isRoundingModeRZ))
626 return emitOpError(
"Only RN and RZ rounding modes are supported "
627 "for conversions from bf16x2 to ")
628 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
629 << mlir::Float8E5M2Type::get(ctx) <<
" types";
631 return emitOpError(
"Only SATFINITE saturation mode is supported "
632 "for conversions from bf16x2 to ")
633 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
634 << mlir::Float8E5M2Type::get(ctx) <<
" types";
637 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
638 if (!(isRoundingModeRZ || isRoundingModeRP))
639 return emitOpError(
"Only RZ and RP rounding modes are supported for "
640 "conversions from bf16x2 to ")
641 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
642 if (hasRelu || hasPzo)
643 return emitOpError() << (hasRelu ?
"relu" :
"pzo")
644 <<
" not supported for conversions to "
645 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
649 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
654LogicalResult ConvertF32x2ToF4x2Op::verify() {
657 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
658 return emitOpError(
"Only ")
659 << mlir::Float4E2M1FNType::get(ctx)
660 <<
" type is supported for conversions from f32x2 to f4x2.";
665LogicalResult ConvertF8x2ToBF16x2Op::verify() {
667 if (llvm::isa<Float8E8M0FNUType>(getSrcType())) {
668 if (getSat() != SaturationMode::NONE)
670 "Only NONE saturation mode is supported for conversions from ")
671 << Float8E8M0FNUType::get(ctx) <<
" type";
672 if (getScaleFactor())
673 return emitOpError(
"scaleFactor not supported for conversions from ")
674 << Float8E8M0FNUType::get(ctx) <<
" type";
676 return emitOpError(
"relu not supported for conversions from ")
677 << Float8E8M0FNUType::get(ctx) <<
" type";
683LogicalResult PermuteOp::verify() {
684 using Mode = NVVM::PermuteMode;
685 bool hasHi =
static_cast<bool>(getHi());
692 return emitError(
"mode '") << getMode() <<
"' requires 'hi' operand.";
700 << getMode() <<
"' does not accept 'hi' operand.";
715 static constexpr FPRoundingMode validRndModes[] = {
716 FPRoundingMode::RN, FPRoundingMode::RZ, FPRoundingMode::RS};
718 if (!llvm::is_contained(validRndModes, rnd)) {
720 "Only RN, RZ, and RS rounding modes are supported for "
721 "conversions from f32x2 to ")
725 if (rnd == FPRoundingMode::RS) {
727 return op->
emitOpError(
"random_bits is required for RS rounding mode.");
729 return op->
emitOpError(
"pzo not supported for RS rounding mode.");
733 "random_bits not supported for RN and RZ rounding modes.");
740LogicalResult ConvertF32x2ToF16x2Op::verify() {
742 "f16x2", getRnd(), getRandomBits() ?
true :
false, getPzo(), *
this);
745LogicalResult ConvertF32x2ToBF16x2Op::verify() {
747 "bf16x2", getRnd(), getRandomBits() ?
true :
false, getPzo(), *
this);
750LogicalResult ConvertF32x4ToF8x4Op::verify() {
753 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy()))
754 return emitOpError(
"Only ")
755 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
756 << mlir::Float8E5M2Type::get(ctx)
757 <<
" types are supported for conversions from f32x4 to f8x4.";
762LogicalResult ConvertF32x4ToF6x4Op::verify() {
765 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy()))
766 return emitOpError(
"Only ")
767 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
768 << mlir::Float6E3M2FNType::get(ctx)
769 <<
" types are supported for conversions from f32x4 to f6x4.";
774LogicalResult ConvertF32x4ToF4x4Op::verify() {
777 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
778 return emitOpError(
"Only ") << mlir::Float4E2M1FNType::get(ctx)
779 <<
" type is supported for conversions from "
785LogicalResult BulkStoreOp::verify() {
786 if (getInitVal() != 0)
787 return emitOpError(
"only 0 is supported for initVal, got ") << getInitVal();
791LogicalResult AsyncStoreGlobalOp::verify() {
792 NVVM::MemScopeKind scope = getScope();
793 bool isMmio = getMmio();
794 bool isMultimem = getMultimem();
796 if (scope != MemScopeKind::SYS && scope != MemScopeKind::GPU)
797 return emitOpError(
"scope must be either SYS or GPU");
799 if (isMmio && scope != MemScopeKind::SYS)
800 return emitOpError(
"mmio is only supported for SYS scope");
802 if (isMmio && isMultimem)
803 return emitOpError(
"multimem is not supported with mmio");
808LogicalResult PMEventOp::verify() {
809 auto eventId = getEventId();
810 auto maskedEventId = getMaskedEventId();
811 if (!maskedEventId && !eventId) {
812 return emitOpError() <<
"either `id` or `mask` must be set";
815 if (maskedEventId && eventId) {
816 return emitOpError() <<
"`id` and `mask` cannot be set at the same time";
820 if (eventId < 0 || eventId > 15) {
821 return emitOpError() <<
"`id` must be between 0 and 15";
825 return llvm::success();
831std::optional<mlir::NVVM::MMATypes>
832MmaOp::inferOperandMMAType(
Type operandElType,
bool isAccumulator) {
834 VectorType::get(2, Float16Type::get(operandElType.
getContext()));
835 if (operandElType.
isF64())
836 return NVVM::MMATypes::f64;
837 if (operandElType.
isF16() || operandElType == half2Type)
838 return NVVM::MMATypes::f16;
839 if (operandElType.
isF32() && isAccumulator)
840 return NVVM::MMATypes::f32;
841 if (operandElType.
isF32() && !isAccumulator)
842 return NVVM::MMATypes::tf32;
843 if (llvm::isa<IntegerType>(operandElType)) {
845 return NVVM::MMATypes::s32;
849 if (
auto structType = llvm::dyn_cast<LLVM::LLVMStructType>(operandElType)) {
850 if (structType.getBody().empty())
852 return inferOperandMMAType(structType.getBody()[0], isAccumulator);
859 return (type == MMATypes::u4 || type == MMATypes::s4);
863 return (type == MMATypes::u8 || type == MMATypes::s8);
868 type == MMATypes::s32;
871MMATypes MmaOp::accumPtxType() {
872 std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
873 getODSOperands(2).getTypes().front(),
true);
874 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
878MMATypes MmaOp::resultPtxType() {
879 std::optional<mlir::NVVM::MMATypes> val =
880 inferOperandMMAType(getResult().
getType(),
true);
881 assert(val.has_value() &&
"result PTX type should always be inferrable");
885template <
typename AttrTy>
887 StringRef keyword, AttrTy value) {
888 printer << (isFirst ?
" " :
", ") << keyword <<
" = ";
893template <
typename AttrTy>
895 StringRef keyword, AttrTy value) {
896 printer << (isFirst ?
" " :
", ") << keyword <<
" = "
897 << NVVM::stringifyEnum(value.getValue());
903 printer << (isFirst ?
" " :
", ") << keyword;
907template <
typename AttrTy>
911 if (attributes.
get(name))
913 "duplicate property '" + name +
"'");
917 attributes.
append(name, value);
921template <
typename AttrTy>
925 if (attributes.
get(name))
927 "duplicate property '" + name +
"'");
931 attributes.
append(name, value);
936 return llvm::is_contained(
938 "shape",
"b1Op",
"intOverflowBehavior",
"layoutA",
"layoutB",
939 "multiplicandAPtxType",
"multiplicandBPtxType",
"orderedMetadata",
940 "kind",
"scaleVecSize",
"blockScaleFormat",
"operandSegmentSizes"},
952 if (!llvm::is_contained(allowedKeywords, keyword))
954 "unknown MMA property '" + keyword +
"'");
956 ParseResult parseResult =
success();
957 if (keyword ==
"shape")
960 else if (keyword ==
"b1_op")
963 else if (keyword ==
"int_overflow")
965 parser, attributes,
"intOverflowBehavior");
966 else if (keyword ==
"layout_a")
969 else if (keyword ==
"layout_b")
972 else if (keyword ==
"multiplicand_a_ptx_type")
974 parser, attributes,
"multiplicandAPtxType");
975 else if (keyword ==
"multiplicand_b_ptx_type")
977 parser, attributes,
"multiplicandBPtxType");
978 else if (keyword ==
"kind") {
979 if (llvm::is_contained(allowedKeywords,
"block_scale_format"))
981 parser, attributes,
"kind");
985 }
else if (keyword ==
"scale_vec_size")
987 parser, attributes,
"scaleVecSize");
988 else if (keyword ==
"block_scale_format")
990 parser, attributes,
"blockScaleFormat");
991 else if (keyword ==
"ordered_metadata") {
992 if (attributes.
get(
"orderedMetadata"))
994 "duplicate property 'orderedMetadata'");
998 "unknown MMA property '" + keyword +
"'");
1000 if (failed(parseResult))
1006 for (StringRef property : requiredProperties) {
1007 if (!attributes.
get(property))
1009 "missing required property '" + property +
"'");
1019 "inherent property '" + attribute.getName().getValue() +
1020 "' must be spelled directly in the operation syntax");
1021 attributes.
append(attribute);
1028 struct MMAOperandFragment {
1029 StringRef operandName;
1030 StringRef ptxTypeAttr;
1031 SmallVector<Value, 4> regs;
1032 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1033 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1036 std::array<MMAOperandFragment, 3> frags{
1037 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1038 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1039 MMAOperandFragment(
"C",
"")};
1041 mlir::NVVM::MmaOp::getOperandSegmentSizeAttr()};
1043 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
1044 auto &frag = frags[fragIdx];
1045 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1046 for (
auto operandIdx = varOperandSpec.first;
1047 operandIdx < varOperandSpec.first + varOperandSpec.second;
1049 frag.regs.push_back(this->getOperand(operandIdx));
1050 if (operandIdx == 0) {
1051 regTypes.push_back(this->getOperand(operandIdx).
getType());
1054 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1055 regTypes.back(), fragIdx >= 2);
1057 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1060 auto printMmaOperand = [&](
const MMAOperandFragment &frag) ->
void {
1061 p <<
" " << frag.operandName;
1067 for (
const auto &frag : frags) {
1068 printMmaOperand(frag);
1071 bool isFirstProperty =
true;
1075 if (getIntOverflowBehaviorAttr())
1077 getIntOverflowBehaviorAttr());
1080 if (getMultiplicandAPtxTypeAttr() &&
1081 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1083 getMultiplicandAPtxTypeAttr());
1084 if (getMultiplicandBPtxTypeAttr() &&
1085 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1087 getMultiplicandBPtxTypeAttr());
1088 llvm::append_range(ignoreAttrNames,
1090 getIntOverflowBehaviorAttrName(),
1091 getLayoutAAttrName(),
1092 getLayoutBAttrName(),
1093 getMultiplicandAPtxTypeAttrName(),
1094 getMultiplicandBPtxTypeAttrName()});
1101 frags[1].regs[0].getType(),
1102 frags[2].regs[0].getType()},
1111 std::optional<MMAIntOverflow> intOverflow,
1112 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
1113 std::optional<std::array<MMALayout, 2>> multiplicandLayouts) {
1115 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1120 result.addOperands(operandA);
1121 result.addOperands(operandB);
1122 result.addOperands(operandC);
1124 if (multiplicandPtxTypes) {
1125 result.addAttribute(
"multiplicandAPtxType",
1126 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1127 result.addAttribute(
"multiplicandBPtxType",
1128 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1130 if (
auto res = inferOperandMMAType(operandA[0].
getType(),
false))
1131 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1132 if (
auto res = inferOperandMMAType(operandB[0].
getType(),
false))
1133 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1136 if (multiplicandLayouts) {
1137 result.addAttribute(
"layoutA",
1138 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[0]));
1139 result.addAttribute(
"layoutB",
1140 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[1]));
1142 result.addAttribute(
"layoutA", MMALayoutAttr::get(ctx, MMALayout::row));
1143 result.addAttribute(
"layoutB", MMALayoutAttr::get(ctx, MMALayout::col));
1146 if (intOverflow.has_value())
1147 result.addAttribute(
"intOverflowBehavior",
1148 MMAIntOverflowAttr::get(ctx, *intOverflow));
1149 if (b1Op.has_value())
1150 result.addAttribute(
"b1Op", MMAB1OpAttr::get(ctx, *b1Op));
1152 result.addTypes(resultType);
1154 MmaOp::getOperandSegmentSizeAttr(),
1156 static_cast<int32_t>(operandB.size()),
1157 static_cast<int32_t>(operandC.size())}));
1166 struct MMAOperandFragment {
1167 std::optional<MMATypes> elemtype;
1168 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1169 SmallVector<Type> regTypes;
1173 std::array<MMAOperandFragment, 4> frags;
1179 MMAOperandFragment &frag) -> LogicalResult {
1198 {
"shape",
"b1_op",
"int_overflow",
"layout_a",
1199 "layout_b",
"multiplicand_a_ptx_type",
1200 "multiplicand_b_ptx_type"},
1201 {
"shape",
"layoutA",
"layoutB"}))
1213 if (operandTypes.size() != 3)
1216 "expected one type for each operand segment but got " +
1217 Twine(operandTypes.size()) +
" types");
1218 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1219 auto &frag = frags[iter.index()];
1220 frag.regTypes.resize(frag.regs.size(), iter.value());
1224 frag.elemtype = inferOperandMMAType(frag.regTypes[0],
1231 frags[3].elemtype = inferOperandMMAType(resultType,
true);
1233 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1234 "multiplicandBPtxType"};
1235 for (
unsigned idx = 0; idx < names.size(); idx++) {
1236 const auto &frag = frags[idx];
1237 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1238 if (!frag.elemtype.has_value() && !attr.has_value()) {
1241 "attribute " + names[idx] +
1242 " is not provided explicitly and cannot be inferred");
1244 if (!attr.has_value())
1246 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1249 result.addTypes(resultType);
1250 if (!namedAttributes.
empty())
1251 result.addAttributes(namedAttributes);
1252 result.addAttribute(MmaOp::getOperandSegmentSizeAttr(),
1254 static_cast<int32_t>(frags[0].regs.size()),
1255 static_cast<int32_t>(frags[1].regs.size()),
1256 static_cast<int32_t>(frags[2].regs.size()),
1261LogicalResult MmaOp::verify() {
1263 auto f16Ty = Float16Type::get(context);
1264 auto i32Ty = IntegerType::get(context, 32);
1265 auto f16x2Ty = VectorType::get(2, f16Ty);
1266 auto f32Ty = Float32Type::get(context);
1267 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1268 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1270 auto s32x4StructTy =
1271 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1272 auto f32x8StructTy =
1274 auto f16x2x2StructTy =
1275 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1276 auto f32x4StructTy =
1277 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1278 auto s32x2StructTy =
1279 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1281 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1282 getShapeAttr().getK()};
1288 AllowedShapes allowedShapes;
1289 AllowedTypes expectedA;
1290 AllowedTypes expectedB;
1291 AllowedTypes expectedC;
1296 if (mmaShape[0] == 16) {
1298 Type multiplicandFragType;
1299 switch (*getMultiplicandAPtxType()) {
1300 case MMATypes::tf32:
1302 multiplicandFragType = i32Ty;
1303 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1304 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1306 case MMATypes::bf16:
1308 multiplicandFragType = i32Ty;
1309 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1310 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1314 multiplicandFragType = f16x2Ty;
1315 expectedResult.push_back(f16x2x2StructTy);
1316 expectedResult.push_back(f32x4StructTy);
1318 case MMATypes::e4m3:
1319 case MMATypes::e5m2:
1323 multiplicandFragType = i32Ty;
1324 expectedResult.push_back(f16x2x2StructTy);
1325 expectedResult.push_back(f32x4StructTy);
1339 return emitError(
"invalid shape or multiplicand type: ")
1340 << getMultiplicandAPtxType().value();
1344 expectedResult.push_back(s32x4StructTy);
1345 expectedC.emplace_back(4, i32Ty);
1346 multiplicandFragType = i32Ty;
1348 expectedC.emplace_back(2, f16x2Ty);
1349 expectedC.emplace_back(4, f32Ty);
1352 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor);
1353 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1354 expectedA.emplace_back(unitA, multiplicandFragType);
1355 expectedB.emplace_back(unitB, multiplicandFragType);
1356 allowedShapes.push_back({16, 8, kFactor});
1357 allowedShapes.push_back({16, 8, kFactor * 2});
1359 if (resultPtxType() != accumPtxType())
1360 return emitOpError(
"ctype does not match dtype");
1364 if (mmaShape[0] == 8) {
1365 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1366 expectedA.emplace_back(2, f16x2Ty);
1367 expectedB.emplace_back(2, f16x2Ty);
1368 expectedResult.push_back(f16x2x4StructTy);
1369 expectedResult.push_back(f32x8StructTy);
1370 expectedC.emplace_back(4, f16x2Ty);
1371 expectedC.emplace_back(8, f32Ty);
1372 allowedShapes.push_back({8, 8, 4});
1374 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1375 Type f64Ty = Float64Type::get(context);
1376 expectedA.emplace_back(1, f64Ty);
1377 expectedB.emplace_back(1, f64Ty);
1378 expectedC.emplace_back(2, f64Ty);
1379 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1381 allowedShapes.push_back({8, 8, 4});
1384 expectedA.push_back({i32Ty});
1385 expectedB.push_back({i32Ty});
1386 expectedC.push_back({i32Ty, i32Ty});
1387 expectedResult.push_back(s32x2StructTy);
1389 allowedShapes.push_back({8, 8, 32});
1391 allowedShapes.push_back({8, 8, 16});
1392 if (getMultiplicandAPtxType().value() == MMATypes::b1)
1393 allowedShapes.push_back({8, 8, 128});
1397 std::string errorMessage;
1398 llvm::raw_string_ostream errorStream(errorMessage);
1401 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1402 !llvm::is_contained(allowedShapes, mmaShape)) {
1403 errorStream <<
"unimplemented variant for MMA shape <";
1404 llvm::interleaveComma(mmaShape, errorStream);
1406 return emitOpError(errorMessage);
1410 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1411 for (
const auto &iter : llvm::enumerate(
1412 std::array<AllowedTypes, 3>{std::move(expectedA),
1413 std::move(expectedB),
1414 std::move(expectedC)})) {
1415 auto spec = this->getODSOperandIndexAndLength(iter.index());
1417 operand_type_begin() + spec.first +
1419 bool match = llvm::is_contained(iter.value(), operandTySeg);
1422 errorStream <<
"Could not match types for the "
1423 << operandNames[iter.index()]
1424 <<
" operands; expected one of ";
1425 for (
const auto &x : iter.value()) {
1426 errorStream << x.size() <<
"x" << x[0] <<
" ";
1428 errorStream <<
"but got ";
1429 llvm::interleaveComma(operandTySeg, errorStream);
1430 return emitOpError(errorMessage);
1435 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1436 return expectedResultType == getResult().getType();
1439 <<
"Could not match allowed types for the result; expected one of ";
1440 llvm::interleaveComma(expectedResult, errorStream);
1441 errorStream <<
" but got " << getResult().getType();
1442 return emitOpError(errorMessage);
1446 if (getMultiplicandAPtxType() == MMATypes::b1 && !getB1Op()) {
1447 return emitOpError(
"op requires " + getB1OpAttrName().strref() +
1455 if (!getIntOverflowBehavior())
1456 return emitOpError(
"op requires " +
1457 getIntOverflowBehaviorAttrName().strref() +
1465 (mmaShape[0] == 8 && mmaShape[1] == 8 && mmaShape[2] == 4 &&
1466 getMultiplicandAPtxType() == MMATypes::f16);
1468 if (!isM8N8K4_F16) {
1470 if (getLayoutA() != MMALayout::row || getLayoutB() != MMALayout::col) {
1471 return emitOpError(
"requires layoutA = #nvvm.mma_layout<row> and "
1472 "layoutB = #nvvm.mma_layout<col> for shape <")
1473 << mmaShape[0] <<
", " << mmaShape[1] <<
", " << mmaShape[2]
1474 <<
"> with element types " << *getMultiplicandAPtxType() <<
" and "
1475 << *getMultiplicandBPtxType()
1476 <<
". Only m8n8k4 with f16 supports other layouts.";
1483MMATypes MmaSpOp::accumPtxType() {
1484 std::optional<mlir::NVVM::MMATypes> val = MmaOp::inferOperandMMAType(
1485 getODSOperands(2).getTypes().front(),
true);
1486 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
1490MMATypes MmaSpOp::resultPtxType() {
1491 std::optional<mlir::NVVM::MMATypes> val =
1492 MmaOp::inferOperandMMAType(getResult().
getType(),
true);
1493 assert(val.has_value() &&
"result PTX type should always be inferrable");
1499 llvm::IRBuilderBase &builder) {
1500 auto thisOp = cast<NVVM::MmaSpOp>(op);
1508 auto intId = MmaSpOp::getIntrinsicID(
1509 thisOp.getShape().getM(), thisOp.getShape().getN(),
1510 thisOp.getShape().getK(), thisOp.getIntOverflowBehavior(),
1511 thisOp.getOrderedMetadata(), thisOp.getKind(),
1512 *thisOp.getMultiplicandAPtxType(), *thisOp.getMultiplicandBPtxType(),
1513 thisOp.accumPtxType(), thisOp.resultPtxType());
1515 return {intId, args};
1520 struct MMAOperandFragment {
1521 StringRef operandName;
1522 StringRef ptxTypeAttr;
1523 SmallVector<Value, 4> regs;
1524 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1525 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1528 std::array<MMAOperandFragment, 5> frags{
1529 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1530 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1531 MMAOperandFragment(
"C",
""), MMAOperandFragment(
"sparseMetadata",
""),
1532 MMAOperandFragment(
"selector",
"")};
1534 mlir::NVVM::MmaSpOp::getOperandSegmentSizeAttr()};
1537 for (
unsigned fragIdx = 0; fragIdx < 3; fragIdx++) {
1538 auto &frag = frags[fragIdx];
1539 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1540 for (
auto operandIdx = varOperandSpec.first;
1541 operandIdx < varOperandSpec.first + varOperandSpec.second;
1543 frag.regs.push_back(this->getOperand(operandIdx));
1544 if (operandIdx == varOperandSpec.first) {
1545 regTypes.push_back(this->getOperand(operandIdx).
getType());
1548 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1549 regTypes.back(), fragIdx >= 2);
1551 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1555 frags[3].regs.push_back(getSparseMetadata());
1556 frags[4].regs.push_back(getSparsitySelector());
1558 auto printMmaSpOperand = [&](
const MMAOperandFragment &frag) ->
void {
1559 p <<
" " << frag.operandName;
1565 for (
const auto &frag : frags)
1566 printMmaSpOperand(frag);
1568 bool isFirstProperty =
true;
1570 if (getIntOverflowBehaviorAttr())
1572 getIntOverflowBehaviorAttr());
1573 if (getMultiplicandAPtxTypeAttr() &&
1574 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1576 getMultiplicandAPtxTypeAttr());
1577 if (getMultiplicandBPtxTypeAttr() &&
1578 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1580 getMultiplicandBPtxTypeAttr());
1581 if (getOrderedMetadata())
1588 getMultiplicandAPtxTypeAttrName(),
1589 getMultiplicandBPtxTypeAttrName(),
1590 getOrderedMetadataAttrName(), getKindAttrName()});
1595 for (
int i = 0; i < 3; ++i) {
1600 p <<
") -> " << getResult().getType();
1607 std::optional<MMAIntOverflow> intOverflow,
1608 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
1610 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1615 result.addOperands(operandA);
1616 result.addOperands(operandB);
1617 result.addOperands(operandC);
1618 result.addOperands(sparseMetadata);
1619 result.addOperands(sparsitySelector);
1621 if (multiplicandPtxTypes) {
1622 result.addAttribute(
"multiplicandAPtxType",
1623 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1624 result.addAttribute(
"multiplicandBPtxType",
1625 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1627 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
1628 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1629 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
1630 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1633 if (intOverflow.has_value())
1634 result.addAttribute(
"intOverflowBehavior",
1635 MMAIntOverflowAttr::get(ctx, *intOverflow));
1637 result.addTypes(resultType);
1639 MmaSpOp::getOperandSegmentSizeAttr(),
1641 static_cast<int32_t>(operandB.size()),
1642 static_cast<int32_t>(operandC.size()), 1,
1647 struct MMAOperandFragment {
1648 std::optional<MMATypes> elemtype;
1649 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1650 SmallVector<Type> regTypes;
1654 std::array<MMAOperandFragment, 6> frags;
1659 auto parseMmaSpOperand = [&](StringRef operandName,
1660 MMAOperandFragment &frag) -> LogicalResult {
1671 if (parseMmaSpOperand(
"A", frags[0]).
failed())
1673 if (parseMmaSpOperand(
"B", frags[1]).
failed())
1675 if (parseMmaSpOperand(
"C", frags[2]).
failed())
1677 if (parseMmaSpOperand(
"sparseMetadata", frags[3]).
failed())
1679 if (parseMmaSpOperand(
"selector", frags[4]).
failed())
1683 {
"shape",
"int_overflow",
"multiplicand_a_ptx_type",
1684 "multiplicand_b_ptx_type",
"ordered_metadata",
1699 if (operandTypes.size() != 3)
1702 "expected one type for each operand segment but got " +
1703 Twine(operandTypes.size()) +
" types");
1704 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1705 auto &frag = frags[iter.index()];
1706 frag.regTypes.resize(frag.regs.size(), iter.value());
1711 MmaOp::inferOperandMMAType(frag.regTypes[0],
1719 MmaOp::inferOperandMMAType(resultType,
true);
1734 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1735 "multiplicandBPtxType"};
1736 for (
unsigned idx = 0; idx < names.size(); idx++) {
1737 const auto &frag = frags[idx];
1738 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1739 if (!frag.elemtype.has_value() && !attr.has_value()) {
1742 "attribute " + names[idx] +
1743 " is not provided explicitly and cannot be inferred");
1745 if (!attr.has_value())
1747 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1750 result.addTypes(resultType);
1751 if (!namedAttributes.
empty())
1752 result.addAttributes(namedAttributes);
1753 result.addAttribute(MmaSpOp::getOperandSegmentSizeAttr(),
1755 static_cast<int32_t>(frags[0].regs.size()),
1756 static_cast<int32_t>(frags[1].regs.size()),
1757 static_cast<int32_t>(frags[2].regs.size()),
1764LogicalResult MmaSpOp::verify() {
1766 auto f16Ty = Float16Type::get(context);
1767 auto i32Ty = IntegerType::get(context, 32);
1768 auto f16x2Ty = VectorType::get(2, f16Ty);
1769 auto f32Ty = Float32Type::get(context);
1770 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1771 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1773 auto s32x4StructTy =
1774 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1775 auto f32x8StructTy =
1777 auto f16x2x2StructTy =
1778 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1779 auto f32x4StructTy =
1780 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1781 auto s32x2StructTy =
1782 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1784 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1785 getShapeAttr().getK()};
1791 AllowedShapes allowedShapes;
1792 AllowedTypes expectedA;
1793 AllowedTypes expectedB;
1794 AllowedTypes expectedC;
1799 if (mmaShape[0] == 16) {
1801 Type multiplicandFragType;
1802 switch (*getMultiplicandAPtxType()) {
1803 case MMATypes::tf32:
1805 multiplicandFragType = i32Ty;
1806 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1807 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1809 allowedShapes.push_back({16, 8, 8});
1810 allowedShapes.push_back({16, 8, 16});
1812 case MMATypes::bf16:
1814 multiplicandFragType = i32Ty;
1815 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1816 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1818 allowedShapes.push_back({16, 8, 16});
1819 allowedShapes.push_back({16, 8, 32});
1823 multiplicandFragType = f16x2Ty;
1824 expectedResult.push_back(f16x2x2StructTy);
1825 expectedResult.push_back(f32x4StructTy);
1827 allowedShapes.push_back({16, 8, 16});
1828 allowedShapes.push_back({16, 8, 32});
1834 allowedShapes.push_back({16, 8, 64});
1835 allowedShapes.push_back({16, 8, 128});
1841 allowedShapes.push_back({16, 8, 32});
1842 allowedShapes.push_back({16, 8, 64});
1844 case MMATypes::e4m3:
1845 case MMATypes::e5m2:
1846 case MMATypes::e3m2:
1847 case MMATypes::e2m3:
1848 case MMATypes::e2m1:
1850 multiplicandFragType = i32Ty;
1851 expectedResult.push_back(f16x2x2StructTy);
1852 expectedResult.push_back(f32x4StructTy);
1854 allowedShapes.push_back({16, 8, 64});
1857 return emitError(
"invalid shape or multiplicand type: ")
1858 << getMultiplicandAPtxType().value();
1862 expectedResult.push_back(s32x4StructTy);
1863 expectedC.emplace_back(4, i32Ty);
1864 multiplicandFragType = i32Ty;
1865 }
else if (*getMultiplicandAPtxType() >= MMATypes::e4m3 &&
1866 *getMultiplicandAPtxType() <= MMATypes::e2m1) {
1868 expectedC.emplace_back(2, f16x2Ty);
1869 expectedC.emplace_back(4, f32Ty);
1871 expectedC.emplace_back(2, f16x2Ty);
1872 expectedC.emplace_back(4, f32Ty);
1877 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor) / 2;
1878 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1879 expectedA.emplace_back(unitA, multiplicandFragType);
1880 expectedB.emplace_back(unitB, multiplicandFragType);
1882 if (resultPtxType() != accumPtxType())
1883 return emitOpError(
"ctype does not match dtype");
1887 if (mmaShape[0] == 8) {
1888 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1889 expectedA.emplace_back(2, f16x2Ty);
1890 expectedB.emplace_back(2, f16x2Ty);
1891 expectedResult.push_back(f16x2x4StructTy);
1892 expectedResult.push_back(f32x8StructTy);
1893 expectedC.emplace_back(4, f16x2Ty);
1894 expectedC.emplace_back(8, f32Ty);
1895 allowedShapes.push_back({8, 8, 4});
1897 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1898 Type f64Ty = Float64Type::get(context);
1899 expectedA.emplace_back(1, f64Ty);
1900 expectedB.emplace_back(1, f64Ty);
1901 expectedC.emplace_back(2, f64Ty);
1902 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1904 allowedShapes.push_back({8, 8, 4});
1907 expectedA.push_back({i32Ty});
1908 expectedB.push_back({i32Ty});
1909 expectedC.push_back({i32Ty, i32Ty});
1910 expectedResult.push_back(s32x2StructTy);
1912 allowedShapes.push_back({8, 8, 32});
1914 allowedShapes.push_back({8, 8, 16});
1918 std::string errorMessage;
1919 llvm::raw_string_ostream errorStream(errorMessage);
1922 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1923 !llvm::is_contained(allowedShapes, mmaShape)) {
1924 errorStream <<
"unimplemented variant for MMA shape <";
1925 llvm::interleaveComma(mmaShape, errorStream);
1927 return emitOpError(errorMessage);
1931 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1932 for (
const auto &iter : llvm::enumerate(
1933 std::array<AllowedTypes, 3>{std::move(expectedA),
1934 std::move(expectedB),
1935 std::move(expectedC)})) {
1936 auto spec = this->getODSOperandIndexAndLength(iter.index());
1938 operand_type_begin() + spec.first +
1940 bool match = llvm::is_contained(iter.value(), operandTySeg);
1943 errorStream <<
"Could not match types for the "
1944 << operandNames[iter.index()]
1945 <<
" operands; expected one of ";
1946 for (
const auto &x : iter.value()) {
1947 errorStream << x.size() <<
"x" << x[0] <<
" ";
1949 errorStream <<
"but got ";
1950 llvm::interleaveComma(operandTySeg, errorStream);
1951 return emitOpError(errorMessage);
1956 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1957 return expectedResultType == getResult().getType();
1960 <<
"Could not match allowed types for the result; expected one of ";
1961 llvm::interleaveComma(expectedResult, errorStream);
1962 errorStream <<
" but got " << getResult().getType();
1963 return emitOpError(errorMessage);
1970 if (!getIntOverflowBehavior())
1971 return emitOpError(
"op requires " +
1972 getIntOverflowBehaviorAttrName().strref() +
1977 if (!getSparseMetadata().
getType().isInteger(32)) {
1978 return emitOpError() <<
"sparse metadata must be i32 type";
1982 if (!getSparsitySelector().
getType().isInteger(32)) {
1983 return emitOpError() <<
"sparsity selector must be i32 type";
1995struct MMAOperandFragment {
1996 StringRef operandName;
1997 StringRef ptxTypeAttr;
1998 SmallVector<Value, 4> regs;
1999 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
2000 : operandName(name), ptxTypeAttr(ptxTypeName) {}
2007 p <<
" " << name <<
"[";
2026template <
typename Op>
2031 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
2032 auto &frag = frags[fragIdx];
2033 auto varOperandSpec = op.getODSOperandIndexAndLength(fragIdx);
2034 for (
auto operandIdx = varOperandSpec.first;
2035 operandIdx < varOperandSpec.first + varOperandSpec.second;
2037 frag.regs.push_back(op.getOperand(operandIdx));
2038 if (fragIdx == 0 && operandIdx == varOperandSpec.first) {
2039 regTypes.push_back(op.getOperand(operandIdx).getType());
2043 regTypes.push_back(frag.regs[0].getType());
2045 std::optional<MMATypes> inferredType =
2046 MmaOp::inferOperandMMAType(regTypes.back(),
2049 ignoreAttrNames.push_back(frag.ptxTypeAttr);
2060 auto typeParser = [&]() {
2064 operandTypes.push_back(ty);
2070 if (operandTypes.size() != 3)
2072 "expected exactly 3 types");
2081 if (!attrs.
get(
"multiplicandAPtxType")) {
2082 if (
auto inferredType =
2083 MmaOp::inferOperandMMAType(operandTypes[0],
false)) {
2084 attrs.
set(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *inferredType));
2087 if (!attrs.
get(
"multiplicandBPtxType")) {
2088 if (
auto inferredType =
2089 MmaOp::inferOperandMMAType(operandTypes[1],
false)) {
2090 attrs.
set(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *inferredType));
2096template <
typename OpType>
2099 ScaleVecSize scaleVecSize,
2100 BlockScaleFormat blockScaleFormat,
2101 MMABlockScaleKind kind) {
2103 auto &properties =
result.getOrAddProperties<
typename OpType::Properties>();
2104 properties.setShape(
2106 properties.setScaleVecSize(ScaleVecSizeAttr::get(ctx, scaleVecSize));
2107 properties.setBlockScaleFormat(
2108 BlockScaleFormatAttr::get(ctx, blockScaleFormat));
2109 properties.setKind(MMABlockScaleKindAttr::get(ctx, kind));
2116 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
2117 if (multiplicandPtxTypes) {
2118 result.addAttribute(
"multiplicandAPtxType",
2119 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
2120 result.addAttribute(
"multiplicandBPtxType",
2121 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
2123 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
2124 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
2125 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
2126 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
2131template <
typename OpTy>
2133 return *MmaOp::inferOperandMMAType(
2134 cast<LLVM::LLVMStructType>(op.getRes().getType()).getBody()[0],
2144 std::array<MMAOperandFragment, 3> frags{
2145 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2146 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2147 MMAOperandFragment(
"C",
"")};
2149 mlir::NVVM::MmaBlockScaleOp::getOperandSegmentSizeAttr()};
2154 for (
const auto &frag : frags)
2159 {getScaleAData(), getByteIdA(), getThreadIdA()});
2161 {getScaleBData(), getByteIdB(), getThreadIdB()});
2163 bool isFirstProperty =
true;
2165 if (getMultiplicandAPtxTypeAttr() &&
2166 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2168 getMultiplicandAPtxTypeAttr());
2169 if (getMultiplicandBPtxTypeAttr() &&
2170 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2172 getMultiplicandBPtxTypeAttr());
2174 getScaleVecSizeAttr());
2176 getBlockScaleFormatAttr());
2181 getMultiplicandBPtxTypeAttrName(),
2182 getScaleVecSizeAttrName(),
2183 getBlockScaleFormatAttrName(), getKindAttrName()});
2190 frags[1].regs[0].getType(),
2191 frags[2].regs[0].getType()},
2197ParseResult MmaBlockScaleOp::parse(
OpAsmParser &parser,
2199 struct LocalOperandFragment {
2200 std::optional<MMATypes> elemtype;
2201 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2205 std::array<LocalOperandFragment, 3> frags;
2221 {
"shape",
"multiplicand_a_ptx_type",
2222 "multiplicand_b_ptx_type",
"scale_vec_size",
2223 "block_scale_format",
"kind"},
2224 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2238 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2239 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2242 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2252 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2262 result.addAttributes(namedAttributes);
2266 result.addTypes(resultTypes);
2267 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2269 static_cast<int32_t>(frags[0].regs.size()),
2270 static_cast<int32_t>(frags[1].regs.size()),
2271 static_cast<int32_t>(frags[2].regs.size()),
2282void MmaBlockScaleOp::build(
2287 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2288 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2289 MMABlockScaleKind kind) {
2290 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2293 blockScaleFormat, kind);
2295 result.addOperands(operandA);
2296 result.addOperands(operandB);
2297 result.addOperands(operandC);
2299 {scaleAData, byteIdA, threadIdA, scaleBData, byteIdB, threadIdB});
2302 multiplicandPtxTypes);
2304 result.addTypes(resultType);
2305 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2307 static_cast<int32_t>(operandA.size()),
2308 static_cast<int32_t>(operandB.size()),
2309 static_cast<int32_t>(operandC.size()),
2321 auto curOp = cast<NVVM::MmaBlockScaleOp>(op);
2325 for (
Value operand : curOp.getOperandA())
2327 for (
Value operand : curOp.getOperandB())
2329 for (
Value operand : curOp.getOperandC())
2333 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2334 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2335 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2336 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2337 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2338 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2340 unsigned intId = MmaBlockScaleOp::getIntrinsicID(
2341 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2342 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2344 curOp.getBlockScaleFormat(), curOp.getKind());
2346 return {intId, args};
2349LogicalResult MmaBlockScaleOp::verify() {
2355 if (m == 16 && n == 8 && k == 64) {
2356 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2357 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2359 "unsupported MMATypes attribute for mma.m16n8k64.(mxf4nvf4|mxf4)");
2360 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2361 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2363 "unsupported ScaleVecSize attribute for mma.m16n8k64.mxf4");
2364 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2366 "unsupported BlockScaleFormat attribute for mma.m16n8k64.mxf4");
2367 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2368 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2369 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2370 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2371 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2372 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2373 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2374 "attributes for mma.m16n8k64.mxf4nvf4");
2376 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k64");
2378 }
else if (m == 16 && n == 8 && k == 32) {
2379 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2380 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2381 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2383 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2384 "attributes for mma.m16n8k32");
2386 result = emitOpError(
"unsupported Geom for mma with block scaling");
2397 std::array<MMAOperandFragment, 3> frags{
2398 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2399 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2400 MMAOperandFragment(
"C",
"")};
2402 mlir::NVVM::MmaSpBlockScaleOp::getOperandSegmentSizeAttr()};
2407 for (
const auto &frag : frags)
2416 {getScaleAData(), getByteIdA(), getThreadIdA()});
2418 {getScaleBData(), getByteIdB(), getThreadIdB()});
2420 bool isFirstProperty =
true;
2422 if (getMultiplicandAPtxTypeAttr() &&
2423 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2425 getMultiplicandAPtxTypeAttr());
2426 if (getMultiplicandBPtxTypeAttr() &&
2427 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2429 getMultiplicandBPtxTypeAttr());
2432 getScaleVecSizeAttr());
2434 getBlockScaleFormatAttr());
2439 getMultiplicandBPtxTypeAttrName(),
2440 getOrderedMetadataAttrName(),
2441 getScaleVecSizeAttrName(),
2442 getBlockScaleFormatAttrName(), getKindAttrName()});
2449 frags[1].regs[0].getType(),
2450 frags[2].regs[0].getType()},
2456ParseResult MmaSpBlockScaleOp::parse(
OpAsmParser &parser,
2458 struct LocalOperandFragment {
2459 std::optional<MMATypes> elemtype;
2460 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2464 std::array<LocalOperandFragment, 3> frags;
2487 {
"shape",
"multiplicand_a_ptx_type",
2488 "multiplicand_b_ptx_type",
"ordered_metadata",
2489 "scale_vec_size",
"block_scale_format",
"kind"},
2490 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2504 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2505 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2508 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2517 .resolveOperands(metadataOperands, i32Type, parser.
getNameLoc(),
2530 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2540 result.addAttributes(namedAttributes);
2545 if (!
result.attributes.get(
"orderedMetadata"))
2548 result.addTypes(resultTypes);
2549 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2551 static_cast<int32_t>(frags[0].regs.size()),
2552 static_cast<int32_t>(frags[1].regs.size()),
2553 static_cast<int32_t>(frags[2].regs.size()),
2566void MmaSpBlockScaleOp::build(
2572 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2573 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2574 MMABlockScaleKind kind) {
2575 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2578 builder,
result,
shape, scaleVecSize, blockScaleFormat, kind);
2581 result.addOperands(operandA);
2582 result.addOperands(operandB);
2583 result.addOperands(operandC);
2584 result.addOperands({sparseMetadata, sparsitySelector, scaleAData, byteIdA,
2585 threadIdA, scaleBData, byteIdB, threadIdB});
2588 multiplicandPtxTypes);
2590 result.addTypes(resultType);
2591 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2593 static_cast<int32_t>(operandA.size()),
2594 static_cast<int32_t>(operandB.size()),
2595 static_cast<int32_t>(operandC.size()),
2609 auto curOp = cast<NVVM::MmaSpBlockScaleOp>(op);
2613 for (
Value operand : curOp.getOperandA())
2615 for (
Value operand : curOp.getOperandB())
2617 for (
Value operand : curOp.getOperandC())
2621 args.push_back(mt.
lookupValue(curOp.getSparseMetadata()));
2622 args.push_back(mt.
lookupValue(curOp.getSparsitySelector()));
2625 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2626 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2627 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2628 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2629 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2630 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2632 unsigned intId = MmaSpBlockScaleOp::getIntrinsicID(
2633 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2634 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2636 curOp.getBlockScaleFormat(), curOp.getKind());
2638 return {intId, args};
2641LogicalResult MmaSpBlockScaleOp::verify() {
2643 if (!getOrderedMetadata()) {
2644 return emitOpError(
"'orderedMetadata' attribute is mandatory");
2652 if (m == 16 && n == 8 && k == 128) {
2653 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2654 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2656 "unsupported MMATypes attribute for mma.m16n8k128.(mxf4nvf4|mxf4)");
2657 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2658 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2660 "unsupported ScaleVecSize attribute for mma.m16n8k128.mxf4");
2661 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2663 "unsupported BlockScaleFormat attribute for mma.m16n8k128.mxf4");
2664 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2665 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2666 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2667 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2668 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2669 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2670 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2671 "attributes for mma.m16n8k128.mxf4nvf4");
2673 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k128");
2675 }
else if (m == 16 && n == 8 && k == 64) {
2676 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2677 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2678 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2680 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2681 "attributes for mma.m16n8k64");
2683 result = emitOpError(
"unsupported Geom for sparse mma with block scaling");
2688LogicalResult ShflOp::verify() {
2689 auto returnStructType = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
2691 auto verifyTypeError = [&](Twine desc,
Type expectedType,
2692 Type actualType) -> LogicalResult {
2693 return emitOpError(
"expected " + desc +
" to be of type ")
2694 << expectedType <<
" but got " << actualType <<
" instead";
2697 if (returnStructType) {
2698 if (!getReturnValueAndIsValid())
2699 return emitOpError(
"\"return_value_and_is_valid\" attribute must be "
2700 "specified when the return type is a struct type");
2702 if (returnStructType.getBody().size() != 2)
2703 return emitOpError(
"expected return type to be a two-element struct");
2706 auto resultType = returnStruct[0];
2707 if (resultType != getVal().
getType())
2708 return verifyTypeError(
"first element in the returned struct",
2709 getVal().
getType(), resultType);
2711 auto predicateType = returnStruct[1];
2712 if (!predicateType.isInteger(1))
2713 return verifyTypeError(
"second element in the returned struct",
2717 if (getReturnValueAndIsValid())
2718 return emitOpError(
"expected return type to be a two-element struct");
2721 return verifyTypeError(
"return type", getVal().
getType(),
getType());
2727ShflOp::inferReturnTypes(
MLIRContext *context, std::optional<Location> location,
2728 ShflOp::Adaptor adaptor,
2730 Type valType = adaptor.getVal().getType();
2731 if (adaptor.getReturnValueAndIsValid())
2732 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2733 context, {valType, IntegerType::get(context, 1)}));
2735 inferredReturnTypes.push_back(valType);
2740 NVVM::MMAFrag frag,
int nRow,
2743 unsigned numberElements = 0;
2746 Type f16x2 = VectorType::get(2, builder.getF16Type());
2747 if (type == NVVM::MMATypes::f16) {
2748 elementType = f16x2;
2749 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2753 }
else if (type == NVVM::MMATypes::f32) {
2754 elementType = builder.getF32Type();
2756 }
else if (type == NVVM::MMATypes::f64) {
2757 elementType = builder.getF64Type();
2758 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2762 }
else if (type == NVVM::MMATypes::tf32) {
2763 elementType = builder.getI32Type();
2765 }
else if (type == NVVM::MMATypes::s8 || type == NVVM::MMATypes::u8) {
2766 elementType = builder.getI32Type();
2767 int parallelSize = 0;
2768 if (frag == NVVM::MMAFrag::a)
2769 parallelSize = nRow;
2770 if (frag == NVVM::MMAFrag::b)
2771 parallelSize = nCol;
2774 if (parallelSize == 16)
2777 else if (parallelSize == 8)
2779 else if (parallelSize == 32)
2781 }
else if (type == NVVM::MMATypes::s32) {
2782 elementType = builder.getI32Type();
2785 assert(numberElements != 0 && elementType !=
nullptr);
2786 return std::make_pair(elementType, numberElements);
2789static std::pair<mlir::Type, unsigned>
2793 if (frag == NVVM::MMAFrag::a) {
2796 }
else if (frag == NVVM::MMAFrag::b) {
2803 assert(nRow && nCol);
2807LogicalResult NVVM::WMMALoadOp::verify() {
2808 unsigned addressSpace =
2809 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2810 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2811 addressSpace != NVVMMemorySpace::Shared)
2812 return emitOpError(
"expected source pointer in memory "
2815 if (NVVM::WMMALoadOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2816 getEltype(), getFrag()) == 0)
2817 return emitOpError() <<
"invalid attribute combination";
2822 if (typeInfo.first == f64Ty && typeInfo.second == 1) {
2824 return emitOpError(
"expected destination type to be f64");
2828 Type dstType = LLVM::LLVMStructType::getLiteral(
2831 return emitOpError(
"expected destination type is a structure of ")
2832 << typeInfo.second <<
" elements of type " << typeInfo.first;
2836LogicalResult NVVM::WMMAStoreOp::verify() {
2837 unsigned addressSpace =
2838 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2839 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2840 addressSpace != NVVMMemorySpace::Shared)
2841 return emitOpError(
"expected operands to be a source pointer in memory "
2844 if (NVVM::WMMAStoreOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2846 return emitOpError() <<
"invalid attribute combination";
2849 if (getArgs().size() != typeInfo.second)
2850 return emitOpError() <<
"expected " << typeInfo.second <<
" data operands";
2851 if (llvm::any_of(getArgs(), [&typeInfo](
Value operands) {
2852 return operands.
getType() != typeInfo.first;
2854 return emitOpError() <<
"expected data operands of type " << typeInfo.first;
2858LogicalResult NVVM::WMMAMmaOp::verify() {
2859 if (NVVM::WMMAMmaOp::getIntrinsicID(
getM(),
getN(), getK(), getLayoutA(),
2860 getLayoutB(), getEltypeA(),
2862 return emitOpError() <<
"invalid attribute combination";
2870 arguments.append(typeInfoA.second, typeInfoA.first);
2871 arguments.append(typeInfoB.second, typeInfoB.first);
2872 arguments.append(typeInfoC.second, typeInfoC.first);
2873 unsigned numArgs = arguments.size();
2874 if (getArgs().size() != numArgs)
2875 return emitOpError() <<
"expected " << numArgs <<
" arguments";
2876 for (
unsigned i = 0; i < numArgs; i++) {
2877 if (getArgs()[i].
getType() != arguments[i])
2878 return emitOpError() <<
"expected argument " << i <<
" to be of type "
2881 Type dstType = LLVM::LLVMStructType::getLiteral(
2884 return emitOpError(
"expected destination type is a structure of ")
2885 << typeInfoC.second <<
" elements of type " << typeInfoC.first;
2889LogicalResult NVVM::LdMatrixOp::verify() {
2891 if (m == 8 && n == 8) {
2892 if (num != 1 && num != 2 && num != 4) {
2893 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x8 "
2896 if (getEltType() != LdStMatrixEltType::B16) {
2897 return emitOpError(
"expected element type to be b16 for 8x8 matrix");
2899 }
else if (m == 8 && n == 16) {
2900 if (num != 1 && num != 2 && num != 4) {
2901 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x16 "
2904 if (getLayout() != MMALayout::row) {
2905 return emitOpError(
"expected layout to be row for 8x16 matrix");
2907 if (getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2908 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2909 return emitOpError(
"expected element type to be b8x16.b4x16_p64 or "
2910 "b8x16.b6x16_p32 for 8x16 matrix");
2912 }
else if (m == 16 && n == 16) {
2913 if (num != 1 && num != 2) {
2914 return emitOpError(
"expected num attribute to be 1 or 2 for 16x16 "
2917 if (getLayout() != MMALayout::col) {
2918 return emitOpError(
"expected layout to be col for 16x16 matrix");
2920 if (getEltType() != LdStMatrixEltType::B8 &&
2921 getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2922 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2923 return emitOpError(
"expected element type to be b8, b8x16.b4x16_p64 or "
2924 "b8x16.b6x16_p32 for 16x16 matrix");
2927 return emitOpError(
"expected shape to be 8x8, 8x16 or 16x16");
2931 uint32_t numElements = (m == 16 && n == 16 ? num * 2 : num);
2932 if (numElements == 1 &&
getType() != i32)
2933 return emitOpError(
"expected destination type is i32");
2934 if (numElements == 2 || numElements == 4) {
2935 Type dstType = LLVM::LLVMStructType::getLiteral(
2938 return emitOpError(
"expected destination type is a structure of ")
2939 << numElements <<
" elements of type i32";
2945LogicalResult LdMatrixOp::inferReturnTypes(
2946 MLIRContext *context, std::optional<Location> location,
2948 uint32_t num = adaptor.getNum();
2949 uint32_t m = adaptor.getShape().getM();
2950 uint32_t n = adaptor.getShape().getN();
2951 uint32_t numElements = (m == 16 && n == 16) ? num * 2 : num;
2953 Type i32 = IntegerType::get(context, 32);
2954 if (numElements == 1)
2955 inferredReturnTypes.push_back(i32);
2957 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2962LogicalResult NVVM::StMatrixOp::verify() {
2963 int numMatrix = getSources().size();
2964 if (numMatrix != 1 && numMatrix != 2 && numMatrix != 4)
2965 return emitOpError(
"expected num attribute to be 1, 2 or 4");
2968 if (m == 8 && n == 8) {
2969 if (getEltType() != NVVM::LdStMatrixEltType::B16) {
2970 return emitOpError(
"expected element type to be B16 for 8x8 matrix");
2972 }
else if (m == 16 && n == 8) {
2973 if (getEltType() != NVVM::LdStMatrixEltType::B8) {
2974 return emitOpError(
"expected element type to be B8 for 16x8 matrix");
2976 if (getLayout() != NVVM::MMALayout::col) {
2977 return emitOpError(
"expected layout to be col for 16x8 matrix");
2980 return emitOpError(
"expected shape to be 8x8 or 16x8");
2986LogicalResult NVVM::MovMatrixOp::verify() {
2988 if (m != 8 || n != 8)
2989 return emitOpError(
"expected shape to be 8x8");
2990 if (getLayout() != NVVM::MMALayout::col)
2991 return emitOpError(
"expected layout to be col");
2992 if (getEltType() != NVVM::LdStMatrixEltType::B16)
2993 return emitOpError(
"expected element type to be b16");
2998 if (typeA == NVVM::WGMMATypes::tf32)
3000 if (typeA == NVVM::WGMMATypes::f16 || typeA == NVVM::WGMMATypes::bf16)
3002 if (typeA == NVVM::WGMMATypes::s8 || typeA == NVVM::WGMMATypes::u8)
3004 if (typeA == NVVM::WGMMATypes::e4m3 || typeA == NVVM::WGMMATypes::e5m2)
3006 if (typeA == NVVM::WGMMATypes::b1)
3012 NVVM::WGMMATypes typeA,
3013 NVVM::WGMMATypes typeB) {
3015 case NVVM::WGMMATypes::f16:
3016 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3017 typeB == NVVM::WGMMATypes::f16)
3020 case NVVM::WGMMATypes::tf32:
3021 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::tf32)
3024 case NVVM::WGMMATypes::u8:
3025 case NVVM::WGMMATypes::s8:
3026 if (typeD == NVVM::WGMMATypes::s32 &&
3027 (typeB == NVVM::WGMMATypes::u8 || typeB == NVVM::WGMMATypes::s8))
3030 case NVVM::WGMMATypes::b1:
3031 if (typeD == NVVM::WGMMATypes::s32 && typeB == NVVM::WGMMATypes::b1)
3034 case NVVM::WGMMATypes::bf16:
3035 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::bf16)
3038 case NVVM::WGMMATypes::e4m3:
3039 case NVVM::WGMMATypes::e5m2:
3040 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3041 (typeB == NVVM::WGMMATypes::e5m2 || typeB == NVVM::WGMMATypes::e4m3))
3044 case WGMMATypes::f32:
3045 case WGMMATypes::s32:
3046 llvm_unreachable(
"unsupported input types");
3054 72, 80, 88, 96, 104, 112, 120, 128,
3055 136, 144, 152, 160, 168, 176, 184, 192,
3056 200, 208, 216, 224, 232, 240, 248, 256};
3058 80, 96, 112, 128, 144, 160,
3059 176, 192, 208, 224, 240, 256};
3061 case WGMMATypes::f16:
3062 case WGMMATypes::tf32:
3063 case WGMMATypes::bf16:
3064 case WGMMATypes::e4m3:
3065 case WGMMATypes::e5m2:
3066 if (llvm::is_contained(allowedN, sizeN))
3069 case WGMMATypes::u8:
3070 case WGMMATypes::s8:
3071 case WGMMATypes::b1:
3072 if (llvm::is_contained(allowedNshort, sizeN))
3075 case WGMMATypes::f32:
3076 case WGMMATypes::s32:
3077 llvm_unreachable(
"unsupported input types");
3083LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
3084 Value outValue = getResults();
3085 auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.
getType());
3087 return emitOpError() <<
"expected results to be struct";
3088 int outputSize = stype.getBody().size();
3089 WGMMATypes typeD = getTypeD();
3090 WGMMATypes typeA = getTypeA();
3091 WGMMATypes typeB = getTypeB();
3093 for (
Type t : stype.getBody()) {
3094 if (t != stype.getBody().front())
3095 return emitOpError()
3096 <<
"all elements in struct must be same type but there is " << t;
3099 if (typeD != WGMMATypes::f32 && typeD != WGMMATypes::f16 &&
3100 typeD != WGMMATypes::s32) {
3101 return emitOpError() <<
"does not support the given output type " << typeD;
3103 if (typeD == WGMMATypes::s32 &&
3104 (getScaleA() == WGMMAScaleIn::neg || getScaleB() == WGMMAScaleIn::neg)) {
3105 return emitOpError() <<
"has s32 output, scaleA and scaleB cannot be neg";
3109 return emitOpError() << typeD <<
" += " << typeA <<
" * " << typeB
3110 <<
", it is not supported.";
3115 return emitOpError() <<
"shape 'm' must be 64";
3120 return emitOpError() <<
"shape 'k' must be " << allowedK.value()
3121 <<
" for input type " << typeA;
3125 return emitOpError() <<
"has input type " << typeA <<
" n is set to "
3126 <<
getShape().getN() <<
", it is not supported.";
3133 if ((typeA != WGMMATypes::f16 && typeA != WGMMATypes::bf16) &&
3134 (getLayoutA() == mlir::NVVM::MMALayout::col ||
3135 getLayoutB() == mlir::NVVM::MMALayout::row)) {
3136 return emitOpError()
3137 <<
"given layouts layout_a = " << getLayoutA()
3138 <<
" and layout_b = " << getLayoutB() <<
" for input types " << typeA
3140 <<
" requires transpose. However, this is only supported for: "
3141 << MMATypes::f16 <<
" and " << MMATypes::bf16;
3145 int expectedOutput = 0;
3146 if (typeD == WGMMATypes::f32 || typeD == WGMMATypes::s32)
3147 expectedOutput =
getShape().getN() / 2;
3148 if (typeD == WGMMATypes::f16)
3149 expectedOutput =
getShape().getN() / 4;
3150 if (outputSize != expectedOutput) {
3151 return emitOpError() <<
"results " << expectedOutput
3152 <<
", however output struct has " << outputSize
3156 if (typeD != WGMMATypes::s32 &&
3157 getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3158 NVVM::MMAIntOverflow::satfinite) {
3159 return emitOpError()
3160 <<
" `satfinite` can be only used with s32 accumulator, however "
3161 "the current accumulator is "
3168std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
3171 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3173 StringRef outputTypeName = stringifyWGMMATypes(getTypeD());
3175 int expectedOutputRegisters = 0;
3176 if (getTypeD() == WGMMATypes::f16)
3177 expectedOutputRegisters =
getShape().getN() / 4;
3179 expectedOutputRegisters =
getShape().getN() / 2;
3182 llvm::raw_string_ostream ss(ptx);
3187 << ((expectedOutputRegisters * 2) + 2)
3189 "wgmma.mma_async.sync.aligned.m"
3190 << m <<
"n" << n <<
"k" << k <<
"." << outputTypeName <<
"." << getTypeA()
3191 <<
"." << getTypeB();
3192 if (getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3193 NVVM::MMAIntOverflow::satfinite)
3197 for (; regCnt < expectedOutputRegisters; ++regCnt) {
3198 ss <<
"$" << regCnt;
3199 if (regCnt != expectedOutputRegisters - 1)
3205 regCnt = (regCnt * 2);
3206 ss <<
" $" << (regCnt) <<
"," <<
" $" << (regCnt + 1) <<
"," <<
" p";
3207 if (getTypeD() != WGMMATypes::s32) {
3208 ss <<
", $" << (regCnt + 3) <<
", $" << (regCnt + 4);
3212 ss <<
", $" << (regCnt + 5) <<
", $" << (regCnt + 6);
3219bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
3223 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3230 asmValues.push_back({makeConstantI32(rewriter,
static_cast<int>(getScaleD())),
3232 if (getTypeD() != WGMMATypes::s32) {
3233 asmValues.push_back(
3234 {makeConstantI32(rewriter,
3235 getScaleA() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3237 asmValues.push_back(
3238 {makeConstantI32(rewriter,
3239 getScaleB() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3243 asmValues.push_back(
3244 {makeConstantI32(rewriter,
static_cast<int>(getLayoutA())),
3246 asmValues.push_back(
3247 {makeConstantI32(rewriter, 1 -
static_cast<int>(getLayoutB())),
3253LogicalResult NVVM::FenceProxyOp::verify() {
3254 if (getKind() == NVVM::ProxyKind::async_shared && !getSpace().has_value()) {
3255 return emitOpError() <<
"async_shared fence requires space attribute";
3257 if (getKind() != NVVM::ProxyKind::async_shared && getSpace().has_value()) {
3258 return emitOpError() <<
"only async_shared fence can have space attribute";
3263LogicalResult NVVM::FenceProxyAcquireOp::verify() {
3264 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3265 return emitOpError(
"uni-directional proxies only support generic for "
3266 "from_proxy attribute");
3268 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3269 return emitOpError(
"uni-directional proxies only support tensormap "
3270 "for to_proxy attribute");
3274LogicalResult NVVM::FenceProxyReleaseOp::verify() {
3275 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3276 return emitOpError(
"uni-directional proxies only support generic for "
3277 "from_proxy attribute");
3279 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3280 return emitOpError(
"uni-directional proxies only support tensormap "
3281 "for to_proxy attribute");
3285LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
3286 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3287 return emitOpError(
"only generic is support for from_proxy attribute");
3289 if (getToProxy() != NVVM::ProxyKind::async)
3290 return emitOpError(
"only async is supported for to_proxy attribute");
3294LogicalResult NVVM::SetMaxRegisterOp::verify() {
3295 if (getRegCount() % 8)
3296 return emitOpError(
"new register size must be multiple of 8");
3297 if (getRegCount() < 24 || getRegCount() > 256)
3298 return emitOpError(
"new register size must be in between 24 to 256");
3302LogicalResult NVVM::Tcgen05CpOp::verify() {
3303 auto mc = getMulticast();
3305 using SH = Tcgen05CpShape;
3306 using MC = Tcgen05CpMulticast;
3308 case SH::SHAPE_128x256b:
3309 case SH::SHAPE_128x128b:
3310 case SH::SHAPE_4x256b:
3312 return emitError(
"Invalid multicast type for tcgen05.cp Op");
3314 case SH::SHAPE_64x128b:
3315 if (mc != MC::WARPX2_01_23 && mc != MC::WARPX2_02_13)
3316 return emitError(
"Shape 64x128b requires multicast warpx2_01_23 or "
3317 "warpx2_02_13 for tcgen05.cp Op");
3319 case SH::SHAPE_32x128b:
3320 if (mc != MC::WARPX4)
3322 "Shape 32x128b requires multicast warpx4 for tcgen05.cp Op");
3328LogicalResult NVVM::MatchSyncOp::verify() {
3329 if (getKind() == NVVM::MatchSyncKind::all) {
3330 auto type = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
3331 if (!type || type.getBody().size() != 2 ||
3332 !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
3333 return emitOpError(
"match.sync 'all' returns a two element struct with "
3334 "first element as i32 and second element as i1");
3337 if (!
getType().isInteger(32)) {
3338 return emitOpError(
"match.sync 'any' returns an i32");
3344LogicalResult MatchSyncOp::inferReturnTypes(
3345 MLIRContext *context, std::optional<Location> location,
3347 if (adaptor.getKind() == NVVM::MatchSyncKind::all)
3348 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
3350 {IntegerType::get(context, 32), IntegerType::get(context, 1)}));
3352 inferredReturnTypes.push_back(IntegerType::get(context, 32));
3356LogicalResult NVVM::VoteSyncOp::verify() {
3357 if (getKind() == NVVM::VoteSyncKind::ballot) {
3358 if (!
getType().isInteger(32)) {
3359 return emitOpError(
"vote.sync 'ballot' returns an i32");
3362 if (!
getType().isInteger(1)) {
3363 return emitOpError(
"vote.sync 'any', 'all' and 'uni' returns an i1");
3369LogicalResult VoteSyncOp::inferReturnTypes(
3370 MLIRContext *context, std::optional<Location> location,
3372 unsigned width = adaptor.getKind() == NVVM::VoteSyncKind::ballot ? 32 : 1;
3373 inferredReturnTypes.push_back(IntegerType::get(context, width));
3377LogicalResult NVVM::PrefetchOp::verify() {
3378 using MemSpace = NVVM::NVVMMemorySpace;
3379 using CacheLevel = NVVM::PrefetchCacheLevel;
3381 unsigned addressSpace =
3382 llvm::cast<LLVM::LLVMPointerType>(getAddr().
getType()).getAddressSpace();
3383 std::optional<NVVM::CacheEvictionPriority> evictPriority = getEvictPriority();
3384 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = getCacheLevel();
3386 if (getTensormap() && cacheLevel)
3387 return emitOpError(
"cannot specify both tensormap and cache level");
3389 if (getTensormap()) {
3390 if (addressSpace != MemSpace::Generic &&
3391 addressSpace != MemSpace::Constant) {
3393 "prefetch tensormap requires a generic or constant pointer");
3396 if (evictPriority) {
3398 "prefetch tensormap does not support eviction priority");
3401 if (getInParamSpace() && addressSpace != MemSpace::Generic) {
3403 "in_param_space can only be specified for a generic pointer");
3406 }
else if (cacheLevel) {
3407 if (addressSpace != MemSpace::Generic && addressSpace != MemSpace::Global &&
3408 addressSpace != MemSpace::Local) {
3409 return emitOpError(
"prefetch to cache level requires a generic, global, "
3410 "or local pointer");
3414 if (*cacheLevel != CacheLevel::L1) {
3416 "unsupported cache level, the only supported uniform "
3417 "cache level is L1");
3420 if (addressSpace != MemSpace::Generic) {
3422 "prefetch to uniform cache requires a generic pointer");
3426 if (evictPriority) {
3427 if (*cacheLevel != CacheLevel::L2)
3429 "cache eviction priority supported only for cache level L2");
3431 if (addressSpace != MemSpace::Global)
3432 return emitOpError(
"cache eviction priority requires a global pointer");
3434 if (*evictPriority != NVVM::CacheEvictionPriority::EvictNormal &&
3435 *evictPriority != NVVM::CacheEvictionPriority::EvictLast)
3437 "unsupported cache eviction priority, only evict_last and "
3438 "evict_normal are supported");
3442 return emitOpError(
"predicate supported only on prefetch tensormap");
3446 "requires specification of either cache level or tensormap");
3452LogicalResult NVVM::ClusterLaunchControlQueryCancelOp::verify() {
3453 switch (getQueryType()) {
3454 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
3456 return emitOpError(
"is_canceled query type returns an i1");
3458 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
3459 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
3460 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
3461 if (!
getType().isInteger(32)) {
3462 return emitOpError(
"get_first_cta_id_x, get_first_cta_id_y, "
3463 "get_first_cta_id_z query types return an i32");
3470LogicalResult ClusterLaunchControlQueryCancelOp::inferReturnTypes(
3471 MLIRContext *context, std::optional<Location> location,
3472 ClusterLaunchControlQueryCancelOp::Adaptor adaptor,
3475 adaptor.getQueryType() == NVVM::ClusterLaunchControlQueryType::IS_CANCELED
3478 inferredReturnTypes.push_back(IntegerType::get(context, width));
3482LogicalResult NVVM::ReduxOp::verify() {
3485 if (!reduxType.
isF32()) {
3487 return emitOpError(
"abs attribute is supported only for f32 type");
3489 return emitOpError(
"nan attribute is supported only for f32 type");
3492 NVVM::ReductionKind kind = getKind();
3494 case NVVM::ReductionKind::ADD:
3495 case NVVM::ReductionKind::AND:
3496 case NVVM::ReductionKind::OR:
3497 case NVVM::ReductionKind::XOR:
3498 case NVVM::ReductionKind::MAX:
3499 case NVVM::ReductionKind::MIN:
3500 case NVVM::ReductionKind::UMAX:
3501 case NVVM::ReductionKind::UMIN:
3503 return emitOpError(
"'")
3504 << kind <<
"' reduction kind unsupported with " << reduxType
3505 <<
" type. Only supported type is 'i32'.";
3507 case NVVM::ReductionKind::FMIN:
3508 case NVVM::ReductionKind::FMAX:
3509 if (!reduxType.isF32())
3510 return emitOpError(
"'")
3511 << kind <<
"' reduction kind unsupported with " << reduxType
3512 <<
" type. Only supported type is 'f32'.";
3519LogicalResult NVVM::TensormapReplaceOp::verify() {
3520 auto ord = getOrd();
3521 Value newVal = getNewValue();
3522 auto newValAttr = getNewValueAttr();
3523 auto fieldName = stringifyEnum(getField());
3525 if (ord && !llvm::is_contained({NVVM::TensormapField::BOX_DIM,
3526 NVVM::TensormapField::GLOBAL_DIM,
3527 NVVM::TensormapField::GLOBAL_STRIDE,
3528 NVVM::TensormapField::ELEMENT_STRIDE},
3530 return emitOpError(
"ordinal is not supported for ")
3531 << fieldName <<
" field";
3533 auto invalidNewVal = [&](llvm::Twine type) -> std::string {
3534 return llvm::Twine(
"new_value must be specified and must be an " + type +
3535 " for " + llvm::Twine(fieldName) +
" field")
3539 auto invalidNewValAttr = [&]() -> std::string {
3540 return (llvm::Twine(
3541 "new_value_attr must be specified and must be a valid ") +
3542 llvm::Twine(fieldName) +
" attribute for " + fieldName +
" field")
3546 switch (getField()) {
3547 case NVVM::TensormapField::GLOBAL_ADDRESS:
3549 return emitOpError(invalidNewVal(
"i64"));
3551 case NVVM::TensormapField::RANK:
3553 return emitOpError(invalidNewVal(
"i32"));
3555 case NVVM::TensormapField::GLOBAL_STRIDE:
3557 return emitOpError(
"ordinal is required for global_stride field");
3559 return emitOpError(invalidNewVal(
"i64"));
3561 case NVVM::TensormapField::BOX_DIM:
3562 case NVVM::TensormapField::GLOBAL_DIM:
3563 case NVVM::TensormapField::ELEMENT_STRIDE:
3565 return emitOpError(
"ordinal is required for ")
3566 << stringifyEnum(getField()) <<
" field";
3568 return emitOpError(invalidNewVal(
"i32"));
3570 case NVVM::TensormapField::ELEMTYPE:
3571 if (!(newValAttr && llvm::isa<TensormapElemtypeAttr>(*newValAttr)))
3572 return emitOpError(invalidNewValAttr());
3574 case NVVM::TensormapField::INTERLEAVE_LAYOUT:
3575 if (!(newValAttr && llvm::isa<TensormapInterleaveLayoutAttr>(*newValAttr)))
3576 return emitOpError(invalidNewValAttr());
3578 case NVVM::TensormapField::SWIZZLE_MODE:
3579 if (!(newValAttr && llvm::isa<TensormapSwizzleModeAttr>(*newValAttr)))
3580 return emitOpError(invalidNewValAttr());
3582 case NVVM::TensormapField::SWIZZLE_ATOMICITY:
3583 if (!(newValAttr && llvm::isa<TensormapSwizzleAtomicityAttr>(*newValAttr)))
3584 return emitOpError(invalidNewValAttr());
3586 case NVVM::TensormapField::FILL_MODE:
3587 if (!(newValAttr && llvm::isa<TensormapFillModeAttr>(*newValAttr)))
3588 return emitOpError(invalidNewValAttr());
3595template <
typename OpType>
3597 mlir::NVVM::FPRoundingMode rndMode = op.getRnd();
3598 mlir::NVVM::SaturationMode satMode = op.getSat();
3599 bool isFTZ = op.getFtz();
3602 mlir::Type opBaseType = isa<VectorType>(opType)
3603 ? cast<VectorType>(opType).getElementType()
3606 if (opBaseType.
isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3607 return op.emitOpError(
"FTZ and saturation are not supported for ")
3608 << opKind <<
" involving f64 type";
3610 if (opBaseType.
isF16() && !(rndMode == NVVM::FPRoundingMode::RN ||
3611 rndMode == NVVM::FPRoundingMode::NONE))
3612 return op.emitOpError(
"only RN rounding mode is supported for f16 and "
3616 if (opBaseType.
isBF16()) {
3617 if (rndMode != NVVM::FPRoundingMode::RN &&
3618 rndMode != NVVM::FPRoundingMode::NONE)
3619 return op.emitOpError(
"only RN rounding mode is supported for bf16 and "
3622 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3623 return op.emitOpError(
"FTZ and saturation are not supported for bf16 and "
3631LogicalResult NVVM::AddFOp::verify() {
3635LogicalResult NVVM::SubFOp::verify() {
3639LogicalResult NVVM::MulFOp::verify() {
3643LogicalResult NVVM::FmaOp::verify() {
3644 auto opType = getRes().getType();
3645 mlir::NVVM::FPRoundingMode rndMode = getRnd();
3646 mlir::NVVM::SaturationMode satMode = getSat();
3647 bool isFTZ = getFtz();
3648 bool isRelu = getRelu();
3649 bool hasOOB = getOob();
3651 auto getBaseFType = [](
Type type) ->
Type {
3652 if (isa<VectorType>(type))
3653 return cast<VectorType>(type).getElementType();
3657 auto opBaseType = getBaseFType(opType);
3659 if (rndMode == NVVM::FPRoundingMode::NONE)
3660 return emitOpError(
"rounding mode must be specified");
3662 if (isRelu && satMode == NVVM::SaturationMode::SAT)
3663 return emitOpError(
"relu and saturation are not supported together");
3665 if (hasOOB && (satMode == NVVM::SaturationMode::SAT || isFTZ))
3666 return emitOpError(
"oob is not supported with saturation or FTZ");
3668 if (!(opBaseType.isF16() || opBaseType.isBF16()) && (isRelu || hasOOB))
3669 return emitOpError(
"relu and oob are only supported for f16 and bf16");
3671 if (opBaseType.isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3672 return emitOpError(
"FTZ and saturation are not supported for f64 type");
3674 if (opBaseType.isF16() && rndMode != NVVM::FPRoundingMode::RN)
3676 "only RN rounding mode is supported for f16 and vector<2xf16>");
3678 if (opBaseType.isBF16()) {
3679 if (rndMode != NVVM::FPRoundingMode::RN)
3681 "only RN rounding mode is supported for bf16 and vector<2xbf16>");
3682 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3684 "FTZ and saturation are not supported for bf16 and vector<2xbf16>");
3690LogicalResult NVVM::SqrtOp::verify() {
3691 if (getRnd() == NVVM::FPRoundingMode::NONE)
3692 return emitOpError(
"rounding mode cannot be None");
3694 if (getRes().
getType().isF64() && getFtz())
3695 return emitOpError(
"FTZ is not supported for f64");
3700LogicalResult NVVM::DivFOp::verify() {
3701 bool isApprox = getApprox();
3702 bool isFull = getFull();
3703 bool isF64 = getRes().getType().isF64();
3704 bool isFtz = getFtz();
3705 NVVM::FPRoundingMode rndMode = getRnd();
3707 if (isApprox && isFull)
3708 return emitOpError(
"'approx' and 'full' are mutually exclusive");
3710 if (isApprox || isFull) {
3712 return emitOpError(
"'approx' and 'full' forms are f32-only");
3713 if (rndMode != NVVM::FPRoundingMode::NONE)
3715 "'approx' and 'full' forms do not accept a rounding mode");
3720 if (rndMode == NVVM::FPRoundingMode::NONE)
3721 return emitOpError(
"rounding mode cannot be None for the rounded divide");
3723 return emitOpError(
"FTZ is not supported for f64");
3734 unsigned sizeInBits,
3736 field = builder.CreateZExtOrBitCast(field, builder.getInt32Ty());
3738 unsigned mask = (sizeInBits < 32 ? ((1u << sizeInBits) - 1) : 0xffffffffu);
3739 if (mask != 0xffffffffu)
3740 field = builder.CreateAnd(field, builder.getInt32(mask));
3742 field = builder.CreateZExtOrBitCast(field, builder.getInt64Ty());
3743 field = builder.CreateShl(field, start);
3745 return builder.CreateOr(
result, field);
3748void Tcgen05MmaSmemDescOp::createSmemDescriptor(
Operation &op,
3750 llvm::IRBuilderBase &builder) {
3751 auto thisOp = cast<NVVM::Tcgen05MmaSmemDescOp>(op);
3752 llvm::Value *smemDesc = builder.getInt64(0);
3757 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimOffset()), 14, 16);
3759 builder, smemDesc, mt.
lookupValue(thisOp.getStrideDimOffset()), 14, 32);
3765 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimMode()), 1, 52);
3769 mt.
mapValue(thisOp.getRes()) = smemDesc;
3776std::string NVVM::MBarrierInitOp::getPtx() {
3779 std::string layout =
3780 getLayout() == 1 ? std::string(
".layout::v1") : std::string();
3782 return llvm::formatv(
"mbarrier.init{0}{1}.b64 [%0], %1;", layout, space)
3786std::string NVVM::MBarrierArriveExpectTxOp::getPtx() {
3789 ? std::string(
"mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;")
3790 : std::string(
"mbarrier.arrive.expect_tx.b64 _, [%0], %1;");
3793std::string NVVM::MBarrierTryWaitParityOp::getPtx() {
3795 llvm::StringRef space = isShared ?
".shared" :
"";
3797 return llvm::formatv(
"{\n\t"
3798 ".reg .pred P1; \n\t"
3800 "mbarrier.try_wait.parity{0}.b64 P1, [%0], %1, %2; \n\t"
3801 "@P1 bra.uni DONE; \n\t"
3802 "bra.uni LAB_WAIT; \n\t"
3819 LLVM::FNegOp::create(rewriter, loc, op.getRhs().getType(), op.getRhs());
3822 op.getRnd(), op.getSat(), op.getFtz());
3841 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count
3842 : llvm::Intrinsic::nvvm_barrier_cta_sync_count;
3844 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all
3845 : llvm::Intrinsic::nvvm_barrier_cta_sync_all;
3850static llvm::Intrinsic::ID
3853 case NVVM::BarrierReduction::AND:
3854 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_and_aligned_all
3855 : llvm::Intrinsic::nvvm_barrier_cta_red_and_all;
3856 case NVVM::BarrierReduction::OR:
3857 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_or_aligned_all
3858 : llvm::Intrinsic::nvvm_barrier_cta_red_or_all;
3859 case NVVM::BarrierReduction::POPC:
3860 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_popc_aligned_all
3861 : llvm::Intrinsic::nvvm_barrier_cta_red_popc_all;
3863 llvm_unreachable(
"unknown BarrierReduction kind");
3868 auto thisOp = cast<NVVM::BarrierOp>(op);
3869 llvm::Value *barrierId = thisOp.getBarrierId()
3871 : builder.getInt32(0);
3872 bool hasCount =
static_cast<bool>(thisOp.getNumberOfThreads());
3873 llvm::Intrinsic::ID
id =
3877 args.push_back(mt.
lookupValue(thisOp.getNumberOfThreads()));
3878 return {id, std::move(args)};
3883 auto thisOp = cast<NVVM::BarrierArriveOp>(op);
3884 llvm::Value *barrierId = thisOp.getBarrierId()
3886 : builder.getInt32(0);
3887 llvm::Value *numThreads = mt.
lookupValue(thisOp.getNumberOfThreads());
3888 llvm::Intrinsic::ID
id =
3890 ? llvm::Intrinsic::nvvm_barrier_cta_arrive_aligned_count
3891 : llvm::Intrinsic::nvvm_barrier_cta_arrive_count;
3892 return {id, {barrierId, numThreads}};
3897 auto thisOp = cast<NVVM::BarrierReductionOp>(op);
3899 thisOp.getAligned(), thisOp.getReductionOp());
3900 llvm::Value *barrierId = thisOp.getBarrierId()
3902 : builder.getInt32(0);
3905 builder.CreateICmpNE(mt.
lookupValue(thisOp.getReductionPredicate()),
3906 builder.getInt32(0))};
3907 return {id, std::move(args)};
3912 llvm::IRBuilderBase &builder) {
3913 auto thisOp = cast<NVVM::CosOp>(op);
3914 llvm::Intrinsic::ID
id = thisOp.getFtz()
3915 ? llvm::Intrinsic::nvvm_cos_approx_ftz_f
3916 : llvm::Intrinsic::nvvm_cos_approx_f;
3922 llvm::IRBuilderBase &builder) {
3923 auto thisOp = cast<NVVM::SinOp>(op);
3924 llvm::Intrinsic::ID
id = thisOp.getFtz()
3925 ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
3926 : llvm::Intrinsic::nvvm_sin_approx_f;
3932 llvm::IRBuilderBase &builder) {
3933 auto thisOp = cast<NVVM::Log2Op>(op);
3934 llvm::Intrinsic::ID
id = thisOp.getFtz()
3935 ? llvm::Intrinsic::nvvm_lg2_approx_ftz_f
3936 : llvm::Intrinsic::nvvm_lg2_approx_f;
3942 llvm::IRBuilderBase &builder) {
3943 auto thisOp = cast<NVVM::Ex2Op>(op);
3944 llvm::Intrinsic::ID
id = thisOp.getFtz()
3945 ? llvm::Intrinsic::nvvm_ex2_approx_ftz
3946 : llvm::Intrinsic::nvvm_ex2_approx;
3950LogicalResult NVVM::Ex2Op::verify() {
3951 auto vectorType = dyn_cast<VectorType>(getSrc().
getType());
3955 if (vectorType.getElementType().isF16() && getFtz())
3956 return emitOpError(
"FTZ is not supported for vector<2xf16>");
3957 if (vectorType.getElementType().isBF16() && !getFtz())
3958 return emitOpError(
"FTZ is required for vector<2xbf16>");
3964 llvm::IRBuilderBase &builder) {
3965 auto thisOp = cast<NVVM::RsqrtOp>(op);
3966 Type t = thisOp.getRes().getType();
3967 bool isFtz = thisOp.getFtz();
3969 llvm::Intrinsic::ID
id = [&] {
3971 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_f
3972 : llvm::Intrinsic::nvvm_rsqrt_approx_f;
3975 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_d
3976 : llvm::Intrinsic::nvvm_rsqrt_approx_d;
3984 llvm::IRBuilderBase &builder) {
3985 auto thisOp = cast<NVVM::SqrtOp>(op);
3986 Type t = thisOp.getRes().getType();
3987 NVVM::FPRoundingMode rndMode = thisOp.getRnd();
3988 bool isFtz = thisOp.getFtz();
3992 unsigned rndIndex =
static_cast<unsigned>(rndMode) - 1;
3994 static constexpr llvm::Intrinsic::ID f32IDs[] = {
3995 llvm::Intrinsic::nvvm_sqrt_rn_f,
3996 llvm::Intrinsic::nvvm_sqrt_rm_f,
3997 llvm::Intrinsic::nvvm_sqrt_rp_f,
3998 llvm::Intrinsic::nvvm_sqrt_rz_f,
4000 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
4001 llvm::Intrinsic::nvvm_sqrt_rn_ftz_f,
4002 llvm::Intrinsic::nvvm_sqrt_rm_ftz_f,
4003 llvm::Intrinsic::nvvm_sqrt_rp_ftz_f,
4004 llvm::Intrinsic::nvvm_sqrt_rz_ftz_f,
4006 static constexpr llvm::Intrinsic::ID f64IDs[] = {
4007 llvm::Intrinsic::nvvm_sqrt_rn_d,
4008 llvm::Intrinsic::nvvm_sqrt_rm_d,
4009 llvm::Intrinsic::nvvm_sqrt_rp_d,
4010 llvm::Intrinsic::nvvm_sqrt_rz_d,
4013 llvm::Intrinsic::ID
id =
4014 t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4022 llvm::IRBuilderBase &builder) {
4023 auto thisOp = cast<NVVM::SqrtApproxOp>(op);
4024 llvm::Intrinsic::ID
id = thisOp.getFtz()
4025 ? llvm::Intrinsic::nvvm_sqrt_approx_ftz_f
4026 : llvm::Intrinsic::nvvm_sqrt_approx_f;
4032 llvm::IRBuilderBase &builder) {
4033 auto thisOp = cast<NVVM::DivFOp>(op);
4034 bool isFtz = thisOp.getFtz();
4036 llvm::Intrinsic::ID id;
4038 if (thisOp.getApprox()) {
4039 id = isFtz ? llvm::Intrinsic::nvvm_div_approx_ftz_f
4040 : llvm::Intrinsic::nvvm_div_approx_f;
4041 }
else if (thisOp.getFull()) {
4044 id = isFtz ? llvm::Intrinsic::nvvm_div_full_ftz
4045 : llvm::Intrinsic::nvvm_div_full;
4048 unsigned rndIndex =
static_cast<unsigned>(thisOp.getRnd()) - 1;
4050 static constexpr llvm::Intrinsic::ID f32IDs[] = {
4051 llvm::Intrinsic::nvvm_div_rn_f,
4052 llvm::Intrinsic::nvvm_div_rm_f,
4053 llvm::Intrinsic::nvvm_div_rp_f,
4054 llvm::Intrinsic::nvvm_div_rz_f,
4056 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
4057 llvm::Intrinsic::nvvm_div_rn_ftz_f,
4058 llvm::Intrinsic::nvvm_div_rm_ftz_f,
4059 llvm::Intrinsic::nvvm_div_rp_ftz_f,
4060 llvm::Intrinsic::nvvm_div_rz_ftz_f,
4062 static constexpr llvm::Intrinsic::ID f64IDs[] = {
4063 llvm::Intrinsic::nvvm_div_rn_d,
4064 llvm::Intrinsic::nvvm_div_rm_d,
4065 llvm::Intrinsic::nvvm_div_rp_d,
4066 llvm::Intrinsic::nvvm_div_rz_d,
4068 Type t = thisOp.getRes().getType();
4069 id = t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4080 auto thisOp = cast<NVVM::AsyncStoreGlobalOp>(op);
4081 mlir::NVVM::MemScopeKind scope = thisOp.getScope();
4082 bool isMmio = thisOp.getMmio();
4084 llvm::Value *addr = mt.
lookupValue(thisOp.getAddr());
4085 llvm::Value *value = mt.
lookupValue(thisOp.getValue());
4086 llvm::Value *isMultimem = builder.getInt1(thisOp.getMultimem());
4088 if (scope == MemScopeKind::SYS) {
4089 return isMmio ?
IDArgPair(llvm::Intrinsic::nvvm_st_async_mmio_sys,
4092 {addr, value, isMultimem});
4093 }
else if (scope == MemScopeKind::GPU) {
4094 return IDArgPair(llvm::Intrinsic::nvvm_st_async_gpu,
4095 {addr, value, isMultimem});
4097 llvm_unreachable(
"unsupported scope for AsyncStoreGlobalOp");
4102 llvm::IRBuilderBase &builder) {
4103 auto thisOp = cast<NVVM::PMEventOp>(op);
4107 llvm::Value *maskVal;
4108 if (
auto eventAttr = thisOp.getEventIdAttr()) {
4109 uint16_t mask =
static_cast<uint16_t
>(1u << eventAttr.getInt());
4110 maskVal = llvm::ConstantInt::get(i16Ty, mask);
4113 llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
4116 return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
4119bool MBarrierInitOp::getAsmValues(
4126 for (
auto val : getOperands())
4134 auto thisOp = cast<NVVM::MBarrierInitOp>(op);
4138 return {llvm::Intrinsic::nvvm_mbarrier_init,
4140 builder.getInt32(thisOp.getLayout())}};
4145 auto thisOp = cast<NVVM::MBarrierInvalOp>(op);
4147 llvm::Intrinsic::ID
id = isShared
4148 ? llvm::Intrinsic::nvvm_mbarrier_inval_shared
4149 : llvm::Intrinsic::nvvm_mbarrier_inval;
4156 auto thisOp = cast<NVVM::MBarrierCheckLayoutOp>(op);
4159 llvm::Intrinsic::nvvm_mbarrier_check_layout,
4160 {mt.
lookupValue(thisOp.getAddr()), builder.getInt32(thisOp.getLayout())}};
4165 auto thisOp = cast<NVVM::MBarrierExpectTxOp>(op);
4168 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4171 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4173 static constexpr llvm::Intrinsic::ID IDs[] = {
4174 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cta,
4175 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cluster,
4176 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cta,
4177 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cluster};
4182 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4184 return {IDs[
index], std::move(args)};
4189 auto thisOp = cast<NVVM::MBarrierCompleteTxOp>(op);
4192 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4195 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4197 static constexpr llvm::Intrinsic::ID IDs[] = {
4198 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cta,
4199 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cluster,
4200 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cta,
4201 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cluster};
4206 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4208 return {IDs[
index], std::move(args)};
4213 auto thisOp = cast<NVVM::MBarrierArriveOp>(op);
4216 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4219 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4221 static constexpr llvm::Intrinsic::ID IDs[] = {
4222 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta,
4223 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cluster,
4224 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cta,
4225 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cluster};
4226 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4227 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cta,
4228 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cluster,
4229 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cta,
4231 nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cluster};
4232 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4236 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4243 bool hasCount =
static_cast<bool>(thisOp.getCount());
4245 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta))
4246 return {llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {mbar}};
4250 llvm::Value *count =
4252 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4253 return {id, {mbar, count}};
4258 auto thisOp = cast<NVVM::MBarrierArriveDropOp>(op);
4261 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4264 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4266 static constexpr llvm::Intrinsic::ID IDs[] = {
4267 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta,
4268 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cluster,
4269 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cta,
4270 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cluster};
4271 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4272 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cta,
4274 nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cluster,
4276 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cta,
4278 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cluster};
4279 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4283 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4290 bool hasCount =
static_cast<bool>(thisOp.getCount());
4292 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta))
4293 return {llvm::Intrinsic::nvvm_mbarrier_arrive_drop_shared, {mbar}};
4297 llvm::Value *count =
4299 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4300 return {id, {mbar, count}};
4303bool MBarrierArriveExpectTxOp::getAsmValues(
4310 for (
auto val : getOperands())
4318 auto thisOp = cast<NVVM::MBarrierArriveExpectTxOp>(op);
4321 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4324 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4327 static constexpr llvm::Intrinsic::ID IDs[] = {
4328 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cta,
4329 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cluster,
4330 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cta,
4331 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cluster};
4332 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4333 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cta,
4334 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cluster,
4335 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cta,
4336 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cluster};
4338 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4341 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4342 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4347 return {id, {mbar, txcount}};
4352 auto thisOp = cast<NVVM::MBarrierArriveDropExpectTxOp>(op);
4355 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4358 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4361 static constexpr llvm::Intrinsic::ID IDs[] = {
4362 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cta,
4363 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cluster,
4364 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cta,
4365 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cluster};
4366 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4367 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cta,
4368 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cluster,
4369 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cta,
4370 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cluster};
4372 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4375 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4376 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4381 return {id, {mbar, txcount}};
4386 auto thisOp = cast<NVVM::MBarrierArriveNocompleteOp>(op);
4388 llvm::Intrinsic::ID
id =
4389 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared
4390 : llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete;
4394 args.push_back(mt.
lookupValue(thisOp.getCount()));
4396 return {id, std::move(args)};
4401 auto thisOp = cast<NVVM::MBarrierArriveDropNocompleteOp>(op);
4403 llvm::Intrinsic::ID
id =
4404 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete_shared
4405 : llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete;
4409 args.push_back(mt.
lookupValue(thisOp.getCount()));
4411 return {id, std::move(args)};
4416 auto thisOp = cast<NVVM::MBarrierTestWaitOp>(op);
4417 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4418 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4421 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isPhaseParity ? 1 : 0);
4424 static constexpr llvm::Intrinsic::ID IDs[] = {
4425 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cta_space_cta,
4426 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cta_space_cta,
4427 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cluster_space_cta,
4428 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cluster_space_cta};
4429 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4430 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cta_space_cta,
4431 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cta_space_cta,
4432 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cluster_space_cta,
4433 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cluster_space_cta};
4435 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4438 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4439 llvm::Value *input = mt.
lookupValue(thisOp.getStateOrPhase());
4444 return {id, {mbar, input}};
4449 auto thisOp = cast<NVVM::MBarrierTryWaitOp>(op);
4450 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4451 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4452 bool hasTicks =
static_cast<bool>(thisOp.getTicks());
4456 size_t index = ((hasTicks ? 1 : 0) << 2) | ((isClusterScope ? 1 : 0) << 1) |
4457 (isPhaseParity ? 1 : 0);
4460 static constexpr llvm::Intrinsic::ID IDs[] = {
4461 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cta_space_cta,
4462 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cta_space_cta,
4463 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cluster_space_cta,
4464 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cluster_space_cta,
4465 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cta_space_cta,
4466 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cta_space_cta,
4467 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cluster_space_cta,
4468 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cluster_space_cta};
4469 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4470 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cta_space_cta,
4471 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cta_space_cta,
4472 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cluster_space_cta,
4473 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cluster_space_cta,
4474 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cta_space_cta,
4475 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cta_space_cta,
4476 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cluster_space_cta,
4477 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cluster_space_cta};
4479 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4482 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4489 args.push_back(mbar);
4490 args.push_back(mt.
lookupValue(thisOp.getStateOrPhase()));
4492 args.push_back(mt.
lookupValue(thisOp.getTicks()));
4494 return {id, std::move(args)};
4499 auto thisOp = cast<NVVM::CpAsyncMBarrierArriveOp>(op);
4502 llvm::Intrinsic::ID id;
4503 if (thisOp.getNoinc()) {
4504 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc_shared
4505 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc;
4507 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_shared
4508 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive;
4516 llvm::IRBuilderBase &builder) {
4517 auto thisOp = cast<NVVM::MovMatrixOp>(op);
4518 return {llvm::Intrinsic::nvvm_movmatrix_sync_aligned_m8n8_trans_b16,
4522#define CP_ASYNC_ID_IMPL(mod, size, suffix) \
4523 llvm::Intrinsic::nvvm_cp_async_##mod##_shared_global_##size##suffix
4525#define GET_CP_ASYNC_ID(mod, size, has_cpsize) \
4526 has_cpsize ? CP_ASYNC_ID_IMPL(mod, size, _s) : CP_ASYNC_ID_IMPL(mod, size, )
4531 llvm::Intrinsic::ID id;
4533 auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
4534 bool hasCpSize =
static_cast<bool>(cpAsyncOp.getCpSize());
4535 switch (cpAsyncOp.getSize()) {
4543 id = (cpAsyncOp.getModifier() == NVVM::LoadCacheModifierKind::CG)
4548 llvm_unreachable(
"Invalid copy size in CpAsyncOp.");
4552 args.push_back(mt.
lookupValue(cpAsyncOp.getDst()));
4553 args.push_back(mt.
lookupValue(cpAsyncOp.getSrc()));
4555 args.push_back(mt.
lookupValue(cpAsyncOp.getCpSize()));
4562 auto thisOp = cast<NVVM::CpAsyncBulkPrefetchOp>(op);
4564 llvm::Intrinsic::ID
id = llvm::Intrinsic::nvvm_cp_async_bulk_prefetch_L2;
4567 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4571 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4572 llvm::Value *i64Unused =
4573 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4574 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4575 args.push_back(builder.getInt1(hasCacheHint));
4577 return {id, std::move(args)};
4582 auto thisOp = cast<NVVM::CpAsyncBulkGlobalToSharedClusterOp>(op);
4586 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4588 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4592 mlir::Value multicastMask = thisOp.getMulticastMask();
4593 const bool hasMulticastMask =
static_cast<bool>(multicastMask);
4596 llvm::Value *i16Unused = llvm::ConstantInt::get(builder.getInt16Ty(), 0);
4597 args.push_back(hasMulticastMask ? mt.
lookupValue(multicastMask)
4600 args.push_back(builder.getInt32(0));
4601 args.push_back(builder.getInt32(0));
4606 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4607 llvm::Value *i64Unused = llvm::ConstantInt::get(builder.getInt64Ty(), 0);
4608 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4612 args.push_back(builder.getInt1(hasCacheHint));
4613 args.push_back(builder.getInt1(
false));
4615 args.push_back(builder.getInt1(hasMulticastMask));
4616 args.push_back(builder.getInt1(hasCacheHint));
4620 args.push_back(builder.getInt32(0));
4622 llvm::Intrinsic::ID
id =
4624 ? llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta
4625 : llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
4627 return {id, std::move(args)};
4632 auto thisOp = cast<NVVM::CpAsyncBulkSharedCTAToGlobalOp>(op);
4634 llvm::Intrinsic::ID
id =
4635 llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global;
4638 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4639 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4643 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4644 llvm::Value *i64Unused =
4645 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4646 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4647 args.push_back(builder.getInt1(hasCacheHint));
4650 if (
mlir::Value byteMask = thisOp.getByteMask()) {
4652 id = llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global_bytemask;
4655 return {id, std::move(args)};
4658bool CpAsyncBulkTensorGlobalToSharedClusterOp::getAsmValues(
4665 for (
auto val : getOperands())
4672CpAsyncBulkTensorGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
4674 auto thisOp = cast<NVVM::CpAsyncBulkTensorGlobalToSharedClusterOp>(op);
4675 const bool isCTAOnly = thisOp.getIsCTAOnly();
4679 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4681 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4691 const bool hasMC =
static_cast<bool>(mcMask);
4692 llvm::Value *i16Zero =
4693 llvm::ConstantInt::get(llvm::Type::getInt16Ty(mt.
getLLVMContext()), 0);
4697 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4698 llvm::Value *i64Zero =
4699 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4705 thisOp.getGroup() ? (
static_cast<int32_t
>(*thisOp.getGroup()) + 1) : 0;
4707 llvm::ConstantInt::get(llvm::Type::getInt32Ty(mt.
getLLVMContext()), val);
4710 llvm::Value *flagValidPattern = builder.getInt32(0);
4714 args.push_back(hasMC ? mt.
lookupValue(mcMask) : i16Zero);
4715 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4716 args.push_back(builder.getInt1(hasMC));
4717 args.push_back(builder.getInt1(hasCacheHint));
4719 args.push_back(flagValidPattern);
4722 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4723 args.push_back(builder.getInt1(hasCacheHint));
4724 args.push_back(flagValidPattern);
4727 constexpr size_t numDims = 5;
4728 constexpr size_t numModes = 5;
4729 using rowTy = std::array<llvm::Intrinsic::ID, numDims + 1>;
4730 using TableTy = std::array<rowTy, numModes>;
4731 static constexpr TableTy IDTable{
4732 {{
notIntrinsic, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d,
4733 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d,
4734 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d,
4735 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d,
4736 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d},
4738 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d,
4739 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d,
4740 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d},
4742 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_3d,
4743 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_4d,
4744 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_5d},
4746 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_3d,
4747 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_4d,
4748 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_5d},
4750 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_gather4_2d}}};
4752 static constexpr TableTy IDTableCTA{
4754 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_1d,
4755 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_2d,
4756 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_3d,
4757 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_4d,
4758 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_5d},
4760 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_3d,
4761 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_4d,
4762 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_5d},
4764 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_3d,
4765 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_4d,
4766 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_5d},
4768 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_3d,
4769 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_4d,
4770 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_5d},
4772 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_gather4_2d}}};
4775 (getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1) &&
4776 (getMaxEnumValForTMALoadMode() == std::size(IDTableCTA) - 1),
4777 "TMALoadModes must match number of rows in IDTable and IDTableCTA");
4778 size_t mode =
static_cast<size_t>(thisOp.getMode());
4779 size_t dim = thisOp.getCoordinates().size();
4780 auto id = isCTAOnly ? IDTableCTA[mode][dim] : IDTable[mode][dim];
4782 "Invalid intrinsic for CpAsyncBulkTensorGlobalToSharedClusterOp.");
4784 return {id, std::move(args)};
4789 auto thisOp = cast<NVVM::CpAsyncBulkTensorPrefetchOp>(op);
4793 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4795 for (
auto v : thisOp.getCoordinates())
4797 for (
auto v : thisOp.getIm2colOffsets())
4801 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4802 llvm::Value *i64Unused =
4803 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4804 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4805 args.push_back(builder.getInt1(hasCacheHint));
4807 const unsigned NI = llvm::Intrinsic::not_intrinsic;
4808 static constexpr llvm::Intrinsic::ID IDTable[][6] = {
4809 {NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_1d,
4810 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_2d,
4811 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_3d,
4812 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_4d,
4813 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_5d},
4815 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_3d,
4816 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_4d,
4817 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_5d},
4819 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_3d,
4820 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_4d,
4821 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_5d},
4823 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_3d,
4824 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_4d,
4825 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_5d},
4826 {NI, NI, NI, NI, NI,
4827 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_gather4_2d}};
4829 static_assert(getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1,
4830 "TMALoadModes must match number of rows in IDTable");
4831 size_t mode =
static_cast<size_t>(thisOp.getMode());
4832 size_t dim = thisOp.getCoordinates().size();
4833 llvm::Intrinsic::ID
id = IDTable[mode][dim];
4834 if (
id == llvm::Intrinsic::not_intrinsic)
4835 llvm_unreachable(
"Invalid intrinsic for CpAsyncBulkTensorPrefetchOp.");
4837 return {id, std::move(args)};
4841CpAsyncBulkTensorSharedCTAToGlobalOp::getIntrinsicIDAndArgs(
4843 auto thisOp = cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp>(op);
4847 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4848 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4850 for (
auto v : thisOp.getCoordinates())
4854 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4855 llvm::Value *i64Unused =
4856 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4857 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4858 args.push_back(builder.getInt1(hasCacheHint));
4860 using namespace llvm::Intrinsic;
4861 const unsigned NI = not_intrinsic;
4862 static constexpr ID IDTable[][6] = {
4863 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_1d,
4864 nvvm_cp_async_bulk_tensor_s2g_tile_2d,
4865 nvvm_cp_async_bulk_tensor_s2g_tile_3d,
4866 nvvm_cp_async_bulk_tensor_s2g_tile_4d,
4867 nvvm_cp_async_bulk_tensor_s2g_tile_5d},
4868 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_3d,
4869 nvvm_cp_async_bulk_tensor_s2g_im2col_4d,
4870 nvvm_cp_async_bulk_tensor_s2g_im2col_5d},
4871 {NI, NI, NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_2d},
4872 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_3d,
4873 nvvm_cp_async_bulk_tensor_s2g_im2col_w_4d,
4874 nvvm_cp_async_bulk_tensor_s2g_im2col_w_5d}};
4876 static_assert(getMaxEnumValForTMAStoreMode() == std::size(IDTable) - 1,
4877 "TMAStoreModes must match number of rows in IDTable");
4878 size_t mode =
static_cast<size_t>(thisOp.getMode());
4879 size_t dim = thisOp.getCoordinates().size();
4880 ID
id = IDTable[mode][dim];
4881 if (
id == llvm::Intrinsic::not_intrinsic)
4883 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOp.");
4885 return {id, std::move(args)};
4889CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
4892 cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp>(op);
4895 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4896 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4897 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
4898 for (
Value v : thisOp.getTensorSize())
4900 for (
Value v : thisOp.getLowerStride())
4902 if (thisOp.getUpperStride())
4903 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
4904 for (
Value v : thisOp.getCoordinates())
4908 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4909 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4910 : builder.getInt64(0));
4911 args.push_back(builder.getInt1(hasCacheHint));
4913 using namespace llvm::Intrinsic;
4914 const unsigned NI = not_intrinsic;
4917 static constexpr ID IDTable[][6] = {
4918 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_1d,
4919 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_2d,
4920 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_3d,
4921 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_4d,
4922 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_5d},
4923 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_3d,
4924 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_4d,
4925 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_5d},
4926 {NI, NI, NI, NI, NI,
4927 nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_override_addr_2d},
4928 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_3d,
4929 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_4d,
4930 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_5d}};
4934 static constexpr ID dimStrideIDTable[] = {
4935 NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_1d,
4936 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_2d,
4937 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_3d,
4938 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_4d,
4939 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_5d};
4942 size_t mode =
static_cast<size_t>(thisOp.getMode());
4943 size_t dim = thisOp.getCoordinates().size();
4944 bool isDimStride = !thisOp.getTensorSize().empty();
4946 assert(mode < std::size(IDTable) &&
4947 "Invalid mode for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4948 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
4949 "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4951 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
4953 intrinsicID != NI &&
4954 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4955 return {intrinsicID, std::move(args)};
4960 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOp>(op);
4963 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4964 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4965 for (
Value v : thisOp.getCoordinates())
4969 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4970 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4971 : builder.getInt64(0));
4972 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
4973 args.push_back(builder.getInt1(hasCacheHint));
4975 using namespace llvm::Intrinsic;
4976 const unsigned NI = not_intrinsic;
4977 static constexpr ID IDTable[][6] = {
4978 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_1d,
4979 nvvm_cp_async_bulk_tensor_reduce_tile_2d,
4980 nvvm_cp_async_bulk_tensor_reduce_tile_3d,
4981 nvvm_cp_async_bulk_tensor_reduce_tile_4d,
4982 nvvm_cp_async_bulk_tensor_reduce_tile_5d},
4983 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_3d,
4984 nvvm_cp_async_bulk_tensor_reduce_im2col_4d,
4985 nvvm_cp_async_bulk_tensor_reduce_im2col_5d},
4986 {NI, NI, NI, NI, NI, NI},
4987 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_3d,
4988 nvvm_cp_async_bulk_tensor_reduce_im2col_w_4d,
4989 nvvm_cp_async_bulk_tensor_reduce_im2col_w_5d}};
4991 size_t mode =
static_cast<size_t>(thisOp.getMode());
4992 size_t dim = thisOp.getCoordinates().size();
4993 assert(mode < std::size(IDTable) &&
4994 "Invalid mode for CpAsyncBulkTensorReduceOp");
4995 assert(dim < std::size(IDTable[mode]) &&
4996 "Invalid dim for CpAsyncBulkTensorReduceOp");
4998 ID intrinsicID = IDTable[mode][dim];
4999 assert(intrinsicID != NI &&
5000 "Invalid intrinsic for CpAsyncBulkTensorReduceOp");
5001 return {intrinsicID, std::move(args)};
5004NVVM::IDArgPair CpAsyncBulkTensorReduceOverrideAddrOp::getIntrinsicIDAndArgs(
5006 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOverrideAddrOp>(op);
5009 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
5010 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
5011 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
5013 for (
Value v : thisOp.getTensorSize())
5015 for (
Value v : thisOp.getLowerStride())
5017 if (thisOp.getUpperStride())
5018 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
5019 for (
Value v : thisOp.getCoordinates())
5023 const bool hasCacheHint =
static_cast<bool>(cacheHint);
5024 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
5025 : builder.getInt64(0));
5026 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
5027 args.push_back(builder.getInt1(hasCacheHint));
5029 using namespace llvm::Intrinsic;
5030 const unsigned NI = not_intrinsic;
5033static constexpr ID IDTable[][6] = {
5034 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_1d,
5035 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_2d,
5036 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_3d,
5037 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_4d,
5038 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_5d},
5039 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_3d,
5040 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_4d,
5041 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_5d},
5042 {NI, NI, NI, NI, NI, NI},
5043 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_3d,
5044 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_4d,
5045 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_5d}};
5049static constexpr ID dimStrideIDTable[] = {
5050 NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_1d,
5051 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_2d,
5052 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_3d,
5053 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_4d,
5054 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_5d};
5057 size_t mode =
static_cast<size_t>(thisOp.getMode());
5058 size_t dim = thisOp.getCoordinates().size();
5059 bool isDimStride = !thisOp.getTensorSize().empty();
5061 assert(mode < std::size(IDTable) &&
5062 "Invalid mode for CpAsyncBulkTensorReduceOverrideAddrOp");
5063 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
5064 "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
5066 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
5067 assert(intrinsicID != NI &&
5068 "Invalid intrinsic for CpAsyncBulkTensorReduceOverrideAddrOp");
5069 return {intrinsicID, std::move(args)};
5074#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5075 hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
5076 : llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
5078#define GET_CVT_F2TF32_ID(rnd, relu, sf) \
5079 hasSatFinite ? CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5080 : CVT_F2TF32_ID_IMPL(rnd, relu, )
5083ConvertFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
5084 NVVM::SaturationMode sat,
bool hasRelu) {
5085 using RndMode = NVVM::FPRoundingMode;
5086 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5095 llvm_unreachable(
"Invalid RoundingMode for CvtFloatToTF32Op");
5100ConvertF32x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF4x2Op op,
5102 llvm::IRBuilderBase &builder) {
5106 args.push_back(builder.getInt1(op.getPzo()));
5109 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5110 static_cast<unsigned>(op.getRelu());
5111 static constexpr llvm::Intrinsic::ID ids[] = {
5112 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_satfinite,
5113 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_relu_satfinite,
5114 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rz_satfinite,
5115 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rz_relu_satfinite,
5118 return {ids[idx], std::move(args)};
5122ConvertF32x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5124 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5125 static_cast<unsigned>(hasRelu);
5126 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5127 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rn_satfinite,
5128 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rn_relu_satfinite,
5129 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rz_satfinite,
5130 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rz_relu_satfinite,
5132 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5133 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rn_satfinite,
5134 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rn_relu_satfinite,
5135 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rz_satfinite,
5136 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rz_relu_satfinite,
5140 .Case([&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5141 .Case([&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5143 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF6x2Op");
5144 return llvm::Intrinsic::not_intrinsic;
5149ConvertF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF16x2ToF4x2Op &op,
5151 llvm::IRBuilderBase &builder) {
5155 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5156 static_cast<unsigned>(op.getRelu());
5157 static constexpr llvm::Intrinsic::ID ids[] = {
5158 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_satfinite,
5159 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_relu_satfinite,
5160 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rz_satfinite,
5161 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rz_relu_satfinite,
5164 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5165 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5170 args.push_back(builder.getInt1(op.getPzo()));
5172 return {intId, std::move(args)};
5176ConvertBF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertBF16x2ToF4x2Op &op,
5178 llvm::IRBuilderBase &builder) {
5182 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5183 static_cast<unsigned>(op.getRelu());
5184 static constexpr llvm::Intrinsic::ID ids[] = {
5185 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_satfinite,
5186 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite,
5187 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rz_satfinite,
5188 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rz_relu_satfinite,
5191 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5192 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5197 args.push_back(builder.getInt1(op.getPzo()));
5199 return {intId, std::move(args)};
5203ConvertF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5205 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5206 static_cast<unsigned>(hasRelu);
5207 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5208 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_satfinite,
5209 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_relu_satfinite,
5210 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rz_satfinite,
5211 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rz_relu_satfinite,
5213 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5214 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_satfinite,
5215 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_relu_satfinite,
5216 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rz_satfinite,
5217 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rz_relu_satfinite,
5221 .Case<mlir::Float6E2M3FNType>(
5222 [&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5223 .Case<mlir::Float6E3M2FNType>(
5224 [&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5226 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF6x2Op");
5227 return llvm::Intrinsic::not_intrinsic;
5232ConvertBF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
5233 NVVM::FPRoundingMode rnd,
bool hasRelu) {
5234 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5235 static_cast<unsigned>(hasRelu);
5236 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5237 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_satfinite,
5238 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite,
5239 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rz_satfinite,
5240 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite,
5242 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5243 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_satfinite,
5244 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite,
5245 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rz_satfinite,
5246 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite,
5250 .Case<mlir::Float6E2M3FNType>(
5251 [&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5252 .Case<mlir::Float6E3M2FNType>(
5253 [&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5255 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF6x2Op");
5256 return llvm::Intrinsic::not_intrinsic;
5261ConvertF32x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5262 NVVM::SaturationMode sat,
bool hasRelu) {
5263 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5264 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5265 static_cast<unsigned>(hasRelu);
5266 unsigned ue8m0Idx = (
static_cast<unsigned>(hasSatFinite) << 1) |
5267 static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RP);
5269 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5270 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rn,
5271 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rn_relu,
5272 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rz,
5273 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rz_relu,
5275 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5276 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rn,
5277 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rn_relu,
5278 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rz,
5279 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rz_relu,
5281 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
5282 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rz,
5283 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rp,
5284 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rz_satfinite,
5285 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rp_satfinite,
5289 .Case([&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5290 .Case([&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5291 .Case([&](mlir::Float8E8M0FNUType) {
5292 if (rnd != NVVM::FPRoundingMode::RZ && rnd != NVVM::FPRoundingMode::RP)
5293 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5294 return ue8m0x2IDs[ue8m0Idx];
5297 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5298 return llvm::Intrinsic::not_intrinsic;
5303ConvertF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5305 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5306 static_cast<unsigned>(hasRelu);
5307 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5308 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rn,
5309 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rn_relu,
5310 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rz,
5311 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rz_relu,
5313 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5314 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rn,
5315 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rn_relu,
5316 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rz,
5317 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rz_relu,
5321 .Case([&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5322 .Case([&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5324 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF8x2Op");
5325 return llvm::Intrinsic::not_intrinsic;
5330ConvertBF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
5331 NVVM::FPRoundingMode rnd,
5332 NVVM::SaturationMode sat,
bool hasRelu) {
5333 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5334 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5335 static_cast<unsigned>(hasRelu);
5336 unsigned ue8m0Idx = (
static_cast<unsigned>(hasSatFinite) << 1) |
5337 static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RP);
5339 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5340 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_satfinite,
5341 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite,
5342 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rz_satfinite,
5343 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite,
5345 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5346 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_satfinite,
5347 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite,
5348 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rz_satfinite,
5349 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite,
5351 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
5352 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz,
5353 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp,
5354 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz_satfinite,
5355 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp_satfinite,
5359 .Case<mlir::Float8E4M3FNType>(
5360 [&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5361 .Case<mlir::Float8E5M2Type>(
5362 [&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5363 .Case<mlir::Float8E8M0FNUType>(
5364 [&](mlir::Float8E8M0FNUType) {
return ue8m0x2IDs[ue8m0Idx]; })
5366 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
5367 return llvm::Intrinsic::not_intrinsic;
5373 auto curOp = cast<NVVM::ConvertF8x2ToF16x2Op>(op);
5375 bool hasRelu = curOp.getRelu();
5377 llvm::Intrinsic::ID intId =
5379 .Case([&](Float8E4M3FNType type) {
5380 return hasRelu ? llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn_relu
5381 : llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn;
5383 .Case([&](Float8E5M2Type type) {
5384 return hasRelu ? llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn_relu
5385 : llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn;
5388 llvm_unreachable(
"Invalid type for ConvertF8x2ToF16x2Op");
5389 return llvm::Intrinsic::not_intrinsic;
5392 llvm::Value *packedI16 =
5393 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5394 llvm::Type::getInt16Ty(builder.getContext()));
5396 return {intId, {packedI16}};
5401 auto curOp = cast<NVVM::ConvertF8x2ToBF16x2Op>(op);
5402 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5403 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5404 bool hasRelu = curOp.getRelu();
5406 static constexpr llvm::Intrinsic::ID E4M3Ids[] = {
5407 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5408 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5409 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5410 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5413 static constexpr llvm::Intrinsic::ID E5M2Ids[] = {
5414 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5415 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5416 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5417 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5420 llvm::Intrinsic::ID intId =
5422 .Case([&](Float8E8M0FNUType type) {
5423 return llvm::Intrinsic::nvvm_ue8m0x2_to_bf16x2;
5425 .Case([&](Float8E4M3FNType type) {
5426 return E4M3Ids[hasSatfinite << 1 | hasRelu];
5428 .Case([&](Float8E5M2Type type) {
5429 return E5M2Ids[hasSatfinite << 1 | hasRelu];
5432 llvm_unreachable(
"Invalid type for ConvertF8x2ToBF16x2Op");
5433 return llvm::Intrinsic::not_intrinsic;
5435 llvm::Value *packedI16 =
5436 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5437 llvm::Type::getInt16Ty(builder.getContext()));
5440 args.push_back(packedI16);
5441 if (!isa<Float8E8M0FNUType>(curOp.getSrcType()))
5444 : builder.getInt16(0x7f7f));
5447 return {intId, std::move(args)};
5452 auto curOp = cast<NVVM::ConvertF6x2ToF16x2Op>(op);
5454 bool hasRelu = curOp.getRelu();
5456 llvm::Intrinsic::ID intId =
5458 .Case([&](Float6E2M3FNType type) {
5459 return hasRelu ? llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn_relu
5460 : llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn;
5462 .Case([&](Float6E3M2FNType type) {
5463 return hasRelu ? llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn_relu
5464 : llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn;
5467 llvm_unreachable(
"Invalid type for ConvertF6x2ToF16x2Op");
5468 return llvm::Intrinsic::not_intrinsic;
5471 llvm::Value *packedI16 =
5472 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5473 llvm::Type::getInt16Ty(builder.getContext()));
5475 return {intId, {packedI16}};
5480 auto curOp = cast<NVVM::ConvertF6x2ToBF16x2Op>(op);
5481 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5482 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5483 bool hasRelu = curOp.getRelu();
5485 static constexpr llvm::Intrinsic::ID E2M3Ids[] = {
5486 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5487 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5488 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5489 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5492 static constexpr llvm::Intrinsic::ID E3M2Ids[] = {
5493 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5494 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5495 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5496 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5499 unsigned idx = (hasSatfinite << 1) | hasRelu;
5500 llvm::Intrinsic::ID intId =
5502 .Case([&](Float6E2M3FNType type) {
return E2M3Ids[idx]; })
5503 .Case([&](Float6E3M2FNType type) {
return E3M2Ids[idx]; })
5505 llvm_unreachable(
"Invalid type for ConvertF6x2ToBF16x2Op");
5506 return llvm::Intrinsic::not_intrinsic;
5509 llvm::Value *packedI16 =
5510 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5511 llvm::Type::getInt16Ty(builder.getContext()));
5514 args.push_back(packedI16);
5521 return {intId, std::move(args)};
5526 auto curOp = cast<NVVM::ConvertF4x2ToF16x2Op>(op);
5528 bool hasRelu = curOp.getRelu();
5530 llvm::Intrinsic::ID intId =
5532 .Case([&](Float4E2M1FNType type) {
5533 return hasRelu ? llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn_relu
5534 : llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn;
5537 llvm_unreachable(
"Invalid type for ConvertF4x2ToF16x2Op");
5538 return llvm::Intrinsic::not_intrinsic;
5541 llvm::Value *extendedI16 =
5542 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5543 llvm::Type::getInt16Ty(builder.getContext()));
5545 return {intId, {extendedI16}};
5550 auto curOp = cast<NVVM::ConvertF4x2ToBF16x2Op>(op);
5551 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5552 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5553 bool hasRelu = curOp.getRelu();
5555 static constexpr llvm::Intrinsic::ID E2M1Ids[] = {
5556 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_scale_n2_ue8m0,
5557 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5558 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5559 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5562 unsigned idx = (hasSatfinite << 1) | hasRelu;
5563 llvm::Intrinsic::ID intId =
5565 .Case([&](Float4E2M1FNType type) {
return E2M1Ids[idx]; })
5567 llvm_unreachable(
"Invalid type for ConvertF4x2ToBF16x2Op");
5568 return llvm::Intrinsic::not_intrinsic;
5571 llvm::Value *extendedI16 =
5572 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5573 llvm::Type::getInt16Ty(builder.getContext()));
5576 args.push_back(extendedI16);
5583 return {intId, std::move(args)};
5588 auto thisOp = cast<NVVM::ConvertF32x2ToS2F6x2Op>(op);
5589 bool hasRelu = thisOp.getRelu();
5590 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5592 llvm::Intrinsic::ID
id =
5594 ? llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5595 : llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5601 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5602 : builder.getInt16(0x7f7f));
5603 return {id, std::move(args)};
5608 auto thisOp = cast<NVVM::ConvertBF16x2ToS2F6x2Op>(op);
5609 bool hasRelu = thisOp.getRelu();
5610 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5612 llvm::Intrinsic::ID
id =
5615 nvvm_bf16x2_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5616 : llvm::Intrinsic::nvvm_bf16x2_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5621 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5622 : builder.getInt16(0x7f7f));
5623 return {id, std::move(args)};
5628 auto thisOp = cast<NVVM::ConvertS2F6x2ToBF16x2Op>(op);
5629 bool hasRelu = thisOp.getRelu();
5630 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5631 bool hasSat = thisOp.getSat() == NVVM::SaturationMode::SATFINITE;
5633 static constexpr llvm::Intrinsic::ID ids[] = {
5634 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_scale_n2_ue8m0,
5635 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5636 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5637 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5640 unsigned idx = (hasSat << 1) | hasRelu;
5644 llvm::Value *packedI16 =
5645 builder.CreateBitCast(mt.
lookupValue(thisOp.getSrc()),
5646 llvm::Type::getInt16Ty(builder.getContext()));
5647 args.push_back(packedI16);
5648 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5649 : builder.getInt16(0x7f7f));
5651 return {ids[idx], std::move(args)};
5656 auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
5657 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5659 llvm::Intrinsic::ID
id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
5660 : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
5664 builder.getInt1(curOp.getIsExclusive())}};
5669 auto curOp = cast<NVVM::Tcgen05DeallocOp>(op);
5670 llvm::Intrinsic::ID
id = (curOp.getGroup() == CTAGroupKind::CTA_1)
5671 ? llvm::Intrinsic::nvvm_tcgen05_dealloc_cg1
5672 : llvm::Intrinsic::nvvm_tcgen05_dealloc_cg2;
5676 builder.getInt1(curOp.getIsExclusive())}};
5680Tcgen05CommitOp::getIntrinsicIDAndArgs(
Operation &op,
5683 auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
5684 bool hasMulticast =
static_cast<bool>(curOp.getMulticastMask());
5685 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5686 bool hasSmemARead = curOp.getSmemARead();
5687 unsigned index = (
static_cast<unsigned>(hasSmemARead) << 1) |
5688 static_cast<unsigned>(is2CTAMode);
5690 using namespace llvm::Intrinsic;
5691 static constexpr ID IDs[] = {
5692 nvvm_tcgen05_commit_cg1,
5693 nvvm_tcgen05_commit_cg2,
5694 nvvm_tcgen05_commit_smem_a_read_cg1,
5695 nvvm_tcgen05_commit_smem_a_read_cg2,
5698 static constexpr ID multicastIDs[] = {
5699 nvvm_tcgen05_commit_mc_cg1,
5700 nvvm_tcgen05_commit_mc_cg2,
5701 nvvm_tcgen05_commit_smem_a_read_mc_cg1,
5702 nvvm_tcgen05_commit_smem_a_read_mc_cg2,
5705 ID
id = hasMulticast ? multicastIDs[
index] : IDs[
index];
5709 args.push_back(mt.
lookupValue(curOp.getMulticastMask()));
5714#define TCGEN05_CP_IMPL(shape_mc, src_fmt, cg) \
5715 llvm::Intrinsic::nvvm_tcgen05_cp##shape_mc##src_fmt##cg
5717#define TCGEN05_CP_2CTA(shape_mc, src_fmt, is_2cta) \
5718 is_2cta ? TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg2) \
5719 : TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
5721#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
5723 if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
5724 return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
5725 if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
5726 return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
5727 return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
5731ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
5733 llvm::IRBuilderBase &builder) {
5734 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5735 llvm::Intrinsic::nvvm_ff2f16x2_rn,
5736 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu,
5737 llvm::Intrinsic::nvvm_ff2f16x2_rn_satfinite,
5738 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu_satfinite,
5740 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5741 llvm::Intrinsic::nvvm_ff2f16x2_rz,
5742 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu,
5743 llvm::Intrinsic::nvvm_ff2f16x2_rz_satfinite,
5744 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu_satfinite,
5746 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5747 llvm::Intrinsic::nvvm_ff2f16x2_rs,
5748 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu,
5749 llvm::Intrinsic::nvvm_ff2f16x2_rs_satfinite,
5750 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu_satfinite,
5753 unsigned hasRelu = op.getRelu() ? 1 : 0;
5754 unsigned hasSatFinite =
5755 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5758 unsigned idx = (hasSatFinite << 1) | hasRelu;
5763 if (op.getRandomBits())
5764 args.push_back(mt.
lookupValue(op.getRandomBits()));
5766 if (op.getRnd() != FPRoundingMode::RS)
5767 args.push_back(builder.getInt1(op.getPzo()));
5769 switch (op.getRnd()) {
5770 case FPRoundingMode::RN:
5771 return {rndRNIds[idx], std::move(args)};
5772 case FPRoundingMode::RZ:
5773 return {rndRZIds[idx], std::move(args)};
5774 case FPRoundingMode::RS:
5775 return {rndRSIds[idx], std::move(args)};
5777 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToF16x2Op");
5782ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
5784 llvm::IRBuilderBase &builder) {
5785 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5786 llvm::Intrinsic::nvvm_ff2bf16x2_rn,
5787 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu,
5788 llvm::Intrinsic::nvvm_ff2bf16x2_rn_satfinite,
5789 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu_satfinite,
5791 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5792 llvm::Intrinsic::nvvm_ff2bf16x2_rz,
5793 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu,
5794 llvm::Intrinsic::nvvm_ff2bf16x2_rz_satfinite,
5795 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu_satfinite,
5797 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5798 llvm::Intrinsic::nvvm_ff2bf16x2_rs,
5799 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu,
5800 llvm::Intrinsic::nvvm_ff2bf16x2_rs_satfinite,
5801 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu_satfinite,
5804 unsigned hasRelu = op.getRelu() ? 1 : 0;
5805 unsigned hasSatFinite =
5806 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5809 unsigned idx = (hasSatFinite << 1) | hasRelu;
5814 if (op.getRandomBits())
5815 args.push_back(mt.
lookupValue(op.getRandomBits()));
5817 if (op.getRnd() != FPRoundingMode::RS)
5818 args.push_back(builder.getInt1(op.getPzo()));
5820 switch (op.getRnd()) {
5821 case FPRoundingMode::RN:
5822 return {rndRNIds[idx], std::move(args)};
5823 case FPRoundingMode::RZ:
5824 return {rndRZIds[idx], std::move(args)};
5825 case FPRoundingMode::RS:
5826 return {rndRSIds[idx], std::move(args)};
5828 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToBF16x2Op");
5832llvm::Intrinsic::ID ConvertF32x4ToF8x4Op::getIntrinsicID() {
5834 bool hasRelu = getRelu();
5837 .Case([&](mlir::Float8E4M3FNType) {
5838 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_relu_satfinite
5839 : llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_satfinite;
5841 .Case([&](mlir::Float8E5M2Type) {
5842 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_relu_satfinite
5843 : llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_satfinite;
5846 llvm_unreachable(
"Invalid F8 type in ConvertF32x4ToF8x4Op");
5847 return llvm::Intrinsic::not_intrinsic;
5851llvm::Intrinsic::ID ConvertF32x4ToF6x4Op::getIntrinsicID() {
5853 bool hasRelu = getRelu();
5856 .Case([&](mlir::Float6E2M3FNType) {
5857 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_relu_satfinite
5858 : llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_satfinite;
5860 .Case([&](mlir::Float6E3M2FNType) {
5861 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_relu_satfinite
5862 : llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_satfinite;
5865 llvm_unreachable(
"Invalid F6 type in ConvertF32x4ToF6x4Op");
5866 return llvm::Intrinsic::not_intrinsic;
5870llvm::Intrinsic::ID ConvertF32x4ToF4x4Op::getIntrinsicID() {
5872 bool hasRelu = getRelu();
5875 .Case([&](mlir::Float4E2M1FNType) {
5876 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_relu_satfinite
5877 : llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_satfinite;
5880 llvm_unreachable(
"Invalid F4 type in ConvertF32x4ToF4x4Op");
5881 return llvm::Intrinsic::not_intrinsic;
5885llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(
Operation &op) {
5886 auto curOp = cast<NVVM::Tcgen05CpOp>(op);
5887 bool is2CTA = curOp.getGroup() == CTAGroupKind::CTA_2;
5888 auto srcFmt = curOp.getSrcFormat();
5889 auto mc = curOp.getMulticast();
5891 switch (curOp.getShape()) {
5892 case Tcgen05CpShape::SHAPE_128x256b:
5894 case Tcgen05CpShape::SHAPE_128x128b:
5896 case Tcgen05CpShape::SHAPE_4x256b:
5898 case Tcgen05CpShape::SHAPE_32x128b:
5900 case Tcgen05CpShape::SHAPE_64x128b:
5901 return (mc == Tcgen05CpMulticast::WARPX2_01_23)
5905 llvm_unreachable(
"Invalid shape in tcgen05 cp Op");
5912 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
5914 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
5919LogicalResult Tcgen05LdOp::verify() {
5921 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5924 if (
getShape() != NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && getOffset())
5925 result =
emitError(
"offset argument is only supported for shape 16x32bx2");
5927 auto resTy = getRes().getType();
5928 unsigned resLen = isa<VectorType>(resTy)
5929 ? llvm::cast<VectorType>(resTy).getNumElements()
5932 result =
emitError(llvm::formatv(
"invalid result type length {0} for shape "
5933 "{1} in tcgen05.ld Op",
5934 resLen, stringifyEnum(
getShape())));
5939LogicalResult Tcgen05StOp::verify() {
5941 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5944 auto valTy = getVal().getType();
5945 unsigned valLen = isa<VectorType>(valTy)
5946 ? llvm::cast<VectorType>(valTy).getNumElements()
5949 result =
emitError(llvm::formatv(
"invalid input length {0} for shape "
5950 "{1} in tcgen05.st Op",
5951 valLen, stringifyEnum(
getShape())));
5959 std::optional<LLVM::ConstantRangeAttr> range,
Value result,
5963 setResultRanges(
result, {range->getLower(), range->getUpper(),
5964 range->getLower(), range->getUpper()});
5973 Operation *op, std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
5977 const llvm::APInt &lower = rangeAttr->getLower();
5978 const llvm::APInt &upper = rangeAttr->getUpper();
5981 if (lower == upper && !lower.isMaxValue() && !lower.isMinValue()) {
5982 unsigned bitWidth = lower.getBitWidth();
5983 llvm::APInt minVal = llvm::APInt::getMinValue(bitWidth);
5984 llvm::APInt maxVal = llvm::APInt::getMaxValue(bitWidth);
5986 "invalid range attribute: Lower == Upper, but they aren't min (")
5987 << llvm::toString(minVal, 10,
false) <<
") or max ("
5988 << llvm::toString(maxVal, 10,
false)
5989 <<
") value! This is an invalid constant range.";
5996 llvm::IRBuilderBase &builder) {
5997 return builder.CreateBitCast(arg,
5998 llvm::Type::getInt32Ty(builder.getContext()));
6003 auto curOp = cast<NVVM::DotAccumulate4WayOp>(op);
6010 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
6011 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
6012 unsigned type = (isASigned << 1) | isBSigned;
6013 const llvm::Intrinsic::ID ids[] = {
6014 llvm::Intrinsic::nvvm_idp4a_u_u,
6015 llvm::Intrinsic::nvvm_idp4a_u_s,
6016 llvm::Intrinsic::nvvm_idp4a_s_u,
6017 llvm::Intrinsic::nvvm_idp4a_s_s,
6019 return {ids[type], args};
6024 auto curOp = cast<NVVM::DotAccumulate2WayOp>(op);
6029 args.push_back(builder.getInt1(curOp.getBHi()));
6032 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
6033 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
6034 unsigned type = (isASigned << 1) | isBSigned;
6035 const llvm::Intrinsic::ID ids[] = {
6036 llvm::Intrinsic::nvvm_idp2a_u_u,
6037 llvm::Intrinsic::nvvm_idp2a_u_s,
6038 llvm::Intrinsic::nvvm_idp2a_s_u,
6039 llvm::Intrinsic::nvvm_idp2a_s_s,
6041 return {ids[type], args};
6045 llvm::IRBuilderBase &builder) {
6046 return builder.CreateAddrSpaceCast(
6047 addr, builder.getPtrTy(llvm::NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM));
6051PrefetchOp::getIntrinsicIDAndArgs(NVVM::PrefetchOp &op,
6053 llvm::IRBuilderBase &builder) {
6054 using MemSpace = NVVM::NVVMMemorySpace;
6055 using CacheLevel = NVVM::PrefetchCacheLevel;
6057 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = op.getCacheLevel();
6058 std::optional<NVVM::CacheEvictionPriority> evictPriority =
6059 op.getEvictPriority();
6060 unsigned addressSpace =
6061 llvm::cast<LLVM::LLVMPointerType>(op.getAddr().getType())
6069 if (op.getTensormap())
6070 return {llvm::Intrinsic::nvvm_prefetch_tensormap, args};
6072 assert(cacheLevel &&
"expected cache level for non-tensormap prefetch");
6074 if (op.getUniform() && *cacheLevel == CacheLevel::L1)
6075 return {llvm::Intrinsic::nvvm_prefetchu_L1, args};
6077 if (evictPriority && *cacheLevel == CacheLevel::L2) {
6078 switch (*evictPriority) {
6079 case NVVM::CacheEvictionPriority::EvictLast:
6080 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_last, args};
6081 case NVVM::CacheEvictionPriority::EvictNormal:
6082 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_normal, args};
6084 llvm_unreachable(
"Invalid cache eviction priority");
6088 switch (
static_cast<MemSpace
>(addressSpace)) {
6089 case MemSpace::Generic:
6090 return *cacheLevel == CacheLevel::L1
6092 :
NVVM::
IDArgPair({llvm::Intrinsic::nvvm_prefetch_L2, args});
6093 case MemSpace::Global:
6094 return *cacheLevel == CacheLevel::L1
6096 {llvm::Intrinsic::nvvm_prefetch_global_L1, args})
6098 {llvm::Intrinsic::nvvm_prefetch_global_L2, args});
6099 case MemSpace::Local:
6100 return *cacheLevel == CacheLevel::L1
6102 {llvm::Intrinsic::nvvm_prefetch_local_L1, args})
6104 {llvm::Intrinsic::nvvm_prefetch_local_L2, args});
6106 llvm_unreachable(
"Invalid pointer address space");
6110bool NVVM::InlinePtxOp::getAsmValues(
6114 for (
auto arg : getReadWriteArgs())
6116 for (
auto arg : getResults())
6118 for (
auto arg : getReadOnlyArgs())
6125NVVM::IDArgPair ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
6127 auto curOp = cast<NVVM::ClusterLaunchControlTryCancelOp>(op);
6129 args.push_back(mt.
lookupValue(curOp.getSmemAddress()));
6130 args.push_back(mt.
lookupValue(curOp.getMbarrier()));
6132 llvm::Intrinsic::ID intrinsicID =
6133 curOp.getMulticast()
6135 nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
6136 : llvm::Intrinsic::nvvm_clusterlaunchcontrol_try_cancel_async_shared;
6138 return {intrinsicID, args};
6141NVVM::IDArgPair ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
6143 auto curOp = cast<NVVM::ClusterLaunchControlQueryCancelOp>(op);
6145 args.push_back(mt.
lookupValue(curOp.getTryCancelResponse()));
6147 llvm::Intrinsic::ID intrinsicID;
6149 switch (curOp.getQueryType()) {
6150 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
6152 llvm::Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled;
6154 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
6155 intrinsicID = llvm::Intrinsic::
6156 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x;
6158 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
6159 intrinsicID = llvm::Intrinsic::
6160 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y;
6162 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
6163 intrinsicID = llvm::Intrinsic::
6164 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z;
6167 return {intrinsicID, args};
6172 llvm::IRBuilderBase &builder) {
6173 auto thisOp = cast<NVVM::PermuteOp>(op);
6174 NVVM::PermuteMode mode = thisOp.getMode();
6176 static constexpr llvm::Intrinsic::ID IDs[] = {
6177 llvm::Intrinsic::nvvm_prmt, llvm::Intrinsic::nvvm_prmt_f4e,
6178 llvm::Intrinsic::nvvm_prmt_b4e, llvm::Intrinsic::nvvm_prmt_rc8,
6179 llvm::Intrinsic::nvvm_prmt_ecl, llvm::Intrinsic::nvvm_prmt_ecr,
6180 llvm::Intrinsic::nvvm_prmt_rc16};
6182 unsigned modeIndex =
static_cast<unsigned>(mode);
6190 args.push_back(mt.
lookupValue(thisOp.getSelector()));
6192 return {IDs[modeIndex], args};
6197 auto thisOp = cast<NVVM::TensormapReplaceOp>(op);
6201 if (thisOp.getOrd())
6202 args.push_back(builder.getInt32(thisOp.getOrd().value()));
6203 if (thisOp.getNewValue())
6204 args.push_back(mt.
lookupValue(thisOp.getNewValue()));
6205 if (
auto attr = thisOp.getNewValueAttr()) {
6208 .Case<TensormapElemtypeAttr, TensormapInterleaveLayoutAttr,
6209 TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
6210 TensormapFillModeAttr>([](
auto attr) {
6211 return static_cast<unsigned>(attr.getValue());
6213 .Default([](
auto attr) {
6214 llvm_unreachable(
"Invalid attribute type");
6217 args.push_back(builder.getInt32(val));
6220 static constexpr llvm::Intrinsic::ID IDs[] = {
6221 llvm::Intrinsic::nvvm_tensormap_replace_global_address,
6222 llvm::Intrinsic::nvvm_tensormap_replace_rank,
6223 llvm::Intrinsic::nvvm_tensormap_replace_box_dim,
6224 llvm::Intrinsic::nvvm_tensormap_replace_global_dim,
6225 llvm::Intrinsic::nvvm_tensormap_replace_global_stride,
6226 llvm::Intrinsic::nvvm_tensormap_replace_element_stride,
6227 llvm::Intrinsic::nvvm_tensormap_replace_elemtype,
6228 llvm::Intrinsic::nvvm_tensormap_replace_interleave_layout,
6229 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_mode,
6230 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_atomicity,
6231 llvm::Intrinsic::nvvm_tensormap_replace_fill_mode,
6234 unsigned fieldIndex =
static_cast<unsigned>(thisOp.getField());
6236 return {IDs[fieldIndex], args};
6243static llvm::nvvm::Tcgen05MMAKind
6246 case NVVM::Tcgen05MMAKind::F16:
6247 return llvm::nvvm::Tcgen05MMAKind::F16;
6248 case NVVM::Tcgen05MMAKind::TF32:
6249 return llvm::nvvm::Tcgen05MMAKind::TF32;
6250 case NVVM::Tcgen05MMAKind::F8F6F4:
6251 return llvm::nvvm::Tcgen05MMAKind::F8F6F4;
6252 case NVVM::Tcgen05MMAKind::I8:
6253 return llvm::nvvm::Tcgen05MMAKind::I8;
6254 case NVVM::Tcgen05MMAKind::TI16:
6255 return llvm::nvvm::Tcgen05MMAKind::TI16;
6256 case NVVM::Tcgen05MMAKind::MXF8F6F4:
6257 case NVVM::Tcgen05MMAKind::MXF4:
6258 case NVVM::Tcgen05MMAKind::MXF4NVF4:
6261 llvm_unreachable(
"Unsupported tcgen05.mma kind");
6267 llvm::IRBuilderBase &builder) {
6269 auto thisOp = cast<NVVM::Tcgen05MMAOp>(op);
6272 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6275 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6278 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6279 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6280 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6282 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6283 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6284 using IsATensorArray = std::array<CtaGroupArray, 2>;
6285 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6286 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6289 static constexpr HasDisableOutputLaneArray tcgen05MMAIDs = {
6295 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic},
6297 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic}}},
6301 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6302 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6306 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6307 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6313 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic},
6315 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic}}},
6319 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6320 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6324 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6325 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6331 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1,
6334 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2,
6339 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1,
6341 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift,
6346 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2,
6348 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift,
6354 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1,
6358 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2,
6363 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1,
6365 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift},
6369 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2,
6371 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift,
6374 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6375 bool hasScaleInputD = ScaleInputD !=
nullptr;
6377 llvm::Value *DisableOutputLane =
6379 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6381 const unsigned ctaGroup =
6384 llvm::Intrinsic::ID ID =
6385 tcgen05MMAIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6386 [ctaGroup - 1][thisOp.getAShift()];
6388 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMAOp.");
6391 args.push_back(ScaleInputD);
6393 if (hasDisableOutputLane)
6394 args.push_back(DisableOutputLane);
6396 args.push_back(builder.getInt32(
6399 if (!hasDisableOutputLane)
6400 args.push_back(builder.getInt32(ctaGroup));
6403 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6406 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6413 NVVM::CTAGroupKind ctaGroup,
bool hasAShift,
6414 NVVM::Tcgen05MMACollectorOp collectorOp,
Location loc) {
6416 if (disableOutputLane) {
6417 mlir::VectorType disableOutputLaneType =
6418 cast<mlir::VectorType>(disableOutputLane.
getType());
6419 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6420 disableOutputLaneType.getNumElements() != 4) ||
6421 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6422 disableOutputLaneType.getNumElements() != 8))
6423 return emitError(loc) <<
"Disable Output Lane of length "
6424 << disableOutputLaneType.getNumElements()
6425 <<
" is incompatible with CtaGroupAttr";
6428 if (hasAShift && !isATensor)
6430 loc,
"A-shift can be applied only when matrix A is in tensor memory");
6432 if (hasAShift ==
true && (collectorOp == Tcgen05MMACollectorOp::FILL ||
6433 collectorOp == Tcgen05MMACollectorOp::USE))
6435 loc,
"Cannot use collector buffer operation fill or use with ashift");
6440LogicalResult Tcgen05MMAOp::verify() {
6442 getDisableOutputLane(), getCtaGroup(), getAShift(),
6443 getCollectorOp(), getLoc());
6453 auto thisOp = cast<NVVM::Tcgen05MMASparseOp>(op);
6456 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6459 bool isATensor = isa<llvm::PointerType>(
A->getType());
6462 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6463 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6464 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6465 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6467 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6468 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6469 using IsATensorArray = std::array<CtaGroupArray, 2>;
6470 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6471 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6474 static constexpr HasDisableOutputLaneArray tcgen05MMASparseIDs = {
6480 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic},
6482 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic}}},
6486 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6487 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6491 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6492 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6498 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6501 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6506 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6507 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6511 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6512 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6519 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1,
6523 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2,
6528 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1,
6530 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift,
6535 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2,
6537 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift,
6543 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1,
6547 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2,
6552 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1,
6554 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift},
6558 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2,
6560 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift,
6563 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6564 bool hasScaleInputD = ScaleInputD !=
nullptr;
6566 llvm::Value *DisableOutputLane =
6568 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6573 llvm::Intrinsic::ID ID =
6574 tcgen05MMASparseIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6575 [ctaGroup - 1][thisOp.getAShift()];
6577 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMASparseOp.");
6580 args.push_back(ScaleInputD);
6582 if (hasDisableOutputLane)
6583 args.push_back(DisableOutputLane);
6585 args.push_back(builder.getInt32(
6588 if (!hasDisableOutputLane)
6589 args.push_back(builder.getInt32(ctaGroup));
6592 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6595 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6600LogicalResult Tcgen05MMASparseOp::verify() {
6602 getDisableOutputLane(), getCtaGroup(), getAShift(),
6603 getCollectorOp(), getLoc());
6613 auto thisOp = cast<NVVM::Tcgen05MMABlockScaleOp>(op);
6616 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6619 bool isATensor = isa<llvm::PointerType>(
A->getType());
6622 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6623 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6624 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6625 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6626 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6627 args.push_back(builder.getInt32(
6630 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6632 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6634 auto kind = thisOp.getKind();
6635 auto blockScale = thisOp.getBlockScale();
6636 llvm::Intrinsic::ID ID = [&]() {
6637 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6638 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6639 return isATensor ? llvm::Intrinsic::
6640 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale
6642 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale;
6643 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6646 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32
6648 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32;
6650 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6651 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6653 ? llvm::Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale
6654 : llvm::Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale;
6655 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6656 return isATensor ? llvm::Intrinsic::
6657 nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32
6659 nvvm_tcgen05_mma_shared_mxf4_block_scale_block32;
6661 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6662 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6665 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32
6667 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32;
6669 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6672 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16
6674 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16;
6677 llvm_unreachable(
"Invalid tcgen05.mma.block_scale attributes");
6684 NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind,
6685 NVVM::Tcgen05MMABlockScale blockScale,
Location loc) {
6686 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT &&
6687 kind == NVVM::Tcgen05MMAKind::MXF4NVF4)
6688 return emitError(loc,
"mxf4nvf4 requires block scale attribute");
6690 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16 &&
6691 kind != NVVM::Tcgen05MMAKind::MXF4NVF4)
6693 llvm::formatv(
"{} kind does not support block16 attribute",
6694 stringifyEnum(kind)));
6699LogicalResult Tcgen05MMABlockScaleOp::verify() {
6701 getBlockScale(), getLoc());
6711 auto thisOp = cast<NVVM::Tcgen05MMASparseBlockScaleOp>(op);
6714 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6717 bool isATensor = isa<llvm::PointerType>(
A->getType());
6720 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6721 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6722 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6723 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6724 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6725 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6726 args.push_back(builder.getInt32(
6729 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6731 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6733 auto kind = thisOp.getKind();
6734 auto blockScale = thisOp.getBlockScale();
6735 llvm::Intrinsic::ID ID = [&]() {
6736 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6737 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6738 return isATensor ? llvm::Intrinsic::
6739 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale
6741 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale;
6742 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6745 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32
6747 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32;
6749 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6750 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6751 return isATensor ? llvm::Intrinsic::
6752 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale
6754 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale;
6755 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6758 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32
6760 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32;
6762 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6763 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6766 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32
6768 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32;
6770 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6773 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16
6775 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16;
6778 llvm_unreachable(
"Invalid tcgen05.mma.sp.block_scale attributes");
6784LogicalResult Tcgen05MMASparseBlockScaleOp::verify() {
6786 getBlockScale(), getLoc());
6796 auto thisOp = cast<NVVM::Tcgen05MMAWsOp>(op);
6799 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6802 bool isATensor = isa<llvm::PointerType>(
A->getType());
6805 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6806 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6807 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6809 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6813 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor_zero_col_mask
6814 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared_zero_col_mask;
6816 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor
6817 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared;
6819 args.push_back(builder.getInt32(
6822 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6824 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6836 auto thisOp = cast<NVVM::Tcgen05MMAWsSparseOp>(op);
6839 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6842 bool isATensor = isa<llvm::PointerType>(
A->getType());
6845 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6846 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6847 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6848 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6850 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6855 ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor_zero_col_mask
6856 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared_zero_col_mask;
6858 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor
6859 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared;
6861 args.push_back(builder.getInt32(
6864 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6866 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6877 auto thisOp = cast<Tcgen05MMADecompressBOp>(op);
6880 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6883 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6886 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6887 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6888 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6889 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6891 llvm::Value *DisableOutputLane =
6893 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6895 NVVM::CTAGroupKind ctaGroup = thisOp.getCtaGroup();
6897 using namespace llvm::Intrinsic;
6898 ID intrinsicID = not_intrinsic;
6900 if (hasDisableOutputLane) {
6901 if (ctaGroup == NVVM::CTAGroupKind::CTA_1) {
6904 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b
6905 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b;
6906 }
else if (ctaGroup == NVVM::CTAGroupKind::CTA_2) {
6909 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b
6910 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b;
6912 llvm_unreachable(
"Unknown ctaGroup for tcgen05.mma.decompress_b");
6915 intrinsicID = isATensor ? nvvm_tcgen05_mma_tensor_f8f6f4_decompress_b
6916 : nvvm_tcgen05_mma_shared_f8f6f4_decompress_b;
6919 assert(intrinsicID != not_intrinsic &&
6920 "Invalid intrinsic for Tcgen05MMADecompressBOp.");
6922 if (hasDisableOutputLane)
6923 args.push_back(DisableOutputLane);
6929 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6931 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6933 return {intrinsicID, args};
6936LogicalResult Tcgen05MMADecompressBOp::verify() {
6937 mlir::Value disableOutputLane = getDisableOutputLane();
6939 if (disableOutputLane) {
6940 NVVM::CTAGroupKind ctaGroup = getCtaGroup();
6942 mlir::VectorType disableOutputLaneType =
6943 cast<mlir::VectorType>(disableOutputLane.
getType());
6944 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6945 disableOutputLaneType.getNumElements() != 4) ||
6946 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6947 disableOutputLaneType.getNumElements() != 8))
6948 return emitOpError() <<
"Disable Output Lane of length "
6949 << disableOutputLaneType.getNumElements()
6950 <<
" is incompatible with CtaGroupAttr";
6962 auto thisOp = cast<Tcgen05MMABlockScaleDecompressBOp>(op);
6965 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6968 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6971 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6972 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6973 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6974 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6975 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6976 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6977 args.push_back(builder.getInt32(
6980 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6982 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6984 llvm::Intrinsic::ID intrinsicID =
6987 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32_decompress_b
6989 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32_decompress_b;
6991 return {intrinsicID, args};
6998#define TCGEN05LDRED(SHAPE, NUM, TYPE) \
6999 llvm::Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_##NUM##_##TYPE
7003 auto thisOp = cast<NVVM::Tcgen05LdRedOp>(op);
7006 mlir::VectorType VecResTy =
7007 cast<mlir::VectorType>(thisOp.getData().getType());
7008 unsigned Num = VecResTy.getNumElements();
7009 bool IsFloat = thisOp.getRedVal().getType().isF32();
7011 llvm::Intrinsic::ID Shape32x32b[][2] = {
7022 llvm::Intrinsic::ID Shape16x32bx2[][2] = {
7033 NVVM::Tcgen05LdStShape
shape = thisOp.getShape();
7034 unsigned ID = [&]() {
7037 unsigned idx = std::log2(Num);
7039 case NVVM::Tcgen05LdStShape::SHAPE_32X32B:
7040 return Shape32x32b[idx][IsFloat];
7041 case NVVM::Tcgen05LdStShape::SHAPE_16X32BX2:
7042 return Shape16x32bx2[idx][IsFloat];
7044 llvm_unreachable(
"unhandled tcgen05.ld lowering");
7050 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2)
7051 args.push_back(mt.
lookupValue(thisOp.getOffset()));
7054 builder.getInt32(thisOp.getOp() == NVVM::ReductionKind::MIN ? 0 : 1));
7057 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getAbs())));
7058 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getNan())));
7063LogicalResult Tcgen05LdRedOp::verify() {
7064 VectorType data = cast<VectorType>(getData().
getType());
7065 Type redVal = getRedVal().getType();
7067 if (data.getElementType() != redVal)
7069 "type of reduction value and element type of vector data should match");
7071 if (getOp() != NVVM::ReductionKind::MIN &&
7072 getOp() != NVVM::ReductionKind::MAX)
7073 return emitError(
"only min and max reduction kinds are supported");
7075 if (redVal.
isInteger() && (getAbs() || getNan())) {
7076 return emitError(
"abs or nan is only applicable for f32 type");
7088 if (
auto vecTy = dyn_cast<VectorType>(type))
7089 return vecTy.getNumElements();
7095 return llvm::isPowerOf2_32(repFactor) && repFactor <= 64;
7098LogicalResult SparseCompressOp::verify() {
7099 auto dataTy = cast<VectorType>(getData().
getType());
7100 auto compressedTy = cast<VectorType>(getCompressedData().
getType());
7101 if (dataTy.getElementType() != compressedTy.getElementType())
7103 "expects data and compressed data to have the same element type");
7106 unsigned elemSize = dataTy.getElementTypeBitWidth();
7107 unsigned dataBits = dataTy.getNumElements() * elemSize;
7108 unsigned repFactor = dataBits / 64;
7110 return emitOpError(
"expects data to occupy 2 * repFactor 32-bit "
7111 "registers, where repFactor is a power of two up to 64");
7114 if (2 * compressedTy.getNumElements() != dataTy.getNumElements())
7116 "expects compressed data to have half as many elements as data");
7118 unsigned numMetadataRegs =
7119 llvm::divideCeil(repFactor * getIdxSize(), elemSize);
7121 return emitOpError(
"expects metadata to occupy ")
7122 << numMetadataRegs <<
" 32-bit register(s)";
7127LogicalResult SparseDecompressOp::verify() {
7128 auto dataTy = cast<VectorType>(getData().
getType());
7129 auto compressedTy = cast<VectorType>(getCompressedData().
getType());
7130 if (dataTy.getElementType() != compressedTy.getElementType())
7132 "expects data and compressed data to have the same element type");
7136 unsigned idxSize = getIdxSize();
7137 unsigned numTgt = getNumTgt();
7138 unsigned numData = dataTy.getNumElements();
7139 unsigned numCompressed = compressedTy.getNumElements();
7140 unsigned repFactor = numData / numTgt;
7142 return emitOpError(
"expects data to have repFactor * num_tgt elements, "
7143 "where repFactor is a power of two up to 64");
7145 unsigned numSrc = numCompressed / repFactor;
7146 if (numCompressed % repFactor != 0 ||
7147 !llvm::is_contained({1u, 2u, 4u}, numSrc) || numSrc >= numTgt)
7148 return emitOpError(
"unsupported sparsity factor: compressed data has ")
7149 << numCompressed <<
" elements for " << repFactor <<
" group(s) of "
7150 << numTgt <<
" elements";
7152 unsigned elemSize = dataTy.getElementTypeBitWidth();
7153 unsigned dataBits = numData * elemSize;
7154 if (dataBits < 32 || dataBits > 4096)
7155 return emitOpError(
"expects data to be between 32 and 4096 bits");
7157 if (numSrc * elemSize > 32)
7158 return emitOpError(
"expects at most 32 bits of compressed data per group");
7160 if (idxSize == 2 && numTgt > 4)
7161 return emitOpError(
"expects num_tgt to be at most 4 if idx_size is 2");
7163 unsigned numMetadataRegs = llvm::divideCeil(numSrc * idxSize * repFactor, 32);
7165 return emitOpError(
"expects metadata to occupy ")
7166 << numMetadataRegs <<
" 32-bit register(s)";
7176struct NVVMInlinerInterface final : DialectInlinerInterface {
7177 using DialectInlinerInterface::DialectInlinerInterface;
7178 bool isLegalToInline(Operation *, Region *,
bool, IRMapping &)
const final {
7185void NVVMDialect::initialize() {
7186 registerNVVMDialectOperations(
this);
7187 registerAttributes();
7191 allowUnknownOperations();
7192 addInterfaces<NVVMInlinerInterface>();
7193 declarePromisedInterface<ConvertToLLVMPatternInterface, NVVMDialect>();
7194 declarePromisedInterface<gpu::TargetAttrInterface, NVVMTargetAttr>();
7197LogicalResult NVVMDialect::verifyOperationAttribute(
Operation *op,
7199 StringAttr attrName = attr.
getName();
7201 if (attrName == NVVMDialect::getKernelFuncAttrName()) {
7202 if (!isa<LLVM::LLVMFuncOp>(op)) {
7203 return op->
emitError() <<
"'" << NVVMDialect::getKernelFuncAttrName()
7204 <<
"' attribute attached to unexpected op";
7209 if (attrName == NVVMDialect::getMaxntidAttrName() ||
7210 attrName == NVVMDialect::getReqntidAttrName() ||
7211 attrName == NVVMDialect::getClusterDimAttrName()) {
7212 auto values = llvm::dyn_cast<DenseI32ArrayAttr>(attr.
getValue());
7213 if (!values || values.empty() || values.size() > 3) {
7216 <<
"' attribute must be integer array with maximum 3 index";
7221 if (attrName == NVVMDialect::getMinctasmAttrName() ||
7222 attrName == NVVMDialect::getMaxnregAttrName() ||
7223 attrName == NVVMDialect::getClusterMaxBlocksAttrName()) {
7224 if (!llvm::dyn_cast<IntegerAttr>(attr.
getValue())) {
7226 <<
"'" << attrName <<
"' attribute must be integer constant";
7230 if (attrName == NVVMDialect::getBlocksAreClustersAttrName()) {
7234 <<
"'" << attrName <<
"' attribute must be used along with " <<
"'"
7235 << NVVMDialect::getReqntidAttrName() <<
"' and " <<
"'"
7236 << NVVMDialect::getClusterDimAttrName() <<
"'";
7243LogicalResult NVVMDialect::verifyRegionArgAttribute(
Operation *op,
7244 unsigned regionIndex,
7247 auto funcOp = dyn_cast<FunctionOpInterface>(op);
7252 StringAttr attrName = argAttr.
getName();
7253 if (attrName == NVVM::NVVMDialect::getGridConstantAttrName()) {
7257 <<
"' attribute must be present only on kernel arguments";
7259 if (!isa<UnitAttr>(argAttr.
getValue()))
7260 return op->
emitError() <<
"'" << attrName <<
"' must be a unit attribute";
7261 if (!funcOp.getArgAttr(argIndex, LLVM::LLVMDialect::getByValAttrName())) {
7264 <<
"' attribute requires the argument to also have attribute '"
7265 << LLVM::LLVMDialect::getByValAttrName() <<
"'";
static bool isLegalToInline(InlinerInterface &interface, Region *src, Region *insertRegion, bool shouldCloneInlinedRegion, IRMapping &valueMapping)
Utility to check that all of the operations within 'src' can be inlined.
#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta)
static LogicalResult verifyTMALoadParams(size_t tensorDims, size_t numIm2colOff, TMALoadMode mode, Location loc)
static LogicalResult verifyTcgen05MMAOp(bool isATensor, mlir::Value disableOutputLane, NVVM::CTAGroupKind ctaGroup, bool hasAShift, NVVM::Tcgen05MMACollectorOp collectorOp, Location loc)
static LogicalResult verifyFloatBinaryOp(OpType op, StringRef opKind)
static bool isPtrInAddrSpace(mlir::Value ptr, NVVMMemorySpace targetAS)
static bool isCompatibleReturnTypesOptionalResult(TypeRange inferred, TypeRange actual)
For ops with optional results, allow the user to omit the result even when inference would produce on...
static LogicalResult verifyConvertF32x2ToFP16x2Op(Twine dstType, FPRoundingMode rnd, bool hasRandomBits, bool hasPzo, Operation *op)
static bool isPtrInSharedCTASpace(mlir::Value ptr)
static LogicalResult isAllowedSizeN(int sizeN, NVVM::WGMMATypes typeA)
static llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(NVVM::CTAGroupKind ctaGroup)
static void addInferredMultiplicandTypes(MLIRContext *ctx, OperationState &result, ValueRange operandA, ValueRange operandB, std::optional< std::array< MMATypes, 2 > > multiplicandPtxTypes)
#define GET_CVT_F2TF32_ID(rnd, relu, sf)
static void addBlockScaleProperties(OpBuilder &builder, OperationState &result, ArrayRef< int64_t > shape, ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat, MMABlockScaleKind kind)
static ParseResult parseEnumKeyword(OpAsmParser &parser, AttrTy &attr)
static llvm::nvvm::Tcgen05MMAKind getNVVMTcgen05MMAKind(NVVM::Tcgen05MMAKind kind)
static llvm::Value * getParamCastedAddr(llvm::Value *addr, llvm::IRBuilderBase &builder)
static LogicalResult verifyTcgen05MMABlockScaleOp(NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind, NVVM::Tcgen05MMABlockScale blockScale, Location loc)
static void printMmaUnitProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword)
static bool isValidSPRepFactor(unsigned repFactor)
static llvm::Value * packValInto64Bits(llvm::IRBuilderBase &builder, llvm::Value *result, llvm::Value *field, unsigned sizeInBits, unsigned start)
Packs the given field into the result.
static void printOperandList(OpAsmPrinter &p, StringRef name, ArrayRef< Value > operands)
static void printMmaProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword, AttrTy value)
static bool isMmaPropertyName(StringRef name)
static llvm::Value * getAsPackedI32(llvm::Value *arg, llvm::IRBuilderBase &builder)
static unsigned getNumSPRegisters(Type type)
static void printMmaEnumProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword, AttrTy value)
static LogicalResult verifyMBarrierArriveLikeOp(Operation *op, Value addr, NVVM::MemScopeKind scope, Value retVal=nullptr)
static llvm::Value * castPtrToAddrSpace(llvm::IRBuilderBase &builder, llvm::Value *ptr, NVVMMemorySpace targetAS)
static LogicalResult isAllowedWGMMADataType(NVVM::WGMMATypes typeD, NVVM::WGMMATypes typeA, NVVM::WGMMATypes typeB)
static llvm::Intrinsic::ID getBarrierReductionIntrinsic(bool aligned, NVVM::BarrierReduction kind)
Maps the (aligned, kind) pair to the @llvm.nvvm.barrier.cta.red.
static ParseResult parseMmaProperties(OpAsmParser &parser, NamedAttrList &attributes, ArrayRef< StringRef > allowedKeywords, ArrayRef< StringRef > requiredProperties)
static void inferAndSetMultiplicandTypes(MLIRContext *ctx, NamedAttrList &attrs, const SmallVectorImpl< Type > &operandTypes)
static LogicalResult parseMmaOperand(OpAsmParser &parser, StringRef operandName, SmallVectorImpl< OpAsmParser::UnresolvedOperand > ®s)
static std::pair< mlir::Type, unsigned > inferMMATypeFromMNK(NVVM::MMATypes type, NVVM::MMAFrag frag, int m, int n, int k, MLIRContext *context)
static bool isInt8PtxType(MMATypes type)
#define TCGEN05LDRED(SHAPE, NUM, TYPE)
static bool isInt4PtxType(MMATypes type)
static bool isIntegerPtxType(MMATypes type)
static MMATypes inferPtxTypeFromResult(OpTy op)
static LogicalResult parseMmaTypeSignature(OpAsmParser &parser, SmallVectorImpl< Type > &operandTypes)
static FailureOr< int > getAllowedSizeK(NVVM::WGMMATypes typeA)
static bool isPtrInSharedClusterSpace(mlir::Value ptr)
LogicalResult CpAsyncBulkTensorOverrideAddrCommonVerifier(OperandRange coordinates, OperandRange tensorSize, OperandRange lowerStride, Value upperStride, bool isTile, Location loc)
static ParseResult parseMmaEnumPropertyValue(OpAsmParser &parser, NamedAttrList &attributes, StringRef name)
#define GET_CP_ASYNC_ID(mod, size, has_cpsize)
static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape shape, unsigned vecLen)
static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims, bool isIm2Col, size_t numIm2ColOffsets, Location loc)
static ParseResult parseMmaPropertyValue(OpAsmParser &parser, NamedAttrList &attributes, StringRef name)
static bool isPtrInGenericSpace(mlir::Value ptr)
static void processOperandFragments(Op &op, std::array< MMAOperandFragment, 3 > &frags, SmallVectorImpl< Type > ®Types, SmallVectorImpl< StringRef > &ignoreAttrNames)
static llvm::Intrinsic::ID getBarrierSyncIntrinsic(bool aligned, bool hasCount)
Maps the (aligned, hasCount) pair to the @llvm.nvvm.barrier.cta.sync.
static constexpr unsigned notIntrinsic
static LogicalResult inferMBarrierArriveResultTypes(MLIRContext *context, Value addr, SmallVectorImpl< Type > &inferredReturnTypes)
Only shared_cluster (ptr<7>) produces zero results; all other address spaces (including generic) retu...
static ArrayRef< int64_t > getShape(Type type)
Returns the shape of the given type.
@ OptionalSquare
Square brackets supporting zero or more ops, or nothing.
virtual Builder & getBuilder() const =0
Return a builder which provides useful access to MLIRContext, global objects like types and attribute...
virtual ParseResult parseCommaSeparatedList(Delimiter delimiter, function_ref< ParseResult()> parseElementFn, StringRef contextMessage=StringRef())=0
Parse a list of comma-separated items with an optional delimiter.
virtual ParseResult parseOptionalAttrDict(NamedAttrList &result)=0
Parse a named dictionary into 'result' if it is present.
MLIRContext * getContext() const
virtual ParseResult parseRParen()=0
Parse a ) token.
virtual InFlightDiagnostic emitError(SMLoc loc, const Twine &message={})=0
Emit a diagnostic at the specified location and return failure.
ParseResult parseKeywordOrString(std::string *result)
Parse a keyword or a quoted string.
virtual ParseResult parseCustomAttributeWithFallback(Attribute &result, Type type, function_ref< ParseResult(Attribute &result, Type type)> parseAttribute)=0
Parse a custom attribute with the provided callback, unless the next token is #, in which case the ge...
virtual ParseResult parseEqual()=0
Parse a = token.
virtual SMLoc getCurrentLocation()=0
Get the location of the next token and store it into the argument.
virtual ParseResult parseOptionalComma()=0
Parse a , token if present.
virtual ParseResult parseColon()=0
Parse a : token.
virtual SMLoc getNameLoc() const =0
Return the location of the original name token.
virtual ParseResult parseArrow()=0
Parse a '->' token.
virtual ParseResult parseLParen()=0
Parse a ( token.
virtual ParseResult parseType(Type &result)=0
Parse a type.
virtual ParseResult parseArrowTypeList(SmallVectorImpl< Type > &result)=0
Parse an arrow followed by a type list.
ParseResult parseTypeList(SmallVectorImpl< Type > &result)
Parse a type list.
ParseResult parseKeyword(StringRef keyword)
Parse a given keyword.
void printArrowTypeList(TypeRange &&types)
void printStrippedAttrOrType(AttrOrType attrOrType)
Print the provided attribute in the context of an operation custom printer/parser: this will invoke d...
This class is a general helper class for creating context-global objects like types,...
DenseI32ArrayAttr getDenseI32ArrayAttr(ArrayRef< int32_t > values)
IntegerType getIntegerType(unsigned width)
MLIRContext * getContext() const
Attr getAttr(Args &&...args)
Get or construct an instance of the attribute Attr with provided arguments.
static IntegerValueRange getMaxRange(Value value)
Create a maximal range ([0, uint_max(t)] / [int_min(t), int_max(t)]) range that is used to mark the v...
Implementation class for module translation.
llvm::Value * lookupValue(Value value) const
Finds an LLVM IR value corresponding to the given MLIR value.
void mapValue(Value mlir, llvm::Value *llvm)
Stores the mapping between an MLIR value and its LLVM IR counterpart.
llvm::LLVMContext & getLLVMContext() const
Returns the LLVM context in which the IR is being constructed.
This class defines the main interface for locations in MLIR and acts as a non-nullable wrapper around...
MLIRContext is the top-level object for a collection of MLIR operations.
NamedAttrList is array of NamedAttributes that tracks whether it is sorted and does some basic work t...
std::optional< NamedAttribute > getNamed(StringRef name) const
Return the specified named attribute if present, std::nullopt otherwise.
Attribute get(StringAttr name) const
Return the specified attribute if present, null otherwise.
void append(StringRef name, Attribute attr)
Add an attribute with the specified name.
Attribute set(StringAttr name, Attribute value)
If the an attribute exists with the specified name, change it to the new value.
NamedAttribute represents a combination of a name and an Attribute value.
StringAttr getName() const
Return the name of the attribute.
Attribute getValue() const
Return the value of the attribute.
The OpAsmParser has methods for interacting with the asm parser: parsing things from it,...
ParseResult resolveOperands(Operands &&operands, Type type, SmallVectorImpl< Value > &result)
Resolve a list of operands to SSA values, emitting an error on failure, or appending the results to t...
virtual ParseResult parseOperandList(SmallVectorImpl< UnresolvedOperand > &result, Delimiter delimiter=Delimiter::None, bool allowResultNumber=true, int requiredOperandCount=-1)=0
Parse zero or more SSA comma-separated operand references with a specified surrounding delimiter,...
This is a pure-virtual base class that exposes the asmprinter hooks necessary to implement a custom p...
void printOperands(OperandRange operands)
Print a comma separated range of operation operands out of line to avoid instantiating the range iter...
virtual void printOptionalAttrDict(ArrayRef< NamedAttribute > attrs, ArrayRef< StringRef > elidedAttrs={})=0
If the specified operation has attributes, print out an attribute dictionary with their values.
This class helps build Operations.
This provides public APIs that all operations should have.
This class implements the operand iterators for the Operation class.
Operation is the basic unit of execution within MLIR.
InFlightDiagnostic emitError(const Twine &message={})
Emit an error about fatal conditions with this operation, reporting up to any diagnostic handlers tha...
bool hasDiscardableAttr(StringRef name)
Return true if this operation has a discardable attribute with the provided name.
InFlightDiagnostic emitOpError(const Twine &message={})
Emit an error with the op name prefixed, like "'dim' op " which is convenient for verifiers.
A special type of RewriterBase that coordinates the application of a rewrite pattern on the current I...
RewritePatternSet & add(ConstructorArg &&arg, ConstructorArgs &&...args)
Add an instance of each of the pattern types 'Ts' to the pattern list with the given arguments.
This class coordinates the application of a rewrite on a set of IR, providing a way for clients to tr...
OpTy replaceOpWithNewOp(Operation *op, Args &&...args)
Replace the results of the given (original) op with a new op that is created without verification (re...
This class provides an abstraction over the various different ranges of value types.
Instances of the Type class are uniqued, have an immutable identifier and an optional mutable compone...
MLIRContext * getContext() const
Return the MLIRContext in which this type was uniqued.
bool isInteger() const
Return true if this is an integer type (with the specified width).
This class provides an abstraction over the different types of ranges over Values.
This class represents an instance of an SSA value in the MLIR system, representing a computable value...
Type getType() const
Return the type of this value.
The OpAsmOpInterface, see OpAsmInterface.td for more details.
SmallVector< int64_t, 4 > getCoordinates(ArrayRef< int64_t > basis, unsigned linearIndex)
@ Write
Write register with '=' modifier.
@ ReadWrite
ReadWrite register with '+' modifier.
@ Read
Read register with no modifier.
mlir::ParseResult parseCTAGroup(mlir::OpAsmParser &parser, mlir::NVVM::CTAGroupKindAttr &groupAttr)
void nvvmInferResultRanges(std::optional< mlir::LLVM::ConstantRangeAttr > range, mlir::Value result, mlir::ArrayRef< mlir::ConstantIntRanges > argRanges, mlir::SetIntRangeFn setResultRanges)
mlir::LogicalResult verifyConstantRangeAttr(mlir::Operation *op, std::optional< mlir::LLVM::ConstantRangeAttr > rangeAttr)
void printCTAGroup(mlir::OpAsmPrinter &printer, mlir::Operation *, mlir::NVVM::CTAGroupKindAttr groupAttr)
std::pair< mlir::Type, unsigned > inferMMAType(mlir::NVVM::MMATypes type, mlir::NVVM::MMAFrag frag, int nRow, int nCol, mlir::MLIRContext *context)
Return the element type and number of elements associated with a wmma matrix of given chracteristics.
std::pair< llvm::Intrinsic::ID, llvm::SmallVector< llvm::Value * > > IDArgPair
A pair type of LLVM's Intrinsic ID and args (which are llvm values).
uint64_t getN(LevelType lt)
uint64_t getM(LevelType lt)
Include the generated interface declarations.
llvm::function_ref< void(Value, const ConstantIntRanges &)> SetIntRangeFn
The type of the setResultRanges callback provided to ops implementing InferIntRangeInterface.
Type getType(OpFoldResult ofr)
Returns the int type of the integer in ofr.
InFlightDiagnostic emitError(Location loc)
Utility method to emit an error message using this location.
LogicalResult matchAndRewrite(SubFOp op, PatternRewriter &rewriter) const override
OpRewritePattern(MLIRContext *context, PatternBenefit benefit=1, ArrayRef< StringRef > generatedNames={})
This represents an operation in an abstracted form, suitable for use with the builder APIs.