34#include "llvm/ADT/STLExtras.h"
35#include "llvm/ADT/TypeSwitch.h"
36#include "llvm/IR/IRBuilder.h"
37#include "llvm/IR/NVVMIntrinsicUtils.h"
38#include "llvm/Support/Casting.h"
39#include "llvm/Support/FormatVariadic.h"
40#include "llvm/Support/NVPTXAddrSpace.h"
41#include "llvm/Support/raw_ostream.h"
52#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.cpp.inc"
53#include "mlir/Dialect/LLVMIR/NVVMOpsEnums.cpp.inc"
55static constexpr unsigned notIntrinsic = llvm::Intrinsic::not_intrinsic;
62 auto ptrTy = llvm::cast<LLVM::LLVMPointerType>(
ptr.getType());
63 return ptrTy.getAddressSpace() ==
static_cast<unsigned>(targetAS);
80 NVVMMemorySpace targetAS) {
81 unsigned AS =
static_cast<unsigned>(targetAS);
82 return builder.CreateAddrSpaceCast(
83 ptr, llvm::PointerType::get(builder.getContext(), AS));
87static llvm::nvvm::CTAGroupKind
90 case NVVM::CTAGroupKind::CTA_1:
91 return llvm::nvvm::CTAGroupKind::CG_1;
92 case NVVM::CTAGroupKind::CTA_2:
93 return llvm::nvvm::CTAGroupKind::CG_2;
95 llvm_unreachable(
"unsupported cta_group value");
99 NVVM::CTAGroupKindAttr &groupAttr) {
103 std::optional<NVVM::CTAGroupKind> group =
104 NVVM::symbolizeCTAGroupKind(keyword);
107 groupAttr = NVVM::CTAGroupKindAttr::get(parser.
getContext(), *group);
112 NVVM::CTAGroupKindAttr groupAttr) {
113 printer << NVVM::stringifyCTAGroupKind(groupAttr.getValue());
116template <
typename AttrTy>
123 using EnumTy =
decltype(attr.getValue());
124 std::optional<EnumTy> value = NVVM::symbolizeEnum<EnumTy>(keyword);
126 return parser.
emitError(loc) <<
"unknown enum value '" << keyword <<
"'";
128 attr = AttrTy::get(parser.
getContext(), *value);
141 size_t numIm2ColOffsets,
143 if (tensorDims < 1 || tensorDims > 5)
144 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
152 "to use im2col mode, the tensor has to be at least 3-dimensional");
154 if (numIm2ColOffsets && (tensorDims != (numIm2ColOffsets + 2)))
156 loc,
"im2col offsets must be 2 less than number of coordinates");
165 if (!tensorSize.empty() && coordinates.size() != tensorSize.size()) {
167 emitError(loc,
"Expected coordinates size to be equal to tensor size");
170 if (!lowerStride.empty() && tensorSize.empty()) {
173 "Expected tensor_size to be present when lower_stride is provided");
174 }
else if (!lowerStride.empty() &&
175 lowerStride.size() != tensorSize.size() - 1) {
178 "Expected lower_stride size to be equal to one less than tensor size");
181 if (!lowerStride.empty() !=
static_cast<bool>(upperStride)) {
183 "Expected lower_stride and upper_stride to be either both "
184 "present or both absent");
187 bool isDimStride = tensorSize.size() > 0;
188 if (!isTile && isDimStride) {
190 loc,
"Only tile mode supports override address with dim and stride");
196LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
197 TMAStoreMode mode = getMode();
201 if (getPredicate()) {
202 if (mode != TMAStoreMode::TILE)
203 return emitError(
"Inline-ptx lowering supported only for Tile mode.");
204 if (getL2CacheHint())
205 return emitError(
"Inline-ptx lowering unsupported with L2 cache-hint.");
210 case TMAStoreMode::TILE:
212 case TMAStoreMode::IM2COL:
213 case TMAStoreMode::IM2COL_W:
215 case TMAStoreMode::TILE_SCATTER4:
217 return emitError(
"Scatter4 mode expects 5 coordinates");
222LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::verify() {
223 TMAStoreMode mode = getMode();
225 mode == TMAStoreMode::IM2COL || mode == TMAStoreMode::IM2COL_W;
226 bool isTile = mode == TMAStoreMode::TILE;
232 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
235 if (mode == TMAStoreMode::TILE_SCATTER4 &&
getCoordinates().size() != 5)
236 overrideAddrRes =
emitError(
"Mode tile scatter4 expects 5 coordinates");
241LogicalResult CpAsyncOp::verify() {
242 if (getModifier() != LoadCacheModifierKind::CG &&
243 getModifier() != LoadCacheModifierKind::CA)
244 return emitError(
"Only CG and CA cache modifiers are supported.");
245 if (getSize() != 4 && getSize() != 8 && getSize() != 16)
246 return emitError(
"expected byte size to be either 4, 8 or 16.");
247 if (getModifier() == LoadCacheModifierKind::CG && getSize() != 16)
248 return emitError(
"CG cache modifier is only support for 16 bytes copy.");
255 if (tensorDims < 1 || tensorDims > 5)
256 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
258 auto checkTMALoadParams = [&](TMALoadMode mode,
bool isIm2col,
259 size_t expectedIm2colOff) -> LogicalResult {
260 if (isIm2col && (tensorDims < 3))
263 <<
" mode, the tensor has to be at least 3-dimensional";
265 if (numIm2colOff != expectedIm2colOff)
266 return emitError(loc) <<
" im2col offsets expected " << expectedIm2colOff
267 <<
" (provided " << numIm2colOff <<
")";
273 case TMALoadMode::TILE:
274 return checkTMALoadParams(mode,
false, 0);
275 case TMALoadMode::IM2COL:
276 return checkTMALoadParams(mode,
true, tensorDims - 2);
277 case TMALoadMode::IM2COL_W:
278 case TMALoadMode::IM2COL_W_128:
279 return checkTMALoadParams(mode,
true, 2);
280 case TMALoadMode::TILE_GATHER4:
281 return (tensorDims == 5)
282 ? checkTMALoadParams(mode,
false, 0)
283 :
emitError(loc,
"Gather4 mode expects 5 coordinates");
288LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
290 getMode(), getLoc());
293LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
294 TMALoadMode mode = getMode();
295 bool isCTAOnly = getIsCTAOnly();
296 if (getPredicate()) {
298 return emitError(
"Predicate is supported only for shared::cluster mode.");
299 if (mode != TMALoadMode::TILE && mode != TMALoadMode::IM2COL)
301 "Predicate is supported only for Tile and Im2col modes.");
303 NVVMMemorySpace expectedAS =
304 isCTAOnly ? NVVMMemorySpace::Shared : NVVMMemorySpace::SharedCluster;
305 unsigned AS = llvm::cast<LLVM::LLVMPointerType>(getDstMem().
getType())
307 if (AS != expectedAS)
310 ?
"Shared::cta destination requires address-space 3."
311 :
"Shared::cluster destination requires address-space 7.");
314 if (getMulticastMask())
315 return emitError(
"Multicast is not supported with shared::cta mode.");
317 return emitError(
"CTAGroup is not supported with shared::cta mode.");
322 getMode(), getLoc());
325LogicalResult CpAsyncBulkTensorReduceOp::verify() {
326 TMAStoreMode mode = getMode();
329 case TMAStoreMode::TILE:
331 case TMAStoreMode::IM2COL:
332 case TMAStoreMode::IM2COL_W:
334 case TMAStoreMode::TILE_SCATTER4:
335 return emitError(
"Scatter mode unsupported for CpAsyncBulkTensorReduceOp");
340LogicalResult CpAsyncBulkTensorReduceOverrideAddrOp::verify() {
342 getMode() == TMAStoreMode::IM2COL || getMode() == TMAStoreMode::IM2COL_W;
343 bool isTile = getMode() == TMAStoreMode::TILE;
349 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
352 if (getMode() == TMAStoreMode::TILE_SCATTER4)
354 "Scatter mode unsupported for CpAsyncBulkTensorReduceOverrideAddrOp");
359LogicalResult CpAsyncBulkGlobalToSharedClusterOp::verify() {
361 if (isSharedCTA && getMulticastMask())
362 return emitError(
"Multicast is not supported with shared::cta mode.");
368 NVVM::MemScopeKind scope,
369 Value retVal =
nullptr) {
370 if (scope != NVVM::MemScopeKind::CTA && scope != NVVM::MemScopeKind::CLUSTER)
371 return op->
emitError(
"mbarrier scope must be either CTA or Cluster");
374 bool hasRetValue =
static_cast<bool>(retVal);
375 if (isSharedCluster && hasRetValue)
377 "mbarrier in shared_cluster space cannot return any value");
382LogicalResult MBarrierArriveOp::verify() {
387LogicalResult MBarrierArriveDropOp::verify() {
392LogicalResult MBarrierArriveExpectTxOp::verify() {
396 if (getPredicate()) {
397 if (getScope() != NVVM::MemScopeKind::CTA)
398 return emitError(
"mbarrier scope must be CTA when using predicate");
401 return emitError(
"mbarrier in shared_cluster space is not supported when "
405 return emitError(
"return-value is not supported when using predicate");
407 if (getRelaxed() ==
true)
408 return emitError(
"mbarrier with relaxed semantics is not supported when "
415LogicalResult MBarrierArriveDropExpectTxOp::verify() {
430 inferredReturnTypes.push_back(IntegerType::get(context, 64));
435MBarrierArriveOp::inferReturnTypes(
MLIRContext *context,
436 std::optional<Location> location,
437 MBarrierArriveOp::Adaptor adaptor,
440 inferredReturnTypes);
443LogicalResult MBarrierArriveDropOp::inferReturnTypes(
444 MLIRContext *context, std::optional<Location> location,
445 MBarrierArriveDropOp::Adaptor adaptor,
448 inferredReturnTypes);
451LogicalResult MBarrierArriveExpectTxOp::inferReturnTypes(
452 MLIRContext *context, std::optional<Location> location,
453 MBarrierArriveExpectTxOp::Adaptor adaptor,
457 if (adaptor.getPredicate())
460 inferredReturnTypes);
463LogicalResult MBarrierArriveDropExpectTxOp::inferReturnTypes(
464 MLIRContext *context, std::optional<Location> location,
465 MBarrierArriveDropExpectTxOp::Adaptor adaptor,
468 inferredReturnTypes);
478 return inferred == actual;
487bool MBarrierArriveExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
491bool MBarrierArriveDropExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
496LogicalResult MBarrierExpectTxOp::verify() {
500LogicalResult MBarrierCompleteTxOp::verify() {
504LogicalResult MBarrierTestWaitOp::verify() {
508LogicalResult MBarrierTryWaitOp::verify() {
512LogicalResult ConvertFloatToTF32Op::verify() {
513 using RndMode = NVVM::FPRoundingMode;
517 return emitError(
"Relu not supported with rna rounding mode.");
524 "Only {rn,rz,rna} rounding modes supported for ConvertFloatToTF32Op.");
529LogicalResult ConvertF32x2ToF6x2Op::verify() {
532 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy())) {
533 return emitOpError(
"Only ")
534 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
535 << mlir::Float6E3M2FNType::get(ctx)
536 <<
" types are supported for conversions from f32x2 to f6x2.";
541LogicalResult ConvertF32x2ToF8x2Op::verify() {
542 using RndMode = NVVM::FPRoundingMode;
543 using SatMode = NVVM::SaturationMode;
545 bool isRoundingModeRN = getRnd() == RndMode::RN;
546 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
547 bool isRoundingModeRP = getRnd() == RndMode::RP;
548 bool isSatFinite = getSat() == SatMode::SATFINITE;
550 bool hasRelu = getRelu();
555 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
557 if (!isRoundingModeRN) {
558 return emitOpError(
"Only RN rounding mode is supported for "
559 "conversions from f32x2 to ")
560 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
561 << mlir::Float8E5M2Type::get(ctx) <<
" types";
564 return emitOpError(
"Only SATFINITE saturation mode is supported "
567 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
568 << mlir::Float8E5M2Type::get(ctx) <<
" types";
572 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
573 if (!(isRoundingModeRZ || isRoundingModeRP)) {
574 return emitOpError(
"Only RZ and RP rounding modes are supported for "
575 "conversions from f32x2 to ")
576 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
579 return emitOpError(
"relu not supported for conversions to ")
580 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
585 return emitOpError(
"Only ")
586 << mlir::Float8E4M3FNType::get(ctx) <<
", "
587 << mlir::Float8E5M2Type::get(ctx) <<
", and "
588 << mlir::Float8E8M0FNUType::get(ctx)
590 "supported for conversions from f32x2 to f8x2";
594LogicalResult ConvertF16x2ToF8x2Op::verify() {
597 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy())) {
598 return emitOpError(
"Only ")
599 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
600 << mlir::Float8E5M2Type::get(ctx)
601 <<
" types are supported for conversions from f16x2 to f8x2.";
606LogicalResult ConvertBF16x2ToF8x2Op::verify() {
607 using RndMode = NVVM::FPRoundingMode;
608 using SatMode = NVVM::SaturationMode;
610 bool isRoundingModeRN = getRnd() == RndMode::RN;
611 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
612 bool isRoundingModeRP = getRnd() == RndMode::RP;
613 bool isSatFinite = getSat() == SatMode::SATFINITE;
614 bool hasRelu = getRelu();
619 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
621 if (!isRoundingModeRN)
622 return emitOpError(
"Only RN rounding mode is supported for "
623 "conversions from bf16x2 to ")
624 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
625 << mlir::Float8E5M2Type::get(ctx) <<
" types";
627 return emitOpError(
"Only SATFINITE saturation mode is supported "
628 "for conversions from bf16x2 to ")
629 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
630 << mlir::Float8E5M2Type::get(ctx) <<
" types";
633 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
634 if (!(isRoundingModeRZ || isRoundingModeRP))
635 return emitOpError(
"Only RZ and RP rounding modes are supported for "
636 "conversions from bf16x2 to ")
637 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
639 return emitOpError(
"relu not supported for conversions to ")
640 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
644 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
649LogicalResult ConvertF32x2ToF4x2Op::verify() {
652 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
653 return emitOpError(
"Only ")
654 << mlir::Float4E2M1FNType::get(ctx)
655 <<
" type is supported for conversions from f32x2 to f4x2.";
660LogicalResult ConvertF8x2ToBF16x2Op::verify() {
662 if (llvm::isa<Float8E8M0FNUType>(getSrcType())) {
663 if (getSat() != SaturationMode::NONE)
665 "Only NONE saturation mode is supported for conversions from ")
666 << Float8E8M0FNUType::get(ctx) <<
" type";
667 if (getScaleFactor())
668 return emitOpError(
"scaleFactor not supported for conversions from ")
669 << Float8E8M0FNUType::get(ctx) <<
" type";
671 return emitOpError(
"relu not supported for conversions from ")
672 << Float8E8M0FNUType::get(ctx) <<
" type";
678LogicalResult PermuteOp::verify() {
679 using Mode = NVVM::PermuteMode;
680 bool hasHi =
static_cast<bool>(getHi());
687 return emitError(
"mode '") << getMode() <<
"' requires 'hi' operand.";
695 << getMode() <<
"' does not accept 'hi' operand.";
710 static constexpr FPRoundingMode validRndModes[] = {
711 FPRoundingMode::RN, FPRoundingMode::RZ, FPRoundingMode::RS};
713 if (!llvm::is_contained(validRndModes, rnd)) {
715 "Only RN, RZ, and RS rounding modes are supported for "
716 "conversions from f32x2 to ")
720 if (rnd == FPRoundingMode::RS) {
721 if (!hasRandomBits) {
722 return op->
emitOpError(
"random_bits is required for RS rounding mode.");
727 "random_bits not supported for RN and RZ rounding modes.");
734LogicalResult ConvertF32x2ToF16x2Op::verify() {
736 getRandomBits() ?
true :
false, *
this);
739LogicalResult ConvertF32x2ToBF16x2Op::verify() {
741 getRandomBits() ?
true :
false, *
this);
744LogicalResult ConvertF32x4ToF8x4Op::verify() {
747 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy()))
748 return emitOpError(
"Only ")
749 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
750 << mlir::Float8E5M2Type::get(ctx)
751 <<
" types are supported for conversions from f32x4 to f8x4.";
756LogicalResult ConvertF32x4ToF6x4Op::verify() {
759 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy()))
760 return emitOpError(
"Only ")
761 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
762 << mlir::Float6E3M2FNType::get(ctx)
763 <<
" types are supported for conversions from f32x4 to f6x4.";
768LogicalResult ConvertF32x4ToF4x4Op::verify() {
771 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
772 return emitOpError(
"Only ") << mlir::Float4E2M1FNType::get(ctx)
773 <<
" type is supported for conversions from "
779LogicalResult BulkStoreOp::verify() {
780 if (getInitVal() != 0)
781 return emitOpError(
"only 0 is supported for initVal, got ") << getInitVal();
785LogicalResult AsyncStoreGlobalOp::verify() {
786 NVVM::MemScopeKind scope = getScope();
787 bool isMmio = getMmio();
788 bool isMultimem = getMultimem();
790 if (scope != MemScopeKind::SYS && scope != MemScopeKind::GPU)
791 return emitOpError(
"scope must be either SYS or GPU");
793 if (isMmio && scope != MemScopeKind::SYS)
794 return emitOpError(
"mmio is only supported for SYS scope");
796 if (isMmio && isMultimem)
797 return emitOpError(
"multimem is not supported with mmio");
802LogicalResult PMEventOp::verify() {
803 auto eventId = getEventId();
804 auto maskedEventId = getMaskedEventId();
805 if (!maskedEventId && !eventId) {
806 return emitOpError() <<
"either `id` or `mask` must be set";
809 if (maskedEventId && eventId) {
810 return emitOpError() <<
"`id` and `mask` cannot be set at the same time";
814 if (eventId < 0 || eventId > 15) {
815 return emitOpError() <<
"`id` must be between 0 and 15";
819 return llvm::success();
825std::optional<mlir::NVVM::MMATypes>
826MmaOp::inferOperandMMAType(
Type operandElType,
bool isAccumulator) {
828 VectorType::get(2, Float16Type::get(operandElType.
getContext()));
829 if (operandElType.
isF64())
830 return NVVM::MMATypes::f64;
831 if (operandElType.
isF16() || operandElType == half2Type)
832 return NVVM::MMATypes::f16;
833 if (operandElType.
isF32() && isAccumulator)
834 return NVVM::MMATypes::f32;
835 if (operandElType.
isF32() && !isAccumulator)
836 return NVVM::MMATypes::tf32;
837 if (llvm::isa<IntegerType>(operandElType)) {
839 return NVVM::MMATypes::s32;
843 if (
auto structType = llvm::dyn_cast<LLVM::LLVMStructType>(operandElType)) {
844 if (structType.getBody().empty())
846 return inferOperandMMAType(structType.getBody()[0], isAccumulator);
853 return (type == MMATypes::u4 || type == MMATypes::s4);
857 return (type == MMATypes::u8 || type == MMATypes::s8);
862 type == MMATypes::s32;
865MMATypes MmaOp::accumPtxType() {
866 std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
867 getODSOperands(2).getTypes().front(),
true);
868 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
872MMATypes MmaOp::resultPtxType() {
873 std::optional<mlir::NVVM::MMATypes> val =
874 inferOperandMMAType(getResult().
getType(),
true);
875 assert(val.has_value() &&
"result PTX type should always be inferrable");
879template <
typename AttrTy>
881 StringRef keyword, AttrTy value) {
882 printer << (isFirst ?
" " :
", ") << keyword <<
" = ";
887template <
typename AttrTy>
889 StringRef keyword, AttrTy value) {
890 printer << (isFirst ?
" " :
", ") << keyword <<
" = "
891 << NVVM::stringifyEnum(value.getValue());
897 printer << (isFirst ?
" " :
", ") << keyword;
901template <
typename AttrTy>
905 if (attributes.
get(name))
907 "duplicate property '" + name +
"'");
911 attributes.
append(name, value);
915template <
typename AttrTy>
919 if (attributes.
get(name))
921 "duplicate property '" + name +
"'");
925 attributes.
append(name, value);
930 return llvm::is_contained(
932 "shape",
"b1Op",
"intOverflowBehavior",
"layoutA",
"layoutB",
933 "multiplicandAPtxType",
"multiplicandBPtxType",
"orderedMetadata",
934 "kind",
"scaleVecSize",
"blockScaleFormat",
"operandSegmentSizes"},
946 if (!llvm::is_contained(allowedKeywords, keyword))
948 "unknown MMA property '" + keyword +
"'");
950 ParseResult parseResult =
success();
951 if (keyword ==
"shape")
954 else if (keyword ==
"b1_op")
957 else if (keyword ==
"int_overflow")
959 parser, attributes,
"intOverflowBehavior");
960 else if (keyword ==
"layout_a")
963 else if (keyword ==
"layout_b")
966 else if (keyword ==
"multiplicand_a_ptx_type")
968 parser, attributes,
"multiplicandAPtxType");
969 else if (keyword ==
"multiplicand_b_ptx_type")
971 parser, attributes,
"multiplicandBPtxType");
972 else if (keyword ==
"kind") {
973 if (llvm::is_contained(allowedKeywords,
"block_scale_format"))
975 parser, attributes,
"kind");
979 }
else if (keyword ==
"scale_vec_size")
981 parser, attributes,
"scaleVecSize");
982 else if (keyword ==
"block_scale_format")
984 parser, attributes,
"blockScaleFormat");
985 else if (keyword ==
"ordered_metadata") {
986 if (attributes.
get(
"orderedMetadata"))
988 "duplicate property 'orderedMetadata'");
992 "unknown MMA property '" + keyword +
"'");
994 if (failed(parseResult))
1000 for (StringRef property : requiredProperties) {
1001 if (!attributes.
get(property))
1003 "missing required property '" + property +
"'");
1013 "inherent property '" + attribute.getName().getValue() +
1014 "' must be spelled directly in the operation syntax");
1015 attributes.
append(attribute);
1022 struct MMAOperandFragment {
1023 StringRef operandName;
1024 StringRef ptxTypeAttr;
1025 SmallVector<Value, 4> regs;
1026 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1027 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1030 std::array<MMAOperandFragment, 3> frags{
1031 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1032 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1033 MMAOperandFragment(
"C",
"")};
1035 mlir::NVVM::MmaOp::getOperandSegmentSizeAttr()};
1037 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
1038 auto &frag = frags[fragIdx];
1039 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1040 for (
auto operandIdx = varOperandSpec.first;
1041 operandIdx < varOperandSpec.first + varOperandSpec.second;
1043 frag.regs.push_back(this->getOperand(operandIdx));
1044 if (operandIdx == 0) {
1045 regTypes.push_back(this->getOperand(operandIdx).
getType());
1048 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1049 regTypes.back(), fragIdx >= 2);
1051 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1054 auto printMmaOperand = [&](
const MMAOperandFragment &frag) ->
void {
1055 p <<
" " << frag.operandName;
1061 for (
const auto &frag : frags) {
1062 printMmaOperand(frag);
1065 bool isFirstProperty =
true;
1069 if (getIntOverflowBehaviorAttr())
1071 getIntOverflowBehaviorAttr());
1074 if (getMultiplicandAPtxTypeAttr() &&
1075 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1077 getMultiplicandAPtxTypeAttr());
1078 if (getMultiplicandBPtxTypeAttr() &&
1079 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1081 getMultiplicandBPtxTypeAttr());
1082 llvm::append_range(ignoreAttrNames,
1084 getIntOverflowBehaviorAttrName(),
1085 getLayoutAAttrName(),
1086 getLayoutBAttrName(),
1087 getMultiplicandAPtxTypeAttrName(),
1088 getMultiplicandBPtxTypeAttrName()});
1095 frags[1].regs[0].getType(),
1096 frags[2].regs[0].getType()},
1105 std::optional<MMAIntOverflow> intOverflow,
1106 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
1107 std::optional<std::array<MMALayout, 2>> multiplicandLayouts) {
1109 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1114 result.addOperands(operandA);
1115 result.addOperands(operandB);
1116 result.addOperands(operandC);
1118 if (multiplicandPtxTypes) {
1119 result.addAttribute(
"multiplicandAPtxType",
1120 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1121 result.addAttribute(
"multiplicandBPtxType",
1122 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1124 if (
auto res = inferOperandMMAType(operandA[0].
getType(),
false))
1125 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1126 if (
auto res = inferOperandMMAType(operandB[0].
getType(),
false))
1127 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1130 if (multiplicandLayouts) {
1131 result.addAttribute(
"layoutA",
1132 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[0]));
1133 result.addAttribute(
"layoutB",
1134 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[1]));
1136 result.addAttribute(
"layoutA", MMALayoutAttr::get(ctx, MMALayout::row));
1137 result.addAttribute(
"layoutB", MMALayoutAttr::get(ctx, MMALayout::col));
1140 if (intOverflow.has_value())
1141 result.addAttribute(
"intOverflowBehavior",
1142 MMAIntOverflowAttr::get(ctx, *intOverflow));
1143 if (b1Op.has_value())
1144 result.addAttribute(
"b1Op", MMAB1OpAttr::get(ctx, *b1Op));
1146 result.addTypes(resultType);
1148 MmaOp::getOperandSegmentSizeAttr(),
1150 static_cast<int32_t>(operandB.size()),
1151 static_cast<int32_t>(operandC.size())}));
1160 struct MMAOperandFragment {
1161 std::optional<MMATypes> elemtype;
1162 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1163 SmallVector<Type> regTypes;
1167 std::array<MMAOperandFragment, 4> frags;
1173 MMAOperandFragment &frag) -> LogicalResult {
1192 {
"shape",
"b1_op",
"int_overflow",
"layout_a",
1193 "layout_b",
"multiplicand_a_ptx_type",
1194 "multiplicand_b_ptx_type"},
1195 {
"shape",
"layoutA",
"layoutB"}))
1207 if (operandTypes.size() != 3)
1210 "expected one type for each operand segment but got " +
1211 Twine(operandTypes.size()) +
" types");
1212 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1213 auto &frag = frags[iter.index()];
1214 frag.regTypes.resize(frag.regs.size(), iter.value());
1218 frag.elemtype = inferOperandMMAType(frag.regTypes[0],
1225 frags[3].elemtype = inferOperandMMAType(resultType,
true);
1227 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1228 "multiplicandBPtxType"};
1229 for (
unsigned idx = 0; idx < names.size(); idx++) {
1230 const auto &frag = frags[idx];
1231 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1232 if (!frag.elemtype.has_value() && !attr.has_value()) {
1235 "attribute " + names[idx] +
1236 " is not provided explicitly and cannot be inferred");
1238 if (!attr.has_value())
1240 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1243 result.addTypes(resultType);
1244 if (!namedAttributes.
empty())
1245 result.addAttributes(namedAttributes);
1246 result.addAttribute(MmaOp::getOperandSegmentSizeAttr(),
1248 static_cast<int32_t>(frags[0].regs.size()),
1249 static_cast<int32_t>(frags[1].regs.size()),
1250 static_cast<int32_t>(frags[2].regs.size()),
1255LogicalResult MmaOp::verify() {
1257 auto f16Ty = Float16Type::get(context);
1258 auto i32Ty = IntegerType::get(context, 32);
1259 auto f16x2Ty = VectorType::get(2, f16Ty);
1260 auto f32Ty = Float32Type::get(context);
1261 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1262 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1264 auto s32x4StructTy =
1265 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1266 auto f32x8StructTy =
1268 auto f16x2x2StructTy =
1269 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1270 auto f32x4StructTy =
1271 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1272 auto s32x2StructTy =
1273 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1275 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1276 getShapeAttr().getK()};
1282 AllowedShapes allowedShapes;
1283 AllowedTypes expectedA;
1284 AllowedTypes expectedB;
1285 AllowedTypes expectedC;
1290 if (mmaShape[0] == 16) {
1292 Type multiplicandFragType;
1293 switch (*getMultiplicandAPtxType()) {
1294 case MMATypes::tf32:
1296 multiplicandFragType = i32Ty;
1297 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1298 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1300 case MMATypes::bf16:
1302 multiplicandFragType = i32Ty;
1303 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1304 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1308 multiplicandFragType = f16x2Ty;
1309 expectedResult.push_back(f16x2x2StructTy);
1310 expectedResult.push_back(f32x4StructTy);
1312 case MMATypes::e4m3:
1313 case MMATypes::e5m2:
1317 multiplicandFragType = i32Ty;
1318 expectedResult.push_back(f16x2x2StructTy);
1319 expectedResult.push_back(f32x4StructTy);
1333 return emitError(
"invalid shape or multiplicand type: ")
1334 << getMultiplicandAPtxType().value();
1338 expectedResult.push_back(s32x4StructTy);
1339 expectedC.emplace_back(4, i32Ty);
1340 multiplicandFragType = i32Ty;
1342 expectedC.emplace_back(2, f16x2Ty);
1343 expectedC.emplace_back(4, f32Ty);
1346 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor);
1347 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1348 expectedA.emplace_back(unitA, multiplicandFragType);
1349 expectedB.emplace_back(unitB, multiplicandFragType);
1350 allowedShapes.push_back({16, 8, kFactor});
1351 allowedShapes.push_back({16, 8, kFactor * 2});
1353 if (resultPtxType() != accumPtxType())
1354 return emitOpError(
"ctype does not match dtype");
1358 if (mmaShape[0] == 8) {
1359 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1360 expectedA.emplace_back(2, f16x2Ty);
1361 expectedB.emplace_back(2, f16x2Ty);
1362 expectedResult.push_back(f16x2x4StructTy);
1363 expectedResult.push_back(f32x8StructTy);
1364 expectedC.emplace_back(4, f16x2Ty);
1365 expectedC.emplace_back(8, f32Ty);
1366 allowedShapes.push_back({8, 8, 4});
1368 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1369 Type f64Ty = Float64Type::get(context);
1370 expectedA.emplace_back(1, f64Ty);
1371 expectedB.emplace_back(1, f64Ty);
1372 expectedC.emplace_back(2, f64Ty);
1373 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1375 allowedShapes.push_back({8, 8, 4});
1378 expectedA.push_back({i32Ty});
1379 expectedB.push_back({i32Ty});
1380 expectedC.push_back({i32Ty, i32Ty});
1381 expectedResult.push_back(s32x2StructTy);
1383 allowedShapes.push_back({8, 8, 32});
1385 allowedShapes.push_back({8, 8, 16});
1386 if (getMultiplicandAPtxType().value() == MMATypes::b1)
1387 allowedShapes.push_back({8, 8, 128});
1391 std::string errorMessage;
1392 llvm::raw_string_ostream errorStream(errorMessage);
1395 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1396 !llvm::is_contained(allowedShapes, mmaShape)) {
1397 errorStream <<
"unimplemented variant for MMA shape <";
1398 llvm::interleaveComma(mmaShape, errorStream);
1400 return emitOpError(errorMessage);
1404 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1405 for (
const auto &iter : llvm::enumerate(
1406 std::array<AllowedTypes, 3>{std::move(expectedA),
1407 std::move(expectedB),
1408 std::move(expectedC)})) {
1409 auto spec = this->getODSOperandIndexAndLength(iter.index());
1411 operand_type_begin() + spec.first +
1413 bool match = llvm::is_contained(iter.value(), operandTySeg);
1416 errorStream <<
"Could not match types for the "
1417 << operandNames[iter.index()]
1418 <<
" operands; expected one of ";
1419 for (
const auto &x : iter.value()) {
1420 errorStream << x.size() <<
"x" << x[0] <<
" ";
1422 errorStream <<
"but got ";
1423 llvm::interleaveComma(operandTySeg, errorStream);
1424 return emitOpError(errorMessage);
1429 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1430 return expectedResultType == getResult().getType();
1433 <<
"Could not match allowed types for the result; expected one of ";
1434 llvm::interleaveComma(expectedResult, errorStream);
1435 errorStream <<
" but got " << getResult().getType();
1436 return emitOpError(errorMessage);
1440 if (getMultiplicandAPtxType() == MMATypes::b1 && !getB1Op()) {
1441 return emitOpError(
"op requires " + getB1OpAttrName().strref() +
1449 if (!getIntOverflowBehavior())
1450 return emitOpError(
"op requires " +
1451 getIntOverflowBehaviorAttrName().strref() +
1459 (mmaShape[0] == 8 && mmaShape[1] == 8 && mmaShape[2] == 4 &&
1460 getMultiplicandAPtxType() == MMATypes::f16);
1462 if (!isM8N8K4_F16) {
1464 if (getLayoutA() != MMALayout::row || getLayoutB() != MMALayout::col) {
1465 return emitOpError(
"requires layoutA = #nvvm.mma_layout<row> and "
1466 "layoutB = #nvvm.mma_layout<col> for shape <")
1467 << mmaShape[0] <<
", " << mmaShape[1] <<
", " << mmaShape[2]
1468 <<
"> with element types " << *getMultiplicandAPtxType() <<
" and "
1469 << *getMultiplicandBPtxType()
1470 <<
". Only m8n8k4 with f16 supports other layouts.";
1477MMATypes MmaSpOp::accumPtxType() {
1478 std::optional<mlir::NVVM::MMATypes> val = MmaOp::inferOperandMMAType(
1479 getODSOperands(2).getTypes().front(),
true);
1480 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
1484MMATypes MmaSpOp::resultPtxType() {
1485 std::optional<mlir::NVVM::MMATypes> val =
1486 MmaOp::inferOperandMMAType(getResult().
getType(),
true);
1487 assert(val.has_value() &&
"result PTX type should always be inferrable");
1493 llvm::IRBuilderBase &builder) {
1494 auto thisOp = cast<NVVM::MmaSpOp>(op);
1502 auto intId = MmaSpOp::getIntrinsicID(
1503 thisOp.getShape().getM(), thisOp.getShape().getN(),
1504 thisOp.getShape().getK(), thisOp.getIntOverflowBehavior(),
1505 thisOp.getOrderedMetadata(), thisOp.getKind(),
1506 *thisOp.getMultiplicandAPtxType(), *thisOp.getMultiplicandBPtxType(),
1507 thisOp.accumPtxType(), thisOp.resultPtxType());
1509 return {intId, args};
1514 struct MMAOperandFragment {
1515 StringRef operandName;
1516 StringRef ptxTypeAttr;
1517 SmallVector<Value, 4> regs;
1518 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1519 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1522 std::array<MMAOperandFragment, 5> frags{
1523 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1524 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1525 MMAOperandFragment(
"C",
""), MMAOperandFragment(
"sparseMetadata",
""),
1526 MMAOperandFragment(
"selector",
"")};
1528 mlir::NVVM::MmaSpOp::getOperandSegmentSizeAttr()};
1531 for (
unsigned fragIdx = 0; fragIdx < 3; fragIdx++) {
1532 auto &frag = frags[fragIdx];
1533 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1534 for (
auto operandIdx = varOperandSpec.first;
1535 operandIdx < varOperandSpec.first + varOperandSpec.second;
1537 frag.regs.push_back(this->getOperand(operandIdx));
1538 if (operandIdx == varOperandSpec.first) {
1539 regTypes.push_back(this->getOperand(operandIdx).
getType());
1542 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1543 regTypes.back(), fragIdx >= 2);
1545 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1549 frags[3].regs.push_back(getSparseMetadata());
1550 frags[4].regs.push_back(getSparsitySelector());
1552 auto printMmaSpOperand = [&](
const MMAOperandFragment &frag) ->
void {
1553 p <<
" " << frag.operandName;
1559 for (
const auto &frag : frags)
1560 printMmaSpOperand(frag);
1562 bool isFirstProperty =
true;
1564 if (getIntOverflowBehaviorAttr())
1566 getIntOverflowBehaviorAttr());
1567 if (getMultiplicandAPtxTypeAttr() &&
1568 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1570 getMultiplicandAPtxTypeAttr());
1571 if (getMultiplicandBPtxTypeAttr() &&
1572 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1574 getMultiplicandBPtxTypeAttr());
1575 if (getOrderedMetadata())
1582 getMultiplicandAPtxTypeAttrName(),
1583 getMultiplicandBPtxTypeAttrName(),
1584 getOrderedMetadataAttrName(), getKindAttrName()});
1589 for (
int i = 0; i < 3; ++i) {
1594 p <<
") -> " << getResult().getType();
1601 std::optional<MMAIntOverflow> intOverflow,
1602 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
1604 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1609 result.addOperands(operandA);
1610 result.addOperands(operandB);
1611 result.addOperands(operandC);
1612 result.addOperands(sparseMetadata);
1613 result.addOperands(sparsitySelector);
1615 if (multiplicandPtxTypes) {
1616 result.addAttribute(
"multiplicandAPtxType",
1617 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1618 result.addAttribute(
"multiplicandBPtxType",
1619 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1621 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
1622 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1623 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
1624 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1627 if (intOverflow.has_value())
1628 result.addAttribute(
"intOverflowBehavior",
1629 MMAIntOverflowAttr::get(ctx, *intOverflow));
1631 result.addTypes(resultType);
1633 MmaSpOp::getOperandSegmentSizeAttr(),
1635 static_cast<int32_t>(operandB.size()),
1636 static_cast<int32_t>(operandC.size()), 1,
1641 struct MMAOperandFragment {
1642 std::optional<MMATypes> elemtype;
1643 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1644 SmallVector<Type> regTypes;
1648 std::array<MMAOperandFragment, 6> frags;
1653 auto parseMmaSpOperand = [&](StringRef operandName,
1654 MMAOperandFragment &frag) -> LogicalResult {
1665 if (parseMmaSpOperand(
"A", frags[0]).
failed())
1667 if (parseMmaSpOperand(
"B", frags[1]).
failed())
1669 if (parseMmaSpOperand(
"C", frags[2]).
failed())
1671 if (parseMmaSpOperand(
"sparseMetadata", frags[3]).
failed())
1673 if (parseMmaSpOperand(
"selector", frags[4]).
failed())
1677 {
"shape",
"int_overflow",
"multiplicand_a_ptx_type",
1678 "multiplicand_b_ptx_type",
"ordered_metadata",
1693 if (operandTypes.size() != 3)
1696 "expected one type for each operand segment but got " +
1697 Twine(operandTypes.size()) +
" types");
1698 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1699 auto &frag = frags[iter.index()];
1700 frag.regTypes.resize(frag.regs.size(), iter.value());
1705 MmaOp::inferOperandMMAType(frag.regTypes[0],
1713 MmaOp::inferOperandMMAType(resultType,
true);
1728 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1729 "multiplicandBPtxType"};
1730 for (
unsigned idx = 0; idx < names.size(); idx++) {
1731 const auto &frag = frags[idx];
1732 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1733 if (!frag.elemtype.has_value() && !attr.has_value()) {
1736 "attribute " + names[idx] +
1737 " is not provided explicitly and cannot be inferred");
1739 if (!attr.has_value())
1741 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1744 result.addTypes(resultType);
1745 if (!namedAttributes.
empty())
1746 result.addAttributes(namedAttributes);
1747 result.addAttribute(MmaSpOp::getOperandSegmentSizeAttr(),
1749 static_cast<int32_t>(frags[0].regs.size()),
1750 static_cast<int32_t>(frags[1].regs.size()),
1751 static_cast<int32_t>(frags[2].regs.size()),
1758LogicalResult MmaSpOp::verify() {
1760 auto f16Ty = Float16Type::get(context);
1761 auto i32Ty = IntegerType::get(context, 32);
1762 auto f16x2Ty = VectorType::get(2, f16Ty);
1763 auto f32Ty = Float32Type::get(context);
1764 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1765 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1767 auto s32x4StructTy =
1768 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1769 auto f32x8StructTy =
1771 auto f16x2x2StructTy =
1772 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1773 auto f32x4StructTy =
1774 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1775 auto s32x2StructTy =
1776 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1778 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1779 getShapeAttr().getK()};
1785 AllowedShapes allowedShapes;
1786 AllowedTypes expectedA;
1787 AllowedTypes expectedB;
1788 AllowedTypes expectedC;
1793 if (mmaShape[0] == 16) {
1795 Type multiplicandFragType;
1796 switch (*getMultiplicandAPtxType()) {
1797 case MMATypes::tf32:
1799 multiplicandFragType = i32Ty;
1800 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1801 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1803 allowedShapes.push_back({16, 8, 8});
1804 allowedShapes.push_back({16, 8, 16});
1806 case MMATypes::bf16:
1808 multiplicandFragType = i32Ty;
1809 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1810 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1812 allowedShapes.push_back({16, 8, 16});
1813 allowedShapes.push_back({16, 8, 32});
1817 multiplicandFragType = f16x2Ty;
1818 expectedResult.push_back(f16x2x2StructTy);
1819 expectedResult.push_back(f32x4StructTy);
1821 allowedShapes.push_back({16, 8, 16});
1822 allowedShapes.push_back({16, 8, 32});
1828 allowedShapes.push_back({16, 8, 64});
1829 allowedShapes.push_back({16, 8, 128});
1835 allowedShapes.push_back({16, 8, 32});
1836 allowedShapes.push_back({16, 8, 64});
1838 case MMATypes::e4m3:
1839 case MMATypes::e5m2:
1840 case MMATypes::e3m2:
1841 case MMATypes::e2m3:
1842 case MMATypes::e2m1:
1844 multiplicandFragType = i32Ty;
1845 expectedResult.push_back(f16x2x2StructTy);
1846 expectedResult.push_back(f32x4StructTy);
1848 allowedShapes.push_back({16, 8, 64});
1851 return emitError(
"invalid shape or multiplicand type: ")
1852 << getMultiplicandAPtxType().value();
1856 expectedResult.push_back(s32x4StructTy);
1857 expectedC.emplace_back(4, i32Ty);
1858 multiplicandFragType = i32Ty;
1859 }
else if (*getMultiplicandAPtxType() >= MMATypes::e4m3 &&
1860 *getMultiplicandAPtxType() <= MMATypes::e2m1) {
1862 expectedC.emplace_back(2, f16x2Ty);
1863 expectedC.emplace_back(4, f32Ty);
1865 expectedC.emplace_back(2, f16x2Ty);
1866 expectedC.emplace_back(4, f32Ty);
1871 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor) / 2;
1872 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1873 expectedA.emplace_back(unitA, multiplicandFragType);
1874 expectedB.emplace_back(unitB, multiplicandFragType);
1876 if (resultPtxType() != accumPtxType())
1877 return emitOpError(
"ctype does not match dtype");
1881 if (mmaShape[0] == 8) {
1882 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1883 expectedA.emplace_back(2, f16x2Ty);
1884 expectedB.emplace_back(2, f16x2Ty);
1885 expectedResult.push_back(f16x2x4StructTy);
1886 expectedResult.push_back(f32x8StructTy);
1887 expectedC.emplace_back(4, f16x2Ty);
1888 expectedC.emplace_back(8, f32Ty);
1889 allowedShapes.push_back({8, 8, 4});
1891 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1892 Type f64Ty = Float64Type::get(context);
1893 expectedA.emplace_back(1, f64Ty);
1894 expectedB.emplace_back(1, f64Ty);
1895 expectedC.emplace_back(2, f64Ty);
1896 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1898 allowedShapes.push_back({8, 8, 4});
1901 expectedA.push_back({i32Ty});
1902 expectedB.push_back({i32Ty});
1903 expectedC.push_back({i32Ty, i32Ty});
1904 expectedResult.push_back(s32x2StructTy);
1906 allowedShapes.push_back({8, 8, 32});
1908 allowedShapes.push_back({8, 8, 16});
1912 std::string errorMessage;
1913 llvm::raw_string_ostream errorStream(errorMessage);
1916 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1917 !llvm::is_contained(allowedShapes, mmaShape)) {
1918 errorStream <<
"unimplemented variant for MMA shape <";
1919 llvm::interleaveComma(mmaShape, errorStream);
1921 return emitOpError(errorMessage);
1925 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1926 for (
const auto &iter : llvm::enumerate(
1927 std::array<AllowedTypes, 3>{std::move(expectedA),
1928 std::move(expectedB),
1929 std::move(expectedC)})) {
1930 auto spec = this->getODSOperandIndexAndLength(iter.index());
1932 operand_type_begin() + spec.first +
1934 bool match = llvm::is_contained(iter.value(), operandTySeg);
1937 errorStream <<
"Could not match types for the "
1938 << operandNames[iter.index()]
1939 <<
" operands; expected one of ";
1940 for (
const auto &x : iter.value()) {
1941 errorStream << x.size() <<
"x" << x[0] <<
" ";
1943 errorStream <<
"but got ";
1944 llvm::interleaveComma(operandTySeg, errorStream);
1945 return emitOpError(errorMessage);
1950 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1951 return expectedResultType == getResult().getType();
1954 <<
"Could not match allowed types for the result; expected one of ";
1955 llvm::interleaveComma(expectedResult, errorStream);
1956 errorStream <<
" but got " << getResult().getType();
1957 return emitOpError(errorMessage);
1964 if (!getIntOverflowBehavior())
1965 return emitOpError(
"op requires " +
1966 getIntOverflowBehaviorAttrName().strref() +
1971 if (!getSparseMetadata().
getType().isInteger(32)) {
1972 return emitOpError() <<
"sparse metadata must be i32 type";
1976 if (!getSparsitySelector().
getType().isInteger(32)) {
1977 return emitOpError() <<
"sparsity selector must be i32 type";
1989struct MMAOperandFragment {
1990 StringRef operandName;
1991 StringRef ptxTypeAttr;
1992 SmallVector<Value, 4> regs;
1993 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1994 : operandName(name), ptxTypeAttr(ptxTypeName) {}
2001 p <<
" " << name <<
"[";
2020template <
typename Op>
2025 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
2026 auto &frag = frags[fragIdx];
2027 auto varOperandSpec = op.getODSOperandIndexAndLength(fragIdx);
2028 for (
auto operandIdx = varOperandSpec.first;
2029 operandIdx < varOperandSpec.first + varOperandSpec.second;
2031 frag.regs.push_back(op.getOperand(operandIdx));
2032 if (fragIdx == 0 && operandIdx == varOperandSpec.first) {
2033 regTypes.push_back(op.getOperand(operandIdx).getType());
2037 regTypes.push_back(frag.regs[0].getType());
2039 std::optional<MMATypes> inferredType =
2040 MmaOp::inferOperandMMAType(regTypes.back(),
2043 ignoreAttrNames.push_back(frag.ptxTypeAttr);
2054 auto typeParser = [&]() {
2058 operandTypes.push_back(ty);
2064 if (operandTypes.size() != 3)
2066 "expected exactly 3 types");
2075 if (!attrs.
get(
"multiplicandAPtxType")) {
2076 if (
auto inferredType =
2077 MmaOp::inferOperandMMAType(operandTypes[0],
false)) {
2078 attrs.
set(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *inferredType));
2081 if (!attrs.
get(
"multiplicandBPtxType")) {
2082 if (
auto inferredType =
2083 MmaOp::inferOperandMMAType(operandTypes[1],
false)) {
2084 attrs.
set(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *inferredType));
2090template <
typename OpType>
2093 ScaleVecSize scaleVecSize,
2094 BlockScaleFormat blockScaleFormat,
2095 MMABlockScaleKind kind) {
2097 auto &properties =
result.getOrAddProperties<
typename OpType::Properties>();
2098 properties.setShape(
2100 properties.setScaleVecSize(ScaleVecSizeAttr::get(ctx, scaleVecSize));
2101 properties.setBlockScaleFormat(
2102 BlockScaleFormatAttr::get(ctx, blockScaleFormat));
2103 properties.setKind(MMABlockScaleKindAttr::get(ctx, kind));
2110 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
2111 if (multiplicandPtxTypes) {
2112 result.addAttribute(
"multiplicandAPtxType",
2113 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
2114 result.addAttribute(
"multiplicandBPtxType",
2115 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
2117 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
2118 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
2119 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
2120 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
2125template <
typename OpTy>
2127 return *MmaOp::inferOperandMMAType(
2128 cast<LLVM::LLVMStructType>(op.getRes().getType()).getBody()[0],
2138 std::array<MMAOperandFragment, 3> frags{
2139 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2140 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2141 MMAOperandFragment(
"C",
"")};
2143 mlir::NVVM::MmaBlockScaleOp::getOperandSegmentSizeAttr()};
2148 for (
const auto &frag : frags)
2153 {getScaleAData(), getByteIdA(), getThreadIdA()});
2155 {getScaleBData(), getByteIdB(), getThreadIdB()});
2157 bool isFirstProperty =
true;
2159 if (getMultiplicandAPtxTypeAttr() &&
2160 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2162 getMultiplicandAPtxTypeAttr());
2163 if (getMultiplicandBPtxTypeAttr() &&
2164 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2166 getMultiplicandBPtxTypeAttr());
2168 getScaleVecSizeAttr());
2170 getBlockScaleFormatAttr());
2175 getMultiplicandBPtxTypeAttrName(),
2176 getScaleVecSizeAttrName(),
2177 getBlockScaleFormatAttrName(), getKindAttrName()});
2184 frags[1].regs[0].getType(),
2185 frags[2].regs[0].getType()},
2191ParseResult MmaBlockScaleOp::parse(
OpAsmParser &parser,
2193 struct LocalOperandFragment {
2194 std::optional<MMATypes> elemtype;
2195 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2199 std::array<LocalOperandFragment, 3> frags;
2215 {
"shape",
"multiplicand_a_ptx_type",
2216 "multiplicand_b_ptx_type",
"scale_vec_size",
2217 "block_scale_format",
"kind"},
2218 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2232 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2233 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2236 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2246 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2256 result.addAttributes(namedAttributes);
2260 result.addTypes(resultTypes);
2261 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2263 static_cast<int32_t>(frags[0].regs.size()),
2264 static_cast<int32_t>(frags[1].regs.size()),
2265 static_cast<int32_t>(frags[2].regs.size()),
2276void MmaBlockScaleOp::build(
2281 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2282 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2283 MMABlockScaleKind kind) {
2284 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2287 blockScaleFormat, kind);
2289 result.addOperands(operandA);
2290 result.addOperands(operandB);
2291 result.addOperands(operandC);
2293 {scaleAData, byteIdA, threadIdA, scaleBData, byteIdB, threadIdB});
2296 multiplicandPtxTypes);
2298 result.addTypes(resultType);
2299 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2301 static_cast<int32_t>(operandA.size()),
2302 static_cast<int32_t>(operandB.size()),
2303 static_cast<int32_t>(operandC.size()),
2315 auto curOp = cast<NVVM::MmaBlockScaleOp>(op);
2319 for (
Value operand : curOp.getOperandA())
2321 for (
Value operand : curOp.getOperandB())
2323 for (
Value operand : curOp.getOperandC())
2327 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2328 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2329 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2330 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2331 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2332 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2334 unsigned intId = MmaBlockScaleOp::getIntrinsicID(
2335 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2336 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2338 curOp.getBlockScaleFormat(), curOp.getKind());
2340 return {intId, args};
2343LogicalResult MmaBlockScaleOp::verify() {
2349 if (m == 16 && n == 8 && k == 64) {
2350 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2351 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2353 "unsupported MMATypes attribute for mma.m16n8k64.(mxf4nvf4|mxf4)");
2354 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2355 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2357 "unsupported ScaleVecSize attribute for mma.m16n8k64.mxf4");
2358 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2360 "unsupported BlockScaleFormat attribute for mma.m16n8k64.mxf4");
2361 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2362 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2363 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2364 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2365 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2366 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2367 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2368 "attributes for mma.m16n8k64.mxf4nvf4");
2370 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k64");
2372 }
else if (m == 16 && n == 8 && k == 32) {
2373 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2374 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2375 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2377 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2378 "attributes for mma.m16n8k32");
2380 result = emitOpError(
"unsupported Geom for mma with block scaling");
2391 std::array<MMAOperandFragment, 3> frags{
2392 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2393 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2394 MMAOperandFragment(
"C",
"")};
2396 mlir::NVVM::MmaSpBlockScaleOp::getOperandSegmentSizeAttr()};
2401 for (
const auto &frag : frags)
2410 {getScaleAData(), getByteIdA(), getThreadIdA()});
2412 {getScaleBData(), getByteIdB(), getThreadIdB()});
2414 bool isFirstProperty =
true;
2416 if (getMultiplicandAPtxTypeAttr() &&
2417 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2419 getMultiplicandAPtxTypeAttr());
2420 if (getMultiplicandBPtxTypeAttr() &&
2421 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2423 getMultiplicandBPtxTypeAttr());
2426 getScaleVecSizeAttr());
2428 getBlockScaleFormatAttr());
2433 getMultiplicandBPtxTypeAttrName(),
2434 getOrderedMetadataAttrName(),
2435 getScaleVecSizeAttrName(),
2436 getBlockScaleFormatAttrName(), getKindAttrName()});
2443 frags[1].regs[0].getType(),
2444 frags[2].regs[0].getType()},
2450ParseResult MmaSpBlockScaleOp::parse(
OpAsmParser &parser,
2452 struct LocalOperandFragment {
2453 std::optional<MMATypes> elemtype;
2454 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2458 std::array<LocalOperandFragment, 3> frags;
2481 {
"shape",
"multiplicand_a_ptx_type",
2482 "multiplicand_b_ptx_type",
"ordered_metadata",
2483 "scale_vec_size",
"block_scale_format",
"kind"},
2484 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2498 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2499 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2502 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2511 .resolveOperands(metadataOperands, i32Type, parser.
getNameLoc(),
2524 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2534 result.addAttributes(namedAttributes);
2539 if (!
result.attributes.get(
"orderedMetadata"))
2542 result.addTypes(resultTypes);
2543 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2545 static_cast<int32_t>(frags[0].regs.size()),
2546 static_cast<int32_t>(frags[1].regs.size()),
2547 static_cast<int32_t>(frags[2].regs.size()),
2560void MmaSpBlockScaleOp::build(
2566 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2567 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2568 MMABlockScaleKind kind) {
2569 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2572 builder,
result,
shape, scaleVecSize, blockScaleFormat, kind);
2575 result.addOperands(operandA);
2576 result.addOperands(operandB);
2577 result.addOperands(operandC);
2578 result.addOperands({sparseMetadata, sparsitySelector, scaleAData, byteIdA,
2579 threadIdA, scaleBData, byteIdB, threadIdB});
2582 multiplicandPtxTypes);
2584 result.addTypes(resultType);
2585 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2587 static_cast<int32_t>(operandA.size()),
2588 static_cast<int32_t>(operandB.size()),
2589 static_cast<int32_t>(operandC.size()),
2603 auto curOp = cast<NVVM::MmaSpBlockScaleOp>(op);
2607 for (
Value operand : curOp.getOperandA())
2609 for (
Value operand : curOp.getOperandB())
2611 for (
Value operand : curOp.getOperandC())
2615 args.push_back(mt.
lookupValue(curOp.getSparseMetadata()));
2616 args.push_back(mt.
lookupValue(curOp.getSparsitySelector()));
2619 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2620 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2621 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2622 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2623 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2624 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2626 unsigned intId = MmaSpBlockScaleOp::getIntrinsicID(
2627 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2628 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2630 curOp.getBlockScaleFormat(), curOp.getKind());
2632 return {intId, args};
2635LogicalResult MmaSpBlockScaleOp::verify() {
2637 if (!getOrderedMetadata()) {
2638 return emitOpError(
"'orderedMetadata' attribute is mandatory");
2646 if (m == 16 && n == 8 && k == 128) {
2647 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2648 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2650 "unsupported MMATypes attribute for mma.m16n8k128.(mxf4nvf4|mxf4)");
2651 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2652 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2654 "unsupported ScaleVecSize attribute for mma.m16n8k128.mxf4");
2655 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2657 "unsupported BlockScaleFormat attribute for mma.m16n8k128.mxf4");
2658 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2659 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2660 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2661 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2662 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2663 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2664 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2665 "attributes for mma.m16n8k128.mxf4nvf4");
2667 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k128");
2669 }
else if (m == 16 && n == 8 && k == 64) {
2670 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2671 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2672 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2674 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2675 "attributes for mma.m16n8k64");
2677 result = emitOpError(
"unsupported Geom for sparse mma with block scaling");
2682LogicalResult ShflOp::verify() {
2683 auto returnStructType = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
2685 auto verifyTypeError = [&](Twine desc,
Type expectedType,
2686 Type actualType) -> LogicalResult {
2687 return emitOpError(
"expected " + desc +
" to be of type ")
2688 << expectedType <<
" but got " << actualType <<
" instead";
2691 if (returnStructType) {
2692 if (!getReturnValueAndIsValid())
2693 return emitOpError(
"\"return_value_and_is_valid\" attribute must be "
2694 "specified when the return type is a struct type");
2696 if (returnStructType.getBody().size() != 2)
2697 return emitOpError(
"expected return type to be a two-element struct");
2700 auto resultType = returnStruct[0];
2701 if (resultType != getVal().
getType())
2702 return verifyTypeError(
"first element in the returned struct",
2703 getVal().
getType(), resultType);
2705 auto predicateType = returnStruct[1];
2706 if (!predicateType.isInteger(1))
2707 return verifyTypeError(
"second element in the returned struct",
2711 if (getReturnValueAndIsValid())
2712 return emitOpError(
"expected return type to be a two-element struct");
2715 return verifyTypeError(
"return type", getVal().
getType(),
getType());
2721ShflOp::inferReturnTypes(
MLIRContext *context, std::optional<Location> location,
2722 ShflOp::Adaptor adaptor,
2724 Type valType = adaptor.getVal().getType();
2725 if (adaptor.getReturnValueAndIsValid())
2726 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2727 context, {valType, IntegerType::get(context, 1)}));
2729 inferredReturnTypes.push_back(valType);
2734 NVVM::MMAFrag frag,
int nRow,
2737 unsigned numberElements = 0;
2740 Type f16x2 = VectorType::get(2, builder.getF16Type());
2741 if (type == NVVM::MMATypes::f16) {
2742 elementType = f16x2;
2743 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2747 }
else if (type == NVVM::MMATypes::f32) {
2748 elementType = builder.getF32Type();
2750 }
else if (type == NVVM::MMATypes::f64) {
2751 elementType = builder.getF64Type();
2752 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2756 }
else if (type == NVVM::MMATypes::tf32) {
2757 elementType = builder.getI32Type();
2759 }
else if (type == NVVM::MMATypes::s8 || type == NVVM::MMATypes::u8) {
2760 elementType = builder.getI32Type();
2761 int parallelSize = 0;
2762 if (frag == NVVM::MMAFrag::a)
2763 parallelSize = nRow;
2764 if (frag == NVVM::MMAFrag::b)
2765 parallelSize = nCol;
2768 if (parallelSize == 16)
2771 else if (parallelSize == 8)
2773 else if (parallelSize == 32)
2775 }
else if (type == NVVM::MMATypes::s32) {
2776 elementType = builder.getI32Type();
2779 assert(numberElements != 0 && elementType !=
nullptr);
2780 return std::make_pair(elementType, numberElements);
2783static std::pair<mlir::Type, unsigned>
2787 if (frag == NVVM::MMAFrag::a) {
2790 }
else if (frag == NVVM::MMAFrag::b) {
2797 assert(nRow && nCol);
2801LogicalResult NVVM::WMMALoadOp::verify() {
2802 unsigned addressSpace =
2803 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2804 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2805 addressSpace != NVVMMemorySpace::Shared)
2806 return emitOpError(
"expected source pointer in memory "
2809 if (NVVM::WMMALoadOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2810 getEltype(), getFrag()) == 0)
2811 return emitOpError() <<
"invalid attribute combination";
2816 if (typeInfo.first == f64Ty && typeInfo.second == 1) {
2818 return emitOpError(
"expected destination type to be f64");
2822 Type dstType = LLVM::LLVMStructType::getLiteral(
2825 return emitOpError(
"expected destination type is a structure of ")
2826 << typeInfo.second <<
" elements of type " << typeInfo.first;
2830LogicalResult NVVM::WMMAStoreOp::verify() {
2831 unsigned addressSpace =
2832 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2833 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2834 addressSpace != NVVMMemorySpace::Shared)
2835 return emitOpError(
"expected operands to be a source pointer in memory "
2838 if (NVVM::WMMAStoreOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2840 return emitOpError() <<
"invalid attribute combination";
2843 if (getArgs().size() != typeInfo.second)
2844 return emitOpError() <<
"expected " << typeInfo.second <<
" data operands";
2845 if (llvm::any_of(getArgs(), [&typeInfo](
Value operands) {
2846 return operands.
getType() != typeInfo.first;
2848 return emitOpError() <<
"expected data operands of type " << typeInfo.first;
2852LogicalResult NVVM::WMMAMmaOp::verify() {
2853 if (NVVM::WMMAMmaOp::getIntrinsicID(
getM(),
getN(), getK(), getLayoutA(),
2854 getLayoutB(), getEltypeA(),
2856 return emitOpError() <<
"invalid attribute combination";
2864 arguments.append(typeInfoA.second, typeInfoA.first);
2865 arguments.append(typeInfoB.second, typeInfoB.first);
2866 arguments.append(typeInfoC.second, typeInfoC.first);
2867 unsigned numArgs = arguments.size();
2868 if (getArgs().size() != numArgs)
2869 return emitOpError() <<
"expected " << numArgs <<
" arguments";
2870 for (
unsigned i = 0; i < numArgs; i++) {
2871 if (getArgs()[i].
getType() != arguments[i])
2872 return emitOpError() <<
"expected argument " << i <<
" to be of type "
2875 Type dstType = LLVM::LLVMStructType::getLiteral(
2878 return emitOpError(
"expected destination type is a structure of ")
2879 << typeInfoC.second <<
" elements of type " << typeInfoC.first;
2883LogicalResult NVVM::LdMatrixOp::verify() {
2885 if (m == 8 && n == 8) {
2886 if (num != 1 && num != 2 && num != 4) {
2887 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x8 "
2890 if (getEltType() != LdStMatrixEltType::B16) {
2891 return emitOpError(
"expected element type to be b16 for 8x8 matrix");
2893 }
else if (m == 8 && n == 16) {
2894 if (num != 1 && num != 2 && num != 4) {
2895 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x16 "
2898 if (getLayout() != MMALayout::row) {
2899 return emitOpError(
"expected layout to be row for 8x16 matrix");
2901 if (getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2902 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2903 return emitOpError(
"expected element type to be b8x16.b4x16_p64 or "
2904 "b8x16.b6x16_p32 for 8x16 matrix");
2906 }
else if (m == 16 && n == 16) {
2907 if (num != 1 && num != 2) {
2908 return emitOpError(
"expected num attribute to be 1 or 2 for 16x16 "
2911 if (getLayout() != MMALayout::col) {
2912 return emitOpError(
"expected layout to be col for 16x16 matrix");
2914 if (getEltType() != LdStMatrixEltType::B8 &&
2915 getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2916 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2917 return emitOpError(
"expected element type to be b8, b8x16.b4x16_p64 or "
2918 "b8x16.b6x16_p32 for 16x16 matrix");
2921 return emitOpError(
"expected shape to be 8x8, 8x16 or 16x16");
2925 uint32_t numElements = (m == 16 && n == 16 ? num * 2 : num);
2926 if (numElements == 1 &&
getType() != i32)
2927 return emitOpError(
"expected destination type is i32");
2928 if (numElements == 2 || numElements == 4) {
2929 Type dstType = LLVM::LLVMStructType::getLiteral(
2932 return emitOpError(
"expected destination type is a structure of ")
2933 << numElements <<
" elements of type i32";
2939LogicalResult LdMatrixOp::inferReturnTypes(
2940 MLIRContext *context, std::optional<Location> location,
2942 uint32_t num = adaptor.getNum();
2943 uint32_t m = adaptor.getShape().getM();
2944 uint32_t n = adaptor.getShape().getN();
2945 uint32_t numElements = (m == 16 && n == 16) ? num * 2 : num;
2947 Type i32 = IntegerType::get(context, 32);
2948 if (numElements == 1)
2949 inferredReturnTypes.push_back(i32);
2951 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2956LogicalResult NVVM::StMatrixOp::verify() {
2957 int numMatrix = getSources().size();
2958 if (numMatrix != 1 && numMatrix != 2 && numMatrix != 4)
2959 return emitOpError(
"expected num attribute to be 1, 2 or 4");
2962 if (m == 8 && n == 8) {
2963 if (getEltType() != NVVM::LdStMatrixEltType::B16) {
2964 return emitOpError(
"expected element type to be B16 for 8x8 matrix");
2966 }
else if (m == 16 && n == 8) {
2967 if (getEltType() != NVVM::LdStMatrixEltType::B8) {
2968 return emitOpError(
"expected element type to be B8 for 16x8 matrix");
2970 if (getLayout() != NVVM::MMALayout::col) {
2971 return emitOpError(
"expected layout to be col for 16x8 matrix");
2974 return emitOpError(
"expected shape to be 8x8 or 16x8");
2980LogicalResult NVVM::MovMatrixOp::verify() {
2982 if (m != 8 || n != 8)
2983 return emitOpError(
"expected shape to be 8x8");
2984 if (getLayout() != NVVM::MMALayout::col)
2985 return emitOpError(
"expected layout to be col");
2986 if (getEltType() != NVVM::LdStMatrixEltType::B16)
2987 return emitOpError(
"expected element type to be b16");
2992 if (typeA == NVVM::WGMMATypes::tf32)
2994 if (typeA == NVVM::WGMMATypes::f16 || typeA == NVVM::WGMMATypes::bf16)
2996 if (typeA == NVVM::WGMMATypes::s8 || typeA == NVVM::WGMMATypes::u8)
2998 if (typeA == NVVM::WGMMATypes::e4m3 || typeA == NVVM::WGMMATypes::e5m2)
3000 if (typeA == NVVM::WGMMATypes::b1)
3006 NVVM::WGMMATypes typeA,
3007 NVVM::WGMMATypes typeB) {
3009 case NVVM::WGMMATypes::f16:
3010 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3011 typeB == NVVM::WGMMATypes::f16)
3014 case NVVM::WGMMATypes::tf32:
3015 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::tf32)
3018 case NVVM::WGMMATypes::u8:
3019 case NVVM::WGMMATypes::s8:
3020 if (typeD == NVVM::WGMMATypes::s32 &&
3021 (typeB == NVVM::WGMMATypes::u8 || typeB == NVVM::WGMMATypes::s8))
3024 case NVVM::WGMMATypes::b1:
3025 if (typeD == NVVM::WGMMATypes::s32 && typeB == NVVM::WGMMATypes::b1)
3028 case NVVM::WGMMATypes::bf16:
3029 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::bf16)
3032 case NVVM::WGMMATypes::e4m3:
3033 case NVVM::WGMMATypes::e5m2:
3034 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3035 (typeB == NVVM::WGMMATypes::e5m2 || typeB == NVVM::WGMMATypes::e4m3))
3038 case WGMMATypes::f32:
3039 case WGMMATypes::s32:
3040 llvm_unreachable(
"unsupported input types");
3048 72, 80, 88, 96, 104, 112, 120, 128,
3049 136, 144, 152, 160, 168, 176, 184, 192,
3050 200, 208, 216, 224, 232, 240, 248, 256};
3052 80, 96, 112, 128, 144, 160,
3053 176, 192, 208, 224, 240, 256};
3055 case WGMMATypes::f16:
3056 case WGMMATypes::tf32:
3057 case WGMMATypes::bf16:
3058 case WGMMATypes::e4m3:
3059 case WGMMATypes::e5m2:
3060 if (llvm::is_contained(allowedN, sizeN))
3063 case WGMMATypes::u8:
3064 case WGMMATypes::s8:
3065 case WGMMATypes::b1:
3066 if (llvm::is_contained(allowedNshort, sizeN))
3069 case WGMMATypes::f32:
3070 case WGMMATypes::s32:
3071 llvm_unreachable(
"unsupported input types");
3077LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
3078 Value outValue = getResults();
3079 auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.
getType());
3081 return emitOpError() <<
"expected results to be struct";
3082 int outputSize = stype.getBody().size();
3083 WGMMATypes typeD = getTypeD();
3084 WGMMATypes typeA = getTypeA();
3085 WGMMATypes typeB = getTypeB();
3087 for (
Type t : stype.getBody()) {
3088 if (t != stype.getBody().front())
3089 return emitOpError()
3090 <<
"all elements in struct must be same type but there is " << t;
3093 if (typeD != WGMMATypes::f32 && typeD != WGMMATypes::f16 &&
3094 typeD != WGMMATypes::s32) {
3095 return emitOpError() <<
"does not support the given output type " << typeD;
3097 if (typeD == WGMMATypes::s32 &&
3098 (getScaleA() == WGMMAScaleIn::neg || getScaleB() == WGMMAScaleIn::neg)) {
3099 return emitOpError() <<
"has s32 output, scaleA and scaleB cannot be neg";
3103 return emitOpError() << typeD <<
" += " << typeA <<
" * " << typeB
3104 <<
", it is not supported.";
3109 return emitOpError() <<
"shape 'm' must be 64";
3114 return emitOpError() <<
"shape 'k' must be " << allowedK.value()
3115 <<
" for input type " << typeA;
3119 return emitOpError() <<
"has input type " << typeA <<
" n is set to "
3120 <<
getShape().getN() <<
", it is not supported.";
3127 if ((typeA != WGMMATypes::f16 && typeA != WGMMATypes::bf16) &&
3128 (getLayoutA() == mlir::NVVM::MMALayout::col ||
3129 getLayoutB() == mlir::NVVM::MMALayout::row)) {
3130 return emitOpError()
3131 <<
"given layouts layout_a = " << getLayoutA()
3132 <<
" and layout_b = " << getLayoutB() <<
" for input types " << typeA
3134 <<
" requires transpose. However, this is only supported for: "
3135 << MMATypes::f16 <<
" and " << MMATypes::bf16;
3139 int expectedOutput = 0;
3140 if (typeD == WGMMATypes::f32 || typeD == WGMMATypes::s32)
3141 expectedOutput =
getShape().getN() / 2;
3142 if (typeD == WGMMATypes::f16)
3143 expectedOutput =
getShape().getN() / 4;
3144 if (outputSize != expectedOutput) {
3145 return emitOpError() <<
"results " << expectedOutput
3146 <<
", however output struct has " << outputSize
3150 if (typeD != WGMMATypes::s32 &&
3151 getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3152 NVVM::MMAIntOverflow::satfinite) {
3153 return emitOpError()
3154 <<
" `satfinite` can be only used with s32 accumulator, however "
3155 "the current accumulator is "
3162std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
3165 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3167 StringRef outputTypeName = stringifyWGMMATypes(getTypeD());
3169 int expectedOutputRegisters = 0;
3170 if (getTypeD() == WGMMATypes::f16)
3171 expectedOutputRegisters =
getShape().getN() / 4;
3173 expectedOutputRegisters =
getShape().getN() / 2;
3176 llvm::raw_string_ostream ss(ptx);
3181 << ((expectedOutputRegisters * 2) + 2)
3183 "wgmma.mma_async.sync.aligned.m"
3184 << m <<
"n" << n <<
"k" << k <<
"." << outputTypeName <<
"." << getTypeA()
3185 <<
"." << getTypeB();
3186 if (getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3187 NVVM::MMAIntOverflow::satfinite)
3191 for (; regCnt < expectedOutputRegisters; ++regCnt) {
3192 ss <<
"$" << regCnt;
3193 if (regCnt != expectedOutputRegisters - 1)
3199 regCnt = (regCnt * 2);
3200 ss <<
" $" << (regCnt) <<
"," <<
" $" << (regCnt + 1) <<
"," <<
" p";
3201 if (getTypeD() != WGMMATypes::s32) {
3202 ss <<
", $" << (regCnt + 3) <<
", $" << (regCnt + 4);
3206 ss <<
", $" << (regCnt + 5) <<
", $" << (regCnt + 6);
3213bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
3217 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3224 asmValues.push_back({makeConstantI32(rewriter,
static_cast<int>(getScaleD())),
3226 if (getTypeD() != WGMMATypes::s32) {
3227 asmValues.push_back(
3228 {makeConstantI32(rewriter,
3229 getScaleA() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3231 asmValues.push_back(
3232 {makeConstantI32(rewriter,
3233 getScaleB() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3237 asmValues.push_back(
3238 {makeConstantI32(rewriter,
static_cast<int>(getLayoutA())),
3240 asmValues.push_back(
3241 {makeConstantI32(rewriter, 1 -
static_cast<int>(getLayoutB())),
3247LogicalResult NVVM::FenceProxyOp::verify() {
3248 if (getKind() == NVVM::ProxyKind::async_shared && !getSpace().has_value()) {
3249 return emitOpError() <<
"async_shared fence requires space attribute";
3251 if (getKind() != NVVM::ProxyKind::async_shared && getSpace().has_value()) {
3252 return emitOpError() <<
"only async_shared fence can have space attribute";
3257LogicalResult NVVM::FenceProxyAcquireOp::verify() {
3258 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3259 return emitOpError(
"uni-directional proxies only support generic for "
3260 "from_proxy attribute");
3262 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3263 return emitOpError(
"uni-directional proxies only support tensormap "
3264 "for to_proxy attribute");
3268LogicalResult NVVM::FenceProxyReleaseOp::verify() {
3269 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3270 return emitOpError(
"uni-directional proxies only support generic for "
3271 "from_proxy attribute");
3273 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3274 return emitOpError(
"uni-directional proxies only support tensormap "
3275 "for to_proxy attribute");
3279LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
3280 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3281 return emitOpError(
"only generic is support for from_proxy attribute");
3283 if (getToProxy() != NVVM::ProxyKind::async)
3284 return emitOpError(
"only async is supported for to_proxy attribute");
3288LogicalResult NVVM::SetMaxRegisterOp::verify() {
3289 if (getRegCount() % 8)
3290 return emitOpError(
"new register size must be multiple of 8");
3291 if (getRegCount() < 24 || getRegCount() > 256)
3292 return emitOpError(
"new register size must be in between 24 to 256");
3296LogicalResult NVVM::Tcgen05CpOp::verify() {
3297 auto mc = getMulticast();
3299 using SH = Tcgen05CpShape;
3300 using MC = Tcgen05CpMulticast;
3302 case SH::SHAPE_128x256b:
3303 case SH::SHAPE_128x128b:
3304 case SH::SHAPE_4x256b:
3306 return emitError(
"Invalid multicast type for tcgen05.cp Op");
3308 case SH::SHAPE_64x128b:
3309 if (mc != MC::WARPX2_01_23 && mc != MC::WARPX2_02_13)
3310 return emitError(
"Shape 64x128b requires multicast warpx2_01_23 or "
3311 "warpx2_02_13 for tcgen05.cp Op");
3313 case SH::SHAPE_32x128b:
3314 if (mc != MC::WARPX4)
3316 "Shape 32x128b requires multicast warpx4 for tcgen05.cp Op");
3322LogicalResult NVVM::MatchSyncOp::verify() {
3323 if (getKind() == NVVM::MatchSyncKind::all) {
3324 auto type = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
3325 if (!type || type.getBody().size() != 2 ||
3326 !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
3327 return emitOpError(
"match.sync 'all' returns a two element struct with "
3328 "first element as i32 and second element as i1");
3331 if (!
getType().isInteger(32)) {
3332 return emitOpError(
"match.sync 'any' returns an i32");
3338LogicalResult MatchSyncOp::inferReturnTypes(
3339 MLIRContext *context, std::optional<Location> location,
3341 if (adaptor.getKind() == NVVM::MatchSyncKind::all)
3342 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
3344 {IntegerType::get(context, 32), IntegerType::get(context, 1)}));
3346 inferredReturnTypes.push_back(IntegerType::get(context, 32));
3350LogicalResult NVVM::VoteSyncOp::verify() {
3351 if (getKind() == NVVM::VoteSyncKind::ballot) {
3352 if (!
getType().isInteger(32)) {
3353 return emitOpError(
"vote.sync 'ballot' returns an i32");
3356 if (!
getType().isInteger(1)) {
3357 return emitOpError(
"vote.sync 'any', 'all' and 'uni' returns an i1");
3363LogicalResult VoteSyncOp::inferReturnTypes(
3364 MLIRContext *context, std::optional<Location> location,
3366 unsigned width = adaptor.getKind() == NVVM::VoteSyncKind::ballot ? 32 : 1;
3367 inferredReturnTypes.push_back(IntegerType::get(context, width));
3371LogicalResult NVVM::PrefetchOp::verify() {
3372 using MemSpace = NVVM::NVVMMemorySpace;
3373 using CacheLevel = NVVM::PrefetchCacheLevel;
3375 unsigned addressSpace =
3376 llvm::cast<LLVM::LLVMPointerType>(getAddr().
getType()).getAddressSpace();
3377 std::optional<NVVM::CacheEvictionPriority> evictPriority = getEvictPriority();
3378 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = getCacheLevel();
3380 if (getTensormap() && cacheLevel)
3381 return emitOpError(
"cannot specify both tensormap and cache level");
3383 if (getTensormap()) {
3384 if (addressSpace != MemSpace::Generic &&
3385 addressSpace != MemSpace::Constant) {
3387 "prefetch tensormap requires a generic or constant pointer");
3390 if (evictPriority) {
3392 "prefetch tensormap does not support eviction priority");
3395 if (getInParamSpace() && addressSpace != MemSpace::Generic) {
3397 "in_param_space can only be specified for a generic pointer");
3400 }
else if (cacheLevel) {
3401 if (addressSpace != MemSpace::Generic && addressSpace != MemSpace::Global &&
3402 addressSpace != MemSpace::Local) {
3403 return emitOpError(
"prefetch to cache level requires a generic, global, "
3404 "or local pointer");
3408 if (*cacheLevel != CacheLevel::L1) {
3410 "unsupported cache level, the only supported uniform "
3411 "cache level is L1");
3414 if (addressSpace != MemSpace::Generic) {
3416 "prefetch to uniform cache requires a generic pointer");
3420 if (evictPriority) {
3421 if (*cacheLevel != CacheLevel::L2)
3423 "cache eviction priority supported only for cache level L2");
3425 if (addressSpace != MemSpace::Global)
3426 return emitOpError(
"cache eviction priority requires a global pointer");
3428 if (*evictPriority != NVVM::CacheEvictionPriority::EvictNormal &&
3429 *evictPriority != NVVM::CacheEvictionPriority::EvictLast)
3431 "unsupported cache eviction priority, only evict_last and "
3432 "evict_normal are supported");
3436 return emitOpError(
"predicate supported only on prefetch tensormap");
3440 "requires specification of either cache level or tensormap");
3446LogicalResult NVVM::ClusterLaunchControlQueryCancelOp::verify() {
3447 switch (getQueryType()) {
3448 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
3450 return emitOpError(
"is_canceled query type returns an i1");
3452 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
3453 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
3454 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
3455 if (!
getType().isInteger(32)) {
3456 return emitOpError(
"get_first_cta_id_x, get_first_cta_id_y, "
3457 "get_first_cta_id_z query types return an i32");
3464LogicalResult ClusterLaunchControlQueryCancelOp::inferReturnTypes(
3465 MLIRContext *context, std::optional<Location> location,
3466 ClusterLaunchControlQueryCancelOp::Adaptor adaptor,
3469 adaptor.getQueryType() == NVVM::ClusterLaunchControlQueryType::IS_CANCELED
3472 inferredReturnTypes.push_back(IntegerType::get(context, width));
3476LogicalResult NVVM::ReduxOp::verify() {
3479 if (!reduxType.
isF32()) {
3481 return emitOpError(
"abs attribute is supported only for f32 type");
3483 return emitOpError(
"nan attribute is supported only for f32 type");
3486 NVVM::ReductionKind kind = getKind();
3488 case NVVM::ReductionKind::ADD:
3489 case NVVM::ReductionKind::AND:
3490 case NVVM::ReductionKind::OR:
3491 case NVVM::ReductionKind::XOR:
3492 case NVVM::ReductionKind::MAX:
3493 case NVVM::ReductionKind::MIN:
3494 case NVVM::ReductionKind::UMAX:
3495 case NVVM::ReductionKind::UMIN:
3497 return emitOpError(
"'")
3498 << kind <<
"' reduction kind unsupported with " << reduxType
3499 <<
" type. Only supported type is 'i32'.";
3501 case NVVM::ReductionKind::FMIN:
3502 case NVVM::ReductionKind::FMAX:
3503 if (!reduxType.isF32())
3504 return emitOpError(
"'")
3505 << kind <<
"' reduction kind unsupported with " << reduxType
3506 <<
" type. Only supported type is 'f32'.";
3513LogicalResult NVVM::TensormapReplaceOp::verify() {
3514 auto ord = getOrd();
3515 Value newVal = getNewValue();
3516 auto newValAttr = getNewValueAttr();
3517 auto fieldName = stringifyEnum(getField());
3519 if (ord && !llvm::is_contained({NVVM::TensormapField::BOX_DIM,
3520 NVVM::TensormapField::GLOBAL_DIM,
3521 NVVM::TensormapField::GLOBAL_STRIDE,
3522 NVVM::TensormapField::ELEMENT_STRIDE},
3524 return emitOpError(
"ordinal is not supported for ")
3525 << fieldName <<
" field";
3527 auto invalidNewVal = [&](llvm::Twine type) -> std::string {
3528 return llvm::Twine(
"new_value must be specified and must be an " + type +
3529 " for " + llvm::Twine(fieldName) +
" field")
3533 auto invalidNewValAttr = [&]() -> std::string {
3534 return (llvm::Twine(
3535 "new_value_attr must be specified and must be a valid ") +
3536 llvm::Twine(fieldName) +
" attribute for " + fieldName +
" field")
3540 switch (getField()) {
3541 case NVVM::TensormapField::GLOBAL_ADDRESS:
3543 return emitOpError(invalidNewVal(
"i64"));
3545 case NVVM::TensormapField::RANK:
3547 return emitOpError(invalidNewVal(
"i32"));
3549 case NVVM::TensormapField::GLOBAL_STRIDE:
3551 return emitOpError(
"ordinal is required for global_stride field");
3553 return emitOpError(invalidNewVal(
"i64"));
3555 case NVVM::TensormapField::BOX_DIM:
3556 case NVVM::TensormapField::GLOBAL_DIM:
3557 case NVVM::TensormapField::ELEMENT_STRIDE:
3559 return emitOpError(
"ordinal is required for ")
3560 << stringifyEnum(getField()) <<
" field";
3562 return emitOpError(invalidNewVal(
"i32"));
3564 case NVVM::TensormapField::ELEMTYPE:
3565 if (!(newValAttr && llvm::isa<TensormapElemtypeAttr>(*newValAttr)))
3566 return emitOpError(invalidNewValAttr());
3568 case NVVM::TensormapField::INTERLEAVE_LAYOUT:
3569 if (!(newValAttr && llvm::isa<TensormapInterleaveLayoutAttr>(*newValAttr)))
3570 return emitOpError(invalidNewValAttr());
3572 case NVVM::TensormapField::SWIZZLE_MODE:
3573 if (!(newValAttr && llvm::isa<TensormapSwizzleModeAttr>(*newValAttr)))
3574 return emitOpError(invalidNewValAttr());
3576 case NVVM::TensormapField::SWIZZLE_ATOMICITY:
3577 if (!(newValAttr && llvm::isa<TensormapSwizzleAtomicityAttr>(*newValAttr)))
3578 return emitOpError(invalidNewValAttr());
3580 case NVVM::TensormapField::FILL_MODE:
3581 if (!(newValAttr && llvm::isa<TensormapFillModeAttr>(*newValAttr)))
3582 return emitOpError(invalidNewValAttr());
3589template <
typename OpType>
3591 mlir::NVVM::FPRoundingMode rndMode = op.getRnd();
3592 mlir::NVVM::SaturationMode satMode = op.getSat();
3593 bool isFTZ = op.getFtz();
3596 mlir::Type opBaseType = isa<VectorType>(opType)
3597 ? cast<VectorType>(opType).getElementType()
3600 if (opBaseType.
isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3601 return op.emitOpError(
"FTZ and saturation are not supported for "
3602 "additions/subtractions involving f64 type");
3604 if (opBaseType.
isF16() && !(rndMode == NVVM::FPRoundingMode::RN ||
3605 rndMode == NVVM::FPRoundingMode::NONE))
3606 return op.emitOpError(
"only RN rounding mode is supported for f16 and "
3607 "vector<2xf16> additions/subtractions");
3609 if (opBaseType.
isBF16()) {
3610 if (rndMode != NVVM::FPRoundingMode::RN &&
3611 rndMode != NVVM::FPRoundingMode::NONE)
3612 return op.emitOpError(
"only RN rounding mode is supported for bf16 and "
3613 "vector<2xbf16> additions/subtractions");
3614 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3615 return op.emitOpError(
"FTZ and saturation are not supported for bf16 and "
3616 "vector<2xbf16> additions/subtractions");
3626LogicalResult NVVM::FmaOp::verify() {
3627 auto opType = getRes().getType();
3628 mlir::NVVM::FPRoundingMode rndMode = getRnd();
3629 mlir::NVVM::SaturationMode satMode = getSat();
3630 bool isFTZ = getFtz();
3631 bool isRelu = getRelu();
3632 bool hasOOB = getOob();
3634 auto getBaseFType = [](
Type type) ->
Type {
3635 if (isa<VectorType>(type))
3636 return cast<VectorType>(type).getElementType();
3640 auto opBaseType = getBaseFType(opType);
3642 if (rndMode == NVVM::FPRoundingMode::NONE)
3643 return emitOpError(
"rounding mode must be specified");
3645 if (isRelu && satMode == NVVM::SaturationMode::SAT)
3646 return emitOpError(
"relu and saturation are not supported together");
3648 if (hasOOB && (satMode == NVVM::SaturationMode::SAT || isFTZ))
3649 return emitOpError(
"oob is not supported with saturation or FTZ");
3651 if (!(opBaseType.isF16() || opBaseType.isBF16()) && (isRelu || hasOOB))
3652 return emitOpError(
"relu and oob are only supported for f16 and bf16");
3654 if (opBaseType.isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3655 return emitOpError(
"FTZ and saturation are not supported for f64 type");
3657 if (opBaseType.isF16() && rndMode != NVVM::FPRoundingMode::RN)
3659 "only RN rounding mode is supported for f16 and vector<2xf16>");
3661 if (opBaseType.isBF16()) {
3662 if (rndMode != NVVM::FPRoundingMode::RN)
3664 "only RN rounding mode is supported for bf16 and vector<2xbf16>");
3665 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3667 "FTZ and saturation are not supported for bf16 and vector<2xbf16>");
3673LogicalResult NVVM::SqrtOp::verify() {
3674 if (getRnd() == NVVM::FPRoundingMode::NONE)
3675 return emitOpError(
"rounding mode cannot be None");
3677 if (getRes().
getType().isF64() && getFtz())
3678 return emitOpError(
"FTZ is not supported for f64");
3683LogicalResult NVVM::DivFOp::verify() {
3684 bool isApprox = getApprox();
3685 bool isFull = getFull();
3686 bool isF64 = getRes().getType().isF64();
3687 bool isFtz = getFtz();
3688 NVVM::FPRoundingMode rndMode = getRnd();
3690 if (isApprox && isFull)
3691 return emitOpError(
"'approx' and 'full' are mutually exclusive");
3693 if (isApprox || isFull) {
3695 return emitOpError(
"'approx' and 'full' forms are f32-only");
3696 if (rndMode != NVVM::FPRoundingMode::NONE)
3698 "'approx' and 'full' forms do not accept a rounding mode");
3703 if (rndMode == NVVM::FPRoundingMode::NONE)
3704 return emitOpError(
"rounding mode cannot be None for the rounded divide");
3706 return emitOpError(
"FTZ is not supported for f64");
3717 unsigned sizeInBits,
3719 field = builder.CreateZExtOrBitCast(field, builder.getInt32Ty());
3721 unsigned mask = (sizeInBits < 32 ? ((1u << sizeInBits) - 1) : 0xffffffffu);
3722 if (mask != 0xffffffffu)
3723 field = builder.CreateAnd(field, builder.getInt32(mask));
3725 field = builder.CreateZExtOrBitCast(field, builder.getInt64Ty());
3726 field = builder.CreateShl(field, start);
3728 return builder.CreateOr(
result, field);
3731void Tcgen05MmaSmemDescOp::createSmemDescriptor(
Operation &op,
3733 llvm::IRBuilderBase &builder) {
3734 auto thisOp = cast<NVVM::Tcgen05MmaSmemDescOp>(op);
3735 llvm::Value *smemDesc = builder.getInt64(0);
3740 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimOffset()), 14, 16);
3742 builder, smemDesc, mt.
lookupValue(thisOp.getStrideDimOffset()), 14, 32);
3748 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimMode()), 1, 52);
3752 mt.
mapValue(thisOp.getRes()) = smemDesc;
3759std::string NVVM::MBarrierInitOp::getPtx() {
3762 std::string layout =
3763 getLayout() == 1 ? std::string(
".layout::v1") : std::string();
3765 return llvm::formatv(
"mbarrier.init{0}{1}.b64 [%0], %1;", layout, space)
3769std::string NVVM::MBarrierArriveExpectTxOp::getPtx() {
3772 ? std::string(
"mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;")
3773 : std::string(
"mbarrier.arrive.expect_tx.b64 _, [%0], %1;");
3776std::string NVVM::MBarrierTryWaitParityOp::getPtx() {
3778 llvm::StringRef space = isShared ?
".shared" :
"";
3780 return llvm::formatv(
"{\n\t"
3781 ".reg .pred P1; \n\t"
3783 "mbarrier.try_wait.parity{0}.b64 P1, [%0], %1, %2; \n\t"
3784 "@P1 bra.uni DONE; \n\t"
3785 "bra.uni LAB_WAIT; \n\t"
3802 LLVM::FNegOp::create(rewriter, loc, op.getRhs().getType(), op.getRhs());
3805 op.getRnd(), op.getSat(), op.getFtz());
3824 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count
3825 : llvm::Intrinsic::nvvm_barrier_cta_sync_count;
3827 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all
3828 : llvm::Intrinsic::nvvm_barrier_cta_sync_all;
3833static llvm::Intrinsic::ID
3836 case NVVM::BarrierReduction::AND:
3837 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_and_aligned_all
3838 : llvm::Intrinsic::nvvm_barrier_cta_red_and_all;
3839 case NVVM::BarrierReduction::OR:
3840 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_or_aligned_all
3841 : llvm::Intrinsic::nvvm_barrier_cta_red_or_all;
3842 case NVVM::BarrierReduction::POPC:
3843 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_popc_aligned_all
3844 : llvm::Intrinsic::nvvm_barrier_cta_red_popc_all;
3846 llvm_unreachable(
"unknown BarrierReduction kind");
3851 auto thisOp = cast<NVVM::BarrierOp>(op);
3852 llvm::Value *barrierId = thisOp.getBarrierId()
3854 : builder.getInt32(0);
3855 bool hasCount =
static_cast<bool>(thisOp.getNumberOfThreads());
3856 llvm::Intrinsic::ID
id =
3860 args.push_back(mt.
lookupValue(thisOp.getNumberOfThreads()));
3861 return {id, std::move(args)};
3866 auto thisOp = cast<NVVM::BarrierArriveOp>(op);
3867 llvm::Value *barrierId = thisOp.getBarrierId()
3869 : builder.getInt32(0);
3870 llvm::Value *numThreads = mt.
lookupValue(thisOp.getNumberOfThreads());
3871 llvm::Intrinsic::ID
id =
3873 ? llvm::Intrinsic::nvvm_barrier_cta_arrive_aligned_count
3874 : llvm::Intrinsic::nvvm_barrier_cta_arrive_count;
3875 return {id, {barrierId, numThreads}};
3880 auto thisOp = cast<NVVM::BarrierReductionOp>(op);
3882 thisOp.getAligned(), thisOp.getReductionOp());
3883 llvm::Value *barrierId = thisOp.getBarrierId()
3885 : builder.getInt32(0);
3888 builder.CreateICmpNE(mt.
lookupValue(thisOp.getReductionPredicate()),
3889 builder.getInt32(0))};
3890 return {id, std::move(args)};
3895 llvm::IRBuilderBase &builder) {
3896 auto thisOp = cast<NVVM::CosOp>(op);
3897 llvm::Intrinsic::ID
id = thisOp.getFtz()
3898 ? llvm::Intrinsic::nvvm_cos_approx_ftz_f
3899 : llvm::Intrinsic::nvvm_cos_approx_f;
3905 llvm::IRBuilderBase &builder) {
3906 auto thisOp = cast<NVVM::SinOp>(op);
3907 llvm::Intrinsic::ID
id = thisOp.getFtz()
3908 ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
3909 : llvm::Intrinsic::nvvm_sin_approx_f;
3915 llvm::IRBuilderBase &builder) {
3916 auto thisOp = cast<NVVM::Log2Op>(op);
3917 llvm::Intrinsic::ID
id = thisOp.getFtz()
3918 ? llvm::Intrinsic::nvvm_lg2_approx_ftz_f
3919 : llvm::Intrinsic::nvvm_lg2_approx_f;
3925 llvm::IRBuilderBase &builder) {
3926 auto thisOp = cast<NVVM::Ex2Op>(op);
3927 llvm::Intrinsic::ID
id = thisOp.getFtz()
3928 ? llvm::Intrinsic::nvvm_ex2_approx_ftz
3929 : llvm::Intrinsic::nvvm_ex2_approx;
3933LogicalResult NVVM::Ex2Op::verify() {
3934 auto vectorType = dyn_cast<VectorType>(getSrc().
getType());
3938 if (vectorType.getElementType().isF16() && getFtz())
3939 return emitOpError(
"FTZ is not supported for vector<2xf16>");
3940 if (vectorType.getElementType().isBF16() && !getFtz())
3941 return emitOpError(
"FTZ is required for vector<2xbf16>");
3947 llvm::IRBuilderBase &builder) {
3948 auto thisOp = cast<NVVM::RsqrtOp>(op);
3949 Type t = thisOp.getRes().getType();
3950 bool isFtz = thisOp.getFtz();
3952 llvm::Intrinsic::ID
id = [&] {
3954 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_f
3955 : llvm::Intrinsic::nvvm_rsqrt_approx_f;
3958 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_d
3959 : llvm::Intrinsic::nvvm_rsqrt_approx_d;
3967 llvm::IRBuilderBase &builder) {
3968 auto thisOp = cast<NVVM::SqrtOp>(op);
3969 Type t = thisOp.getRes().getType();
3970 NVVM::FPRoundingMode rndMode = thisOp.getRnd();
3971 bool isFtz = thisOp.getFtz();
3975 unsigned rndIndex =
static_cast<unsigned>(rndMode) - 1;
3977 static constexpr llvm::Intrinsic::ID f32IDs[] = {
3978 llvm::Intrinsic::nvvm_sqrt_rn_f,
3979 llvm::Intrinsic::nvvm_sqrt_rm_f,
3980 llvm::Intrinsic::nvvm_sqrt_rp_f,
3981 llvm::Intrinsic::nvvm_sqrt_rz_f,
3983 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
3984 llvm::Intrinsic::nvvm_sqrt_rn_ftz_f,
3985 llvm::Intrinsic::nvvm_sqrt_rm_ftz_f,
3986 llvm::Intrinsic::nvvm_sqrt_rp_ftz_f,
3987 llvm::Intrinsic::nvvm_sqrt_rz_ftz_f,
3989 static constexpr llvm::Intrinsic::ID f64IDs[] = {
3990 llvm::Intrinsic::nvvm_sqrt_rn_d,
3991 llvm::Intrinsic::nvvm_sqrt_rm_d,
3992 llvm::Intrinsic::nvvm_sqrt_rp_d,
3993 llvm::Intrinsic::nvvm_sqrt_rz_d,
3996 llvm::Intrinsic::ID
id =
3997 t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4005 llvm::IRBuilderBase &builder) {
4006 auto thisOp = cast<NVVM::SqrtApproxOp>(op);
4007 llvm::Intrinsic::ID
id = thisOp.getFtz()
4008 ? llvm::Intrinsic::nvvm_sqrt_approx_ftz_f
4009 : llvm::Intrinsic::nvvm_sqrt_approx_f;
4015 llvm::IRBuilderBase &builder) {
4016 auto thisOp = cast<NVVM::DivFOp>(op);
4017 bool isFtz = thisOp.getFtz();
4019 llvm::Intrinsic::ID id;
4021 if (thisOp.getApprox()) {
4022 id = isFtz ? llvm::Intrinsic::nvvm_div_approx_ftz_f
4023 : llvm::Intrinsic::nvvm_div_approx_f;
4024 }
else if (thisOp.getFull()) {
4027 id = isFtz ? llvm::Intrinsic::nvvm_div_full_ftz
4028 : llvm::Intrinsic::nvvm_div_full;
4031 unsigned rndIndex =
static_cast<unsigned>(thisOp.getRnd()) - 1;
4033 static constexpr llvm::Intrinsic::ID f32IDs[] = {
4034 llvm::Intrinsic::nvvm_div_rn_f,
4035 llvm::Intrinsic::nvvm_div_rm_f,
4036 llvm::Intrinsic::nvvm_div_rp_f,
4037 llvm::Intrinsic::nvvm_div_rz_f,
4039 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
4040 llvm::Intrinsic::nvvm_div_rn_ftz_f,
4041 llvm::Intrinsic::nvvm_div_rm_ftz_f,
4042 llvm::Intrinsic::nvvm_div_rp_ftz_f,
4043 llvm::Intrinsic::nvvm_div_rz_ftz_f,
4045 static constexpr llvm::Intrinsic::ID f64IDs[] = {
4046 llvm::Intrinsic::nvvm_div_rn_d,
4047 llvm::Intrinsic::nvvm_div_rm_d,
4048 llvm::Intrinsic::nvvm_div_rp_d,
4049 llvm::Intrinsic::nvvm_div_rz_d,
4051 Type t = thisOp.getRes().getType();
4052 id = t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4063 auto thisOp = cast<NVVM::AsyncStoreGlobalOp>(op);
4064 mlir::NVVM::MemScopeKind scope = thisOp.getScope();
4065 bool isMmio = thisOp.getMmio();
4067 llvm::Value *addr = mt.
lookupValue(thisOp.getAddr());
4068 llvm::Value *value = mt.
lookupValue(thisOp.getValue());
4069 llvm::Value *isMultimem = builder.getInt1(thisOp.getMultimem());
4071 if (scope == MemScopeKind::SYS) {
4072 return isMmio ?
IDArgPair(llvm::Intrinsic::nvvm_st_async_mmio_sys,
4075 {addr, value, isMultimem});
4076 }
else if (scope == MemScopeKind::GPU) {
4077 return IDArgPair(llvm::Intrinsic::nvvm_st_async_gpu,
4078 {addr, value, isMultimem});
4080 llvm_unreachable(
"unsupported scope for AsyncStoreGlobalOp");
4085 llvm::IRBuilderBase &builder) {
4086 auto thisOp = cast<NVVM::PMEventOp>(op);
4090 llvm::Value *maskVal;
4091 if (
auto eventAttr = thisOp.getEventIdAttr()) {
4092 uint16_t mask =
static_cast<uint16_t
>(1u << eventAttr.getInt());
4093 maskVal = llvm::ConstantInt::get(i16Ty, mask);
4096 llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
4099 return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
4102bool MBarrierInitOp::getAsmValues(
4109 for (
auto val : getOperands())
4117 auto thisOp = cast<NVVM::MBarrierInitOp>(op);
4121 return {llvm::Intrinsic::nvvm_mbarrier_init,
4123 builder.getInt32(thisOp.getLayout())}};
4128 auto thisOp = cast<NVVM::MBarrierInvalOp>(op);
4130 llvm::Intrinsic::ID
id = isShared
4131 ? llvm::Intrinsic::nvvm_mbarrier_inval_shared
4132 : llvm::Intrinsic::nvvm_mbarrier_inval;
4139 auto thisOp = cast<NVVM::MBarrierCheckLayoutOp>(op);
4142 llvm::Intrinsic::nvvm_mbarrier_check_layout,
4143 {mt.
lookupValue(thisOp.getAddr()), builder.getInt32(thisOp.getLayout())}};
4148 auto thisOp = cast<NVVM::MBarrierExpectTxOp>(op);
4151 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4154 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4156 static constexpr llvm::Intrinsic::ID IDs[] = {
4157 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cta,
4158 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cluster,
4159 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cta,
4160 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cluster};
4165 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4167 return {IDs[
index], std::move(args)};
4172 auto thisOp = cast<NVVM::MBarrierCompleteTxOp>(op);
4175 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4178 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4180 static constexpr llvm::Intrinsic::ID IDs[] = {
4181 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cta,
4182 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cluster,
4183 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cta,
4184 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cluster};
4189 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4191 return {IDs[
index], std::move(args)};
4196 auto thisOp = cast<NVVM::MBarrierArriveOp>(op);
4199 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4202 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4204 static constexpr llvm::Intrinsic::ID IDs[] = {
4205 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta,
4206 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cluster,
4207 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cta,
4208 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cluster};
4209 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4210 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cta,
4211 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cluster,
4212 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cta,
4214 nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cluster};
4215 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4219 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4226 bool hasCount =
static_cast<bool>(thisOp.getCount());
4228 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta))
4229 return {llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {mbar}};
4233 llvm::Value *count =
4235 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4236 return {id, {mbar, count}};
4241 auto thisOp = cast<NVVM::MBarrierArriveDropOp>(op);
4244 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4247 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4249 static constexpr llvm::Intrinsic::ID IDs[] = {
4250 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta,
4251 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cluster,
4252 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cta,
4253 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cluster};
4254 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4255 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cta,
4257 nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cluster,
4259 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cta,
4261 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cluster};
4262 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4266 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4273 bool hasCount =
static_cast<bool>(thisOp.getCount());
4275 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta))
4276 return {llvm::Intrinsic::nvvm_mbarrier_arrive_drop_shared, {mbar}};
4280 llvm::Value *count =
4282 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4283 return {id, {mbar, count}};
4286bool MBarrierArriveExpectTxOp::getAsmValues(
4293 for (
auto val : getOperands())
4301 auto thisOp = cast<NVVM::MBarrierArriveExpectTxOp>(op);
4304 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4307 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4310 static constexpr llvm::Intrinsic::ID IDs[] = {
4311 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cta,
4312 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cluster,
4313 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cta,
4314 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cluster};
4315 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4316 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cta,
4317 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cluster,
4318 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cta,
4319 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cluster};
4321 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4324 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4325 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4330 return {id, {mbar, txcount}};
4335 auto thisOp = cast<NVVM::MBarrierArriveDropExpectTxOp>(op);
4338 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4341 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4344 static constexpr llvm::Intrinsic::ID IDs[] = {
4345 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cta,
4346 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cluster,
4347 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cta,
4348 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cluster};
4349 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4350 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cta,
4351 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cluster,
4352 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cta,
4353 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cluster};
4355 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4358 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4359 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4364 return {id, {mbar, txcount}};
4369 auto thisOp = cast<NVVM::MBarrierArriveNocompleteOp>(op);
4371 llvm::Intrinsic::ID
id =
4372 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared
4373 : llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete;
4377 args.push_back(mt.
lookupValue(thisOp.getCount()));
4379 return {id, std::move(args)};
4384 auto thisOp = cast<NVVM::MBarrierArriveDropNocompleteOp>(op);
4386 llvm::Intrinsic::ID
id =
4387 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete_shared
4388 : llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete;
4392 args.push_back(mt.
lookupValue(thisOp.getCount()));
4394 return {id, std::move(args)};
4399 auto thisOp = cast<NVVM::MBarrierTestWaitOp>(op);
4400 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4401 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4404 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isPhaseParity ? 1 : 0);
4407 static constexpr llvm::Intrinsic::ID IDs[] = {
4408 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cta_space_cta,
4409 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cta_space_cta,
4410 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cluster_space_cta,
4411 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cluster_space_cta};
4412 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4413 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cta_space_cta,
4414 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cta_space_cta,
4415 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cluster_space_cta,
4416 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cluster_space_cta};
4418 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4421 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4422 llvm::Value *input = mt.
lookupValue(thisOp.getStateOrPhase());
4427 return {id, {mbar, input}};
4432 auto thisOp = cast<NVVM::MBarrierTryWaitOp>(op);
4433 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4434 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4435 bool hasTicks =
static_cast<bool>(thisOp.getTicks());
4439 size_t index = ((hasTicks ? 1 : 0) << 2) | ((isClusterScope ? 1 : 0) << 1) |
4440 (isPhaseParity ? 1 : 0);
4443 static constexpr llvm::Intrinsic::ID IDs[] = {
4444 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cta_space_cta,
4445 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cta_space_cta,
4446 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cluster_space_cta,
4447 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cluster_space_cta,
4448 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cta_space_cta,
4449 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cta_space_cta,
4450 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cluster_space_cta,
4451 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cluster_space_cta};
4452 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4453 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cta_space_cta,
4454 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cta_space_cta,
4455 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cluster_space_cta,
4456 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cluster_space_cta,
4457 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cta_space_cta,
4458 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cta_space_cta,
4459 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cluster_space_cta,
4460 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cluster_space_cta};
4462 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4465 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4472 args.push_back(mbar);
4473 args.push_back(mt.
lookupValue(thisOp.getStateOrPhase()));
4475 args.push_back(mt.
lookupValue(thisOp.getTicks()));
4477 return {id, std::move(args)};
4482 auto thisOp = cast<NVVM::CpAsyncMBarrierArriveOp>(op);
4485 llvm::Intrinsic::ID id;
4486 if (thisOp.getNoinc()) {
4487 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc_shared
4488 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc;
4490 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_shared
4491 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive;
4499 llvm::IRBuilderBase &builder) {
4500 auto thisOp = cast<NVVM::MovMatrixOp>(op);
4501 return {llvm::Intrinsic::nvvm_movmatrix_sync_aligned_m8n8_trans_b16,
4505#define CP_ASYNC_ID_IMPL(mod, size, suffix) \
4506 llvm::Intrinsic::nvvm_cp_async_##mod##_shared_global_##size##suffix
4508#define GET_CP_ASYNC_ID(mod, size, has_cpsize) \
4509 has_cpsize ? CP_ASYNC_ID_IMPL(mod, size, _s) : CP_ASYNC_ID_IMPL(mod, size, )
4514 llvm::Intrinsic::ID id;
4516 auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
4517 bool hasCpSize =
static_cast<bool>(cpAsyncOp.getCpSize());
4518 switch (cpAsyncOp.getSize()) {
4526 id = (cpAsyncOp.getModifier() == NVVM::LoadCacheModifierKind::CG)
4531 llvm_unreachable(
"Invalid copy size in CpAsyncOp.");
4535 args.push_back(mt.
lookupValue(cpAsyncOp.getDst()));
4536 args.push_back(mt.
lookupValue(cpAsyncOp.getSrc()));
4538 args.push_back(mt.
lookupValue(cpAsyncOp.getCpSize()));
4545 auto thisOp = cast<NVVM::CpAsyncBulkPrefetchOp>(op);
4547 llvm::Intrinsic::ID
id = llvm::Intrinsic::nvvm_cp_async_bulk_prefetch_L2;
4550 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4554 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4555 llvm::Value *i64Unused =
4556 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4557 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4558 args.push_back(builder.getInt1(hasCacheHint));
4560 return {id, std::move(args)};
4565 auto thisOp = cast<NVVM::CpAsyncBulkGlobalToSharedClusterOp>(op);
4569 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4571 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4575 mlir::Value multicastMask = thisOp.getMulticastMask();
4576 const bool hasMulticastMask =
static_cast<bool>(multicastMask);
4579 llvm::Value *i16Unused = llvm::ConstantInt::get(builder.getInt16Ty(), 0);
4580 args.push_back(hasMulticastMask ? mt.
lookupValue(multicastMask)
4583 args.push_back(builder.getInt32(0));
4584 args.push_back(builder.getInt32(0));
4589 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4590 llvm::Value *i64Unused = llvm::ConstantInt::get(builder.getInt64Ty(), 0);
4591 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4595 args.push_back(builder.getInt1(hasCacheHint));
4596 args.push_back(builder.getInt1(
false));
4598 args.push_back(builder.getInt1(hasMulticastMask));
4599 args.push_back(builder.getInt1(hasCacheHint));
4603 args.push_back(builder.getInt32(0));
4605 llvm::Intrinsic::ID
id =
4607 ? llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta
4608 : llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
4610 return {id, std::move(args)};
4615 auto thisOp = cast<NVVM::CpAsyncBulkSharedCTAToGlobalOp>(op);
4617 llvm::Intrinsic::ID
id =
4618 llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global;
4621 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4622 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4626 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4627 llvm::Value *i64Unused =
4628 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4629 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4630 args.push_back(builder.getInt1(hasCacheHint));
4633 if (
mlir::Value byteMask = thisOp.getByteMask()) {
4635 id = llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global_bytemask;
4638 return {id, std::move(args)};
4641bool CpAsyncBulkTensorGlobalToSharedClusterOp::getAsmValues(
4648 for (
auto val : getOperands())
4655CpAsyncBulkTensorGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
4657 auto thisOp = cast<NVVM::CpAsyncBulkTensorGlobalToSharedClusterOp>(op);
4658 const bool isCTAOnly = thisOp.getIsCTAOnly();
4662 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4664 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4674 const bool hasMC =
static_cast<bool>(mcMask);
4675 llvm::Value *i16Zero =
4676 llvm::ConstantInt::get(llvm::Type::getInt16Ty(mt.
getLLVMContext()), 0);
4680 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4681 llvm::Value *i64Zero =
4682 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4688 thisOp.getGroup() ? (
static_cast<int32_t
>(*thisOp.getGroup()) + 1) : 0;
4690 llvm::ConstantInt::get(llvm::Type::getInt32Ty(mt.
getLLVMContext()), val);
4693 llvm::Value *flagValidPattern = builder.getInt32(0);
4697 args.push_back(hasMC ? mt.
lookupValue(mcMask) : i16Zero);
4698 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4699 args.push_back(builder.getInt1(hasMC));
4700 args.push_back(builder.getInt1(hasCacheHint));
4702 args.push_back(flagValidPattern);
4705 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4706 args.push_back(builder.getInt1(hasCacheHint));
4707 args.push_back(flagValidPattern);
4710 constexpr size_t numDims = 5;
4711 constexpr size_t numModes = 5;
4712 using rowTy = std::array<llvm::Intrinsic::ID, numDims + 1>;
4713 using TableTy = std::array<rowTy, numModes>;
4714 static constexpr TableTy IDTable{
4715 {{
notIntrinsic, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d,
4716 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d,
4717 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d,
4718 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d,
4719 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d},
4721 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d,
4722 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d,
4723 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d},
4725 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_3d,
4726 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_4d,
4727 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_5d},
4729 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_3d,
4730 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_4d,
4731 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_5d},
4733 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_gather4_2d}}};
4735 static constexpr TableTy IDTableCTA{
4737 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_1d,
4738 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_2d,
4739 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_3d,
4740 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_4d,
4741 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_5d},
4743 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_3d,
4744 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_4d,
4745 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_5d},
4747 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_3d,
4748 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_4d,
4749 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_5d},
4751 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_3d,
4752 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_4d,
4753 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_5d},
4755 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_gather4_2d}}};
4758 (getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1) &&
4759 (getMaxEnumValForTMALoadMode() == std::size(IDTableCTA) - 1),
4760 "TMALoadModes must match number of rows in IDTable and IDTableCTA");
4761 size_t mode =
static_cast<size_t>(thisOp.getMode());
4762 size_t dim = thisOp.getCoordinates().size();
4763 auto id = isCTAOnly ? IDTableCTA[mode][dim] : IDTable[mode][dim];
4765 "Invalid intrinsic for CpAsyncBulkTensorGlobalToSharedClusterOp.");
4767 return {id, std::move(args)};
4772 auto thisOp = cast<NVVM::CpAsyncBulkTensorPrefetchOp>(op);
4776 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4778 for (
auto v : thisOp.getCoordinates())
4780 for (
auto v : thisOp.getIm2colOffsets())
4784 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4785 llvm::Value *i64Unused =
4786 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4787 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4788 args.push_back(builder.getInt1(hasCacheHint));
4790 const unsigned NI = llvm::Intrinsic::not_intrinsic;
4791 static constexpr llvm::Intrinsic::ID IDTable[][6] = {
4792 {NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_1d,
4793 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_2d,
4794 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_3d,
4795 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_4d,
4796 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_5d},
4798 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_3d,
4799 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_4d,
4800 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_5d},
4802 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_3d,
4803 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_4d,
4804 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_5d},
4806 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_3d,
4807 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_4d,
4808 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_5d},
4809 {NI, NI, NI, NI, NI,
4810 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_gather4_2d}};
4812 static_assert(getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1,
4813 "TMALoadModes must match number of rows in IDTable");
4814 size_t mode =
static_cast<size_t>(thisOp.getMode());
4815 size_t dim = thisOp.getCoordinates().size();
4816 llvm::Intrinsic::ID
id = IDTable[mode][dim];
4817 if (
id == llvm::Intrinsic::not_intrinsic)
4818 llvm_unreachable(
"Invalid intrinsic for CpAsyncBulkTensorPrefetchOp.");
4820 return {id, std::move(args)};
4824CpAsyncBulkTensorSharedCTAToGlobalOp::getIntrinsicIDAndArgs(
4826 auto thisOp = cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp>(op);
4830 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4831 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4833 for (
auto v : thisOp.getCoordinates())
4837 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4838 llvm::Value *i64Unused =
4839 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4840 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4841 args.push_back(builder.getInt1(hasCacheHint));
4843 using namespace llvm::Intrinsic;
4844 const unsigned NI = not_intrinsic;
4845 static constexpr ID IDTable[][6] = {
4846 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_1d,
4847 nvvm_cp_async_bulk_tensor_s2g_tile_2d,
4848 nvvm_cp_async_bulk_tensor_s2g_tile_3d,
4849 nvvm_cp_async_bulk_tensor_s2g_tile_4d,
4850 nvvm_cp_async_bulk_tensor_s2g_tile_5d},
4851 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_3d,
4852 nvvm_cp_async_bulk_tensor_s2g_im2col_4d,
4853 nvvm_cp_async_bulk_tensor_s2g_im2col_5d},
4854 {NI, NI, NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_2d},
4855 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_3d,
4856 nvvm_cp_async_bulk_tensor_s2g_im2col_w_4d,
4857 nvvm_cp_async_bulk_tensor_s2g_im2col_w_5d}};
4859 static_assert(getMaxEnumValForTMAStoreMode() == std::size(IDTable) - 1,
4860 "TMAStoreModes must match number of rows in IDTable");
4861 size_t mode =
static_cast<size_t>(thisOp.getMode());
4862 size_t dim = thisOp.getCoordinates().size();
4863 ID
id = IDTable[mode][dim];
4864 if (
id == llvm::Intrinsic::not_intrinsic)
4866 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOp.");
4868 return {id, std::move(args)};
4872CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
4875 cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp>(op);
4878 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4879 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4880 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
4881 for (
Value v : thisOp.getTensorSize())
4883 for (
Value v : thisOp.getLowerStride())
4885 if (thisOp.getUpperStride())
4886 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
4887 for (
Value v : thisOp.getCoordinates())
4891 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4892 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4893 : builder.getInt64(0));
4894 args.push_back(builder.getInt1(hasCacheHint));
4896 using namespace llvm::Intrinsic;
4897 const unsigned NI = not_intrinsic;
4900 static constexpr ID IDTable[][6] = {
4901 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_1d,
4902 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_2d,
4903 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_3d,
4904 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_4d,
4905 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_5d},
4906 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_3d,
4907 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_4d,
4908 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_5d},
4909 {NI, NI, NI, NI, NI,
4910 nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_override_addr_2d},
4911 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_3d,
4912 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_4d,
4913 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_5d}};
4917 static constexpr ID dimStrideIDTable[] = {
4918 NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_1d,
4919 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_2d,
4920 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_3d,
4921 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_4d,
4922 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_5d};
4925 size_t mode =
static_cast<size_t>(thisOp.getMode());
4926 size_t dim = thisOp.getCoordinates().size();
4927 bool isDimStride = !thisOp.getTensorSize().empty();
4929 assert(mode < std::size(IDTable) &&
4930 "Invalid mode for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4931 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
4932 "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4934 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
4936 intrinsicID != NI &&
4937 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4938 return {intrinsicID, std::move(args)};
4943 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOp>(op);
4946 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4947 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4948 for (
Value v : thisOp.getCoordinates())
4952 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4953 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4954 : builder.getInt64(0));
4955 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
4956 args.push_back(builder.getInt1(hasCacheHint));
4958 using namespace llvm::Intrinsic;
4959 const unsigned NI = not_intrinsic;
4960 static constexpr ID IDTable[][6] = {
4961 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_1d,
4962 nvvm_cp_async_bulk_tensor_reduce_tile_2d,
4963 nvvm_cp_async_bulk_tensor_reduce_tile_3d,
4964 nvvm_cp_async_bulk_tensor_reduce_tile_4d,
4965 nvvm_cp_async_bulk_tensor_reduce_tile_5d},
4966 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_3d,
4967 nvvm_cp_async_bulk_tensor_reduce_im2col_4d,
4968 nvvm_cp_async_bulk_tensor_reduce_im2col_5d},
4969 {NI, NI, NI, NI, NI, NI},
4970 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_3d,
4971 nvvm_cp_async_bulk_tensor_reduce_im2col_w_4d,
4972 nvvm_cp_async_bulk_tensor_reduce_im2col_w_5d}};
4974 size_t mode =
static_cast<size_t>(thisOp.getMode());
4975 size_t dim = thisOp.getCoordinates().size();
4976 assert(mode < std::size(IDTable) &&
4977 "Invalid mode for CpAsyncBulkTensorReduceOp");
4978 assert(dim < std::size(IDTable[mode]) &&
4979 "Invalid dim for CpAsyncBulkTensorReduceOp");
4981 ID intrinsicID = IDTable[mode][dim];
4982 assert(intrinsicID != NI &&
4983 "Invalid intrinsic for CpAsyncBulkTensorReduceOp");
4984 return {intrinsicID, std::move(args)};
4987NVVM::IDArgPair CpAsyncBulkTensorReduceOverrideAddrOp::getIntrinsicIDAndArgs(
4989 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOverrideAddrOp>(op);
4992 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4993 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4994 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
4996 for (
Value v : thisOp.getTensorSize())
4998 for (
Value v : thisOp.getLowerStride())
5000 if (thisOp.getUpperStride())
5001 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
5002 for (
Value v : thisOp.getCoordinates())
5006 const bool hasCacheHint =
static_cast<bool>(cacheHint);
5007 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
5008 : builder.getInt64(0));
5009 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
5010 args.push_back(builder.getInt1(hasCacheHint));
5012 using namespace llvm::Intrinsic;
5013 const unsigned NI = not_intrinsic;
5016static constexpr ID IDTable[][6] = {
5017 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_1d,
5018 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_2d,
5019 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_3d,
5020 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_4d,
5021 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_5d},
5022 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_3d,
5023 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_4d,
5024 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_5d},
5025 {NI, NI, NI, NI, NI, NI},
5026 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_3d,
5027 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_4d,
5028 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_5d}};
5032static constexpr ID dimStrideIDTable[] = {
5033 NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_1d,
5034 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_2d,
5035 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_3d,
5036 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_4d,
5037 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_5d};
5040 size_t mode =
static_cast<size_t>(thisOp.getMode());
5041 size_t dim = thisOp.getCoordinates().size();
5042 bool isDimStride = !thisOp.getTensorSize().empty();
5044 assert(mode < std::size(IDTable) &&
5045 "Invalid mode for CpAsyncBulkTensorReduceOverrideAddrOp");
5046 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
5047 "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
5049 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
5050 assert(intrinsicID != NI &&
5051 "Invalid intrinsic for CpAsyncBulkTensorReduceOverrideAddrOp");
5052 return {intrinsicID, std::move(args)};
5057#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5058 hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
5059 : llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
5061#define GET_CVT_F2TF32_ID(rnd, relu, sf) \
5062 hasSatFinite ? CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5063 : CVT_F2TF32_ID_IMPL(rnd, relu, )
5066ConvertFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
5067 NVVM::SaturationMode sat,
bool hasRelu) {
5068 using RndMode = NVVM::FPRoundingMode;
5069 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5078 llvm_unreachable(
"Invalid RoundingMode for CvtFloatToTF32Op");
5083ConvertF32x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF4x2Op op,
5085 llvm::IRBuilderBase &builder) {
5090 bool hasRelu = op.getRelu();
5092 llvm::Intrinsic::ID intId =
5093 hasRelu ? llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_relu_satfinite
5094 : llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_satfinite;
5096 return {intId, std::move(args)};
5099#define GET_F32x2_TO_F6x2_ID(type, has_relu) \
5100 has_relu ? llvm::Intrinsic::nvvm_ff_to_##type##_rn_relu_satfinite \
5101 : llvm::Intrinsic::nvvm_ff_to_##type##_rn_satfinite
5103llvm::Intrinsic::ID ConvertF32x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
5106 .Case([&](mlir::Float6E2M3FNType) {
5109 .Case([&](mlir::Float6E3M2FNType) {
5113 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF6x2Op");
5114 return llvm::Intrinsic::not_intrinsic;
5119ConvertF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF16x2ToF4x2Op &op,
5121 llvm::IRBuilderBase &builder) {
5123 bool hasRelu = op.getRelu();
5125 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5127 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5128 intId = hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_relu_satfinite
5129 : llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_satfinite;
5134 return {intId, std::move(args)};
5138ConvertBF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertBF16x2ToF4x2Op &op,
5140 llvm::IRBuilderBase &builder) {
5142 bool hasRelu = op.getRelu();
5144 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5146 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5147 intId = hasRelu ? llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite
5148 : llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_satfinite;
5153 return {intId, std::move(args)};
5156llvm::Intrinsic::ID ConvertF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
5159 .Case<mlir::Float6E2M3FNType>([&](mlir::Float6E2M3FNType) {
5160 return hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_relu_satfinite
5161 : llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_satfinite;
5163 .Case<mlir::Float6E3M2FNType>([&](mlir::Float6E3M2FNType) {
5164 return hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_relu_satfinite
5165 : llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_satfinite;
5168 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF6x2Op");
5169 return llvm::Intrinsic::not_intrinsic;
5173llvm::Intrinsic::ID ConvertBF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
5176 .Case<mlir::Float6E2M3FNType>([&](mlir::Float6E2M3FNType) {
5178 ? llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite
5179 : llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_satfinite;
5181 .Case<mlir::Float6E3M2FNType>([&](mlir::Float6E3M2FNType) {
5183 ? llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite
5184 : llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_satfinite;
5187 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF6x2Op");
5188 return llvm::Intrinsic::not_intrinsic;
5192#define GET_F32x2_TO_F8X2_US_ID(rnd, has_satf) \
5193 has_satf ? llvm::Intrinsic::nvvm_ff_to_ue8m0x2_##rnd##_satfinite \
5194 : llvm::Intrinsic::nvvm_ff_to_ue8m0x2_##rnd
5196#define GET_F32x2_TO_F8X2_S_ID(type, has_relu) \
5197 has_relu ? llvm::Intrinsic::nvvm_ff_to_##type##_rn_relu \
5198 : llvm::Intrinsic::nvvm_ff_to_##type##_rn
5201ConvertF32x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5202 NVVM::SaturationMode sat,
bool hasRelu) {
5203 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5204 bool hasRoundingModeRZ = (rnd == NVVM::FPRoundingMode::RZ);
5205 bool hasRoundingModeRP = (rnd == NVVM::FPRoundingMode::RP);
5208 .Case([&](mlir::Float8E4M3FNType) {
5211 .Case([&](mlir::Float8E5M2Type) {
5214 .Case([&](mlir::Float8E8M0FNUType) {
5215 if (hasRoundingModeRZ)
5217 else if (hasRoundingModeRP)
5220 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5223 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5224 return llvm::Intrinsic::not_intrinsic;
5228#define GET_F16x2_TO_F8X2_ID(type, has_relu) \
5229 has_relu ? llvm::Intrinsic::nvvm_f16x2_to_##type##_rn_relu \
5230 : llvm::Intrinsic::nvvm_f16x2_to_##type##_rn
5232llvm::Intrinsic::ID ConvertF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
5235 .Case([&](mlir::Float8E4M3FNType) {
5238 .Case([&](mlir::Float8E5M2Type) {
5242 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF8x2Op");
5243 return llvm::Intrinsic::not_intrinsic;
5248ConvertBF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
5249 NVVM::FPRoundingMode rnd,
5250 NVVM::SaturationMode sat,
bool hasRelu) {
5251 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5253 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
5254 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz,
5255 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp,
5256 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz_satfinite,
5257 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp_satfinite,
5261 .Case<mlir::Float8E4M3FNType>([&](mlir::Float8E4M3FNType) {
5263 ? llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite
5264 : llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_satfinite;
5266 .Case<mlir::Float8E5M2Type>([&](mlir::Float8E5M2Type) {
5268 ? llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite
5269 : llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_satfinite;
5271 .Case<mlir::Float8E8M0FNUType>([&](mlir::Float8E8M0FNUType) {
5272 bool hasRoundingModeRP = (rnd == NVVM::FPRoundingMode::RP);
5273 unsigned index = (hasSatFinite << 1) | hasRoundingModeRP;
5274 return ue8m0x2IDs[
index];
5277 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
5278 return llvm::Intrinsic::not_intrinsic;
5284 auto curOp = cast<NVVM::ConvertF8x2ToF16x2Op>(op);
5286 bool hasRelu = curOp.getRelu();
5288 llvm::Intrinsic::ID intId =
5290 .Case([&](Float8E4M3FNType type) {
5291 return hasRelu ? llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn_relu
5292 : llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn;
5294 .Case([&](Float8E5M2Type type) {
5295 return hasRelu ? llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn_relu
5296 : llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn;
5299 llvm_unreachable(
"Invalid type for ConvertF8x2ToF16x2Op");
5300 return llvm::Intrinsic::not_intrinsic;
5303 llvm::Value *packedI16 =
5304 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5305 llvm::Type::getInt16Ty(builder.getContext()));
5307 return {intId, {packedI16}};
5312 auto curOp = cast<NVVM::ConvertF8x2ToBF16x2Op>(op);
5313 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5314 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5315 bool hasRelu = curOp.getRelu();
5317 static constexpr llvm::Intrinsic::ID E4M3Ids[] = {
5318 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5319 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5320 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5321 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5324 static constexpr llvm::Intrinsic::ID E5M2Ids[] = {
5325 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5326 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5327 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5328 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5331 llvm::Intrinsic::ID intId =
5333 .Case([&](Float8E8M0FNUType type) {
5334 return llvm::Intrinsic::nvvm_ue8m0x2_to_bf16x2;
5336 .Case([&](Float8E4M3FNType type) {
5337 return E4M3Ids[hasSatfinite << 1 | hasRelu];
5339 .Case([&](Float8E5M2Type type) {
5340 return E5M2Ids[hasSatfinite << 1 | hasRelu];
5343 llvm_unreachable(
"Invalid type for ConvertF8x2ToBF16x2Op");
5344 return llvm::Intrinsic::not_intrinsic;
5346 llvm::Value *packedI16 =
5347 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5348 llvm::Type::getInt16Ty(builder.getContext()));
5351 args.push_back(packedI16);
5352 if (!isa<Float8E8M0FNUType>(curOp.getSrcType()))
5355 : builder.getInt16(0x7f7f));
5358 return {intId, std::move(args)};
5363 auto curOp = cast<NVVM::ConvertF6x2ToF16x2Op>(op);
5365 bool hasRelu = curOp.getRelu();
5367 llvm::Intrinsic::ID intId =
5369 .Case([&](Float6E2M3FNType type) {
5370 return hasRelu ? llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn_relu
5371 : llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn;
5373 .Case([&](Float6E3M2FNType type) {
5374 return hasRelu ? llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn_relu
5375 : llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn;
5378 llvm_unreachable(
"Invalid type for ConvertF6x2ToF16x2Op");
5379 return llvm::Intrinsic::not_intrinsic;
5382 llvm::Value *packedI16 =
5383 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5384 llvm::Type::getInt16Ty(builder.getContext()));
5386 return {intId, {packedI16}};
5391 auto curOp = cast<NVVM::ConvertF6x2ToBF16x2Op>(op);
5392 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5393 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5394 bool hasRelu = curOp.getRelu();
5396 static constexpr llvm::Intrinsic::ID E2M3Ids[] = {
5397 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5398 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5399 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5400 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5403 static constexpr llvm::Intrinsic::ID E3M2Ids[] = {
5404 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5405 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5406 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5407 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5410 unsigned idx = (hasSatfinite << 1) | hasRelu;
5411 llvm::Intrinsic::ID intId =
5413 .Case([&](Float6E2M3FNType type) {
return E2M3Ids[idx]; })
5414 .Case([&](Float6E3M2FNType type) {
return E3M2Ids[idx]; })
5416 llvm_unreachable(
"Invalid type for ConvertF6x2ToBF16x2Op");
5417 return llvm::Intrinsic::not_intrinsic;
5420 llvm::Value *packedI16 =
5421 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5422 llvm::Type::getInt16Ty(builder.getContext()));
5425 args.push_back(packedI16);
5432 return {intId, std::move(args)};
5437 auto curOp = cast<NVVM::ConvertF4x2ToF16x2Op>(op);
5439 bool hasRelu = curOp.getRelu();
5441 llvm::Intrinsic::ID intId =
5443 .Case([&](Float4E2M1FNType type) {
5444 return hasRelu ? llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn_relu
5445 : llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn;
5448 llvm_unreachable(
"Invalid type for ConvertF4x2ToF16x2Op");
5449 return llvm::Intrinsic::not_intrinsic;
5452 llvm::Value *extendedI16 =
5453 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5454 llvm::Type::getInt16Ty(builder.getContext()));
5456 return {intId, {extendedI16}};
5461 auto curOp = cast<NVVM::ConvertF4x2ToBF16x2Op>(op);
5462 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5463 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5464 bool hasRelu = curOp.getRelu();
5466 static constexpr llvm::Intrinsic::ID E2M1Ids[] = {
5467 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_scale_n2_ue8m0,
5468 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5469 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5470 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5473 unsigned idx = (hasSatfinite << 1) | hasRelu;
5474 llvm::Intrinsic::ID intId =
5476 .Case([&](Float4E2M1FNType type) {
return E2M1Ids[idx]; })
5478 llvm_unreachable(
"Invalid type for ConvertF4x2ToBF16x2Op");
5479 return llvm::Intrinsic::not_intrinsic;
5482 llvm::Value *extendedI16 =
5483 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5484 llvm::Type::getInt16Ty(builder.getContext()));
5487 args.push_back(extendedI16);
5494 return {intId, std::move(args)};
5499 auto thisOp = cast<NVVM::ConvertF32x2ToS2F6x2Op>(op);
5500 bool hasRelu = thisOp.getRelu();
5501 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5503 llvm::Intrinsic::ID
id =
5505 ? llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5506 : llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5512 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5513 : builder.getInt16(0x7f7f));
5514 return {id, std::move(args)};
5519 auto thisOp = cast<NVVM::ConvertBF16x2ToS2F6x2Op>(op);
5520 bool hasRelu = thisOp.getRelu();
5521 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5523 llvm::Intrinsic::ID
id =
5526 nvvm_bf16x2_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5527 : llvm::Intrinsic::nvvm_bf16x2_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5532 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5533 : builder.getInt16(0x7f7f));
5534 return {id, std::move(args)};
5539 auto thisOp = cast<NVVM::ConvertS2F6x2ToBF16x2Op>(op);
5540 bool hasRelu = thisOp.getRelu();
5541 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5542 bool hasSat = thisOp.getSat() == NVVM::SaturationMode::SATFINITE;
5544 static constexpr llvm::Intrinsic::ID ids[] = {
5545 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_scale_n2_ue8m0,
5546 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5547 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5548 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5551 unsigned idx = (hasSat << 1) | hasRelu;
5555 llvm::Value *packedI16 =
5556 builder.CreateBitCast(mt.
lookupValue(thisOp.getSrc()),
5557 llvm::Type::getInt16Ty(builder.getContext()));
5558 args.push_back(packedI16);
5559 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5560 : builder.getInt16(0x7f7f));
5562 return {ids[idx], std::move(args)};
5567 auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
5568 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5570 llvm::Intrinsic::ID
id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
5571 : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
5575 builder.getInt1(curOp.getIsExclusive())}};
5580 auto curOp = cast<NVVM::Tcgen05DeallocOp>(op);
5581 llvm::Intrinsic::ID
id = (curOp.getGroup() == CTAGroupKind::CTA_1)
5582 ? llvm::Intrinsic::nvvm_tcgen05_dealloc_cg1
5583 : llvm::Intrinsic::nvvm_tcgen05_dealloc_cg2;
5587 builder.getInt1(curOp.getIsExclusive())}};
5591Tcgen05CommitOp::getIntrinsicIDAndArgs(
Operation &op,
5594 auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
5595 bool hasMulticast =
static_cast<bool>(curOp.getMulticastMask());
5596 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5597 bool hasSmemARead = curOp.getSmemARead();
5598 unsigned index = (
static_cast<unsigned>(hasSmemARead) << 1) |
5599 static_cast<unsigned>(is2CTAMode);
5601 using namespace llvm::Intrinsic;
5602 static constexpr ID IDs[] = {
5603 nvvm_tcgen05_commit_cg1,
5604 nvvm_tcgen05_commit_cg2,
5605 nvvm_tcgen05_commit_smem_a_read_cg1,
5606 nvvm_tcgen05_commit_smem_a_read_cg2,
5609 static constexpr ID multicastIDs[] = {
5610 nvvm_tcgen05_commit_mc_cg1,
5611 nvvm_tcgen05_commit_mc_cg2,
5612 nvvm_tcgen05_commit_smem_a_read_mc_cg1,
5613 nvvm_tcgen05_commit_smem_a_read_mc_cg2,
5616 ID
id = hasMulticast ? multicastIDs[
index] : IDs[
index];
5620 args.push_back(mt.
lookupValue(curOp.getMulticastMask()));
5625#define TCGEN05_CP_IMPL(shape_mc, src_fmt, cg) \
5626 llvm::Intrinsic::nvvm_tcgen05_cp##shape_mc##src_fmt##cg
5628#define TCGEN05_CP_2CTA(shape_mc, src_fmt, is_2cta) \
5629 is_2cta ? TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg2) \
5630 : TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
5632#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
5634 if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
5635 return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
5636 if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
5637 return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
5638 return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
5642ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
5644 llvm::IRBuilderBase &builder) {
5645 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5646 llvm::Intrinsic::nvvm_ff2f16x2_rn,
5647 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu,
5648 llvm::Intrinsic::nvvm_ff2f16x2_rn_satfinite,
5649 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu_satfinite,
5651 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5652 llvm::Intrinsic::nvvm_ff2f16x2_rz,
5653 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu,
5654 llvm::Intrinsic::nvvm_ff2f16x2_rz_satfinite,
5655 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu_satfinite,
5657 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5658 llvm::Intrinsic::nvvm_ff2f16x2_rs,
5659 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu,
5660 llvm::Intrinsic::nvvm_ff2f16x2_rs_satfinite,
5661 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu_satfinite,
5664 unsigned hasRelu = op.getRelu() ? 1 : 0;
5665 unsigned hasSatFinite =
5666 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5669 unsigned idx = (hasSatFinite << 1) | hasRelu;
5674 if (op.getRandomBits())
5675 args.push_back(mt.
lookupValue(op.getRandomBits()));
5678 args.push_back(builder.getInt1(
false));
5680 switch (op.getRnd()) {
5681 case FPRoundingMode::RN:
5682 return {rndRNIds[idx], std::move(args)};
5683 case FPRoundingMode::RZ:
5684 return {rndRZIds[idx], std::move(args)};
5685 case FPRoundingMode::RS:
5686 return {rndRSIds[idx], std::move(args)};
5688 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToF16x2Op");
5693ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
5695 llvm::IRBuilderBase &builder) {
5696 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5697 llvm::Intrinsic::nvvm_ff2bf16x2_rn,
5698 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu,
5699 llvm::Intrinsic::nvvm_ff2bf16x2_rn_satfinite,
5700 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu_satfinite,
5702 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5703 llvm::Intrinsic::nvvm_ff2bf16x2_rz,
5704 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu,
5705 llvm::Intrinsic::nvvm_ff2bf16x2_rz_satfinite,
5706 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu_satfinite,
5708 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5709 llvm::Intrinsic::nvvm_ff2bf16x2_rs,
5710 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu,
5711 llvm::Intrinsic::nvvm_ff2bf16x2_rs_satfinite,
5712 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu_satfinite,
5715 unsigned hasRelu = op.getRelu() ? 1 : 0;
5716 unsigned hasSatFinite =
5717 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5720 unsigned idx = (hasSatFinite << 1) | hasRelu;
5725 if (op.getRandomBits())
5726 args.push_back(mt.
lookupValue(op.getRandomBits()));
5729 args.push_back(builder.getInt1(
false));
5731 switch (op.getRnd()) {
5732 case FPRoundingMode::RN:
5733 return {rndRNIds[idx], std::move(args)};
5734 case FPRoundingMode::RZ:
5735 return {rndRZIds[idx], std::move(args)};
5736 case FPRoundingMode::RS:
5737 return {rndRSIds[idx], std::move(args)};
5739 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToBF16x2Op");
5743llvm::Intrinsic::ID ConvertF32x4ToF8x4Op::getIntrinsicID() {
5745 bool hasRelu = getRelu();
5748 .Case([&](mlir::Float8E4M3FNType) {
5749 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_relu_satfinite
5750 : llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_satfinite;
5752 .Case([&](mlir::Float8E5M2Type) {
5753 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_relu_satfinite
5754 : llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_satfinite;
5757 llvm_unreachable(
"Invalid F8 type in ConvertF32x4ToF8x4Op");
5758 return llvm::Intrinsic::not_intrinsic;
5762llvm::Intrinsic::ID ConvertF32x4ToF6x4Op::getIntrinsicID() {
5764 bool hasRelu = getRelu();
5767 .Case([&](mlir::Float6E2M3FNType) {
5768 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_relu_satfinite
5769 : llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_satfinite;
5771 .Case([&](mlir::Float6E3M2FNType) {
5772 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_relu_satfinite
5773 : llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_satfinite;
5776 llvm_unreachable(
"Invalid F6 type in ConvertF32x4ToF6x4Op");
5777 return llvm::Intrinsic::not_intrinsic;
5781llvm::Intrinsic::ID ConvertF32x4ToF4x4Op::getIntrinsicID() {
5783 bool hasRelu = getRelu();
5786 .Case([&](mlir::Float4E2M1FNType) {
5787 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_relu_satfinite
5788 : llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_satfinite;
5791 llvm_unreachable(
"Invalid F4 type in ConvertF32x4ToF4x4Op");
5792 return llvm::Intrinsic::not_intrinsic;
5796llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(
Operation &op) {
5797 auto curOp = cast<NVVM::Tcgen05CpOp>(op);
5798 bool is2CTA = curOp.getGroup() == CTAGroupKind::CTA_2;
5799 auto srcFmt = curOp.getSrcFormat();
5800 auto mc = curOp.getMulticast();
5802 switch (curOp.getShape()) {
5803 case Tcgen05CpShape::SHAPE_128x256b:
5805 case Tcgen05CpShape::SHAPE_128x128b:
5807 case Tcgen05CpShape::SHAPE_4x256b:
5809 case Tcgen05CpShape::SHAPE_32x128b:
5811 case Tcgen05CpShape::SHAPE_64x128b:
5812 return (mc == Tcgen05CpMulticast::WARPX2_01_23)
5816 llvm_unreachable(
"Invalid shape in tcgen05 cp Op");
5823 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
5825 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
5830LogicalResult Tcgen05LdOp::verify() {
5832 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5835 if (
getShape() != NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && getOffset())
5836 result =
emitError(
"offset argument is only supported for shape 16x32bx2");
5838 auto resTy = getRes().getType();
5839 unsigned resLen = isa<VectorType>(resTy)
5840 ? llvm::cast<VectorType>(resTy).getNumElements()
5843 result =
emitError(llvm::formatv(
"invalid result type length {0} for shape "
5844 "{1} in tcgen05.ld Op",
5845 resLen, stringifyEnum(
getShape())));
5850LogicalResult Tcgen05StOp::verify() {
5852 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5855 auto valTy = getVal().getType();
5856 unsigned valLen = isa<VectorType>(valTy)
5857 ? llvm::cast<VectorType>(valTy).getNumElements()
5860 result =
emitError(llvm::formatv(
"invalid input length {0} for shape "
5861 "{1} in tcgen05.st Op",
5862 valLen, stringifyEnum(
getShape())));
5870 std::optional<LLVM::ConstantRangeAttr> range,
Value result,
5874 setResultRanges(
result, {range->getLower(), range->getUpper(),
5875 range->getLower(), range->getUpper()});
5884 Operation *op, std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
5888 const llvm::APInt &lower = rangeAttr->getLower();
5889 const llvm::APInt &upper = rangeAttr->getUpper();
5892 if (lower == upper && !lower.isMaxValue() && !lower.isMinValue()) {
5893 unsigned bitWidth = lower.getBitWidth();
5894 llvm::APInt minVal = llvm::APInt::getMinValue(bitWidth);
5895 llvm::APInt maxVal = llvm::APInt::getMaxValue(bitWidth);
5897 "invalid range attribute: Lower == Upper, but they aren't min (")
5898 << llvm::toString(minVal, 10,
false) <<
") or max ("
5899 << llvm::toString(maxVal, 10,
false)
5900 <<
") value! This is an invalid constant range.";
5907 llvm::IRBuilderBase &builder) {
5908 return builder.CreateBitCast(arg,
5909 llvm::Type::getInt32Ty(builder.getContext()));
5914 auto curOp = cast<NVVM::DotAccumulate4WayOp>(op);
5921 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
5922 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
5923 unsigned type = (isASigned << 1) | isBSigned;
5924 const llvm::Intrinsic::ID ids[] = {
5925 llvm::Intrinsic::nvvm_idp4a_u_u,
5926 llvm::Intrinsic::nvvm_idp4a_u_s,
5927 llvm::Intrinsic::nvvm_idp4a_s_u,
5928 llvm::Intrinsic::nvvm_idp4a_s_s,
5930 return {ids[type], args};
5935 auto curOp = cast<NVVM::DotAccumulate2WayOp>(op);
5940 args.push_back(builder.getInt1(curOp.getBHi()));
5943 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
5944 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
5945 unsigned type = (isASigned << 1) | isBSigned;
5946 const llvm::Intrinsic::ID ids[] = {
5947 llvm::Intrinsic::nvvm_idp2a_u_u,
5948 llvm::Intrinsic::nvvm_idp2a_u_s,
5949 llvm::Intrinsic::nvvm_idp2a_s_u,
5950 llvm::Intrinsic::nvvm_idp2a_s_s,
5952 return {ids[type], args};
5956 llvm::IRBuilderBase &builder) {
5957 return builder.CreateAddrSpaceCast(
5958 addr, builder.getPtrTy(llvm::NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM));
5962PrefetchOp::getIntrinsicIDAndArgs(NVVM::PrefetchOp &op,
5964 llvm::IRBuilderBase &builder) {
5965 using MemSpace = NVVM::NVVMMemorySpace;
5966 using CacheLevel = NVVM::PrefetchCacheLevel;
5968 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = op.getCacheLevel();
5969 std::optional<NVVM::CacheEvictionPriority> evictPriority =
5970 op.getEvictPriority();
5971 unsigned addressSpace =
5972 llvm::cast<LLVM::LLVMPointerType>(op.getAddr().getType())
5980 if (op.getTensormap())
5981 return {llvm::Intrinsic::nvvm_prefetch_tensormap, args};
5983 assert(cacheLevel &&
"expected cache level for non-tensormap prefetch");
5985 if (op.getUniform() && *cacheLevel == CacheLevel::L1)
5986 return {llvm::Intrinsic::nvvm_prefetchu_L1, args};
5988 if (evictPriority && *cacheLevel == CacheLevel::L2) {
5989 switch (*evictPriority) {
5990 case NVVM::CacheEvictionPriority::EvictLast:
5991 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_last, args};
5992 case NVVM::CacheEvictionPriority::EvictNormal:
5993 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_normal, args};
5995 llvm_unreachable(
"Invalid cache eviction priority");
5999 switch (
static_cast<MemSpace
>(addressSpace)) {
6000 case MemSpace::Generic:
6001 return *cacheLevel == CacheLevel::L1
6003 :
NVVM::
IDArgPair({llvm::Intrinsic::nvvm_prefetch_L2, args});
6004 case MemSpace::Global:
6005 return *cacheLevel == CacheLevel::L1
6007 {llvm::Intrinsic::nvvm_prefetch_global_L1, args})
6009 {llvm::Intrinsic::nvvm_prefetch_global_L2, args});
6010 case MemSpace::Local:
6011 return *cacheLevel == CacheLevel::L1
6013 {llvm::Intrinsic::nvvm_prefetch_local_L1, args})
6015 {llvm::Intrinsic::nvvm_prefetch_local_L2, args});
6017 llvm_unreachable(
"Invalid pointer address space");
6021bool NVVM::InlinePtxOp::getAsmValues(
6025 for (
auto arg : getReadWriteArgs())
6027 for (
auto arg : getResults())
6029 for (
auto arg : getReadOnlyArgs())
6036NVVM::IDArgPair ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
6038 auto curOp = cast<NVVM::ClusterLaunchControlTryCancelOp>(op);
6040 args.push_back(mt.
lookupValue(curOp.getSmemAddress()));
6041 args.push_back(mt.
lookupValue(curOp.getMbarrier()));
6043 llvm::Intrinsic::ID intrinsicID =
6044 curOp.getMulticast()
6046 nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
6047 : llvm::Intrinsic::nvvm_clusterlaunchcontrol_try_cancel_async_shared;
6049 return {intrinsicID, args};
6052NVVM::IDArgPair ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
6054 auto curOp = cast<NVVM::ClusterLaunchControlQueryCancelOp>(op);
6056 args.push_back(mt.
lookupValue(curOp.getTryCancelResponse()));
6058 llvm::Intrinsic::ID intrinsicID;
6060 switch (curOp.getQueryType()) {
6061 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
6063 llvm::Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled;
6065 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
6066 intrinsicID = llvm::Intrinsic::
6067 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x;
6069 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
6070 intrinsicID = llvm::Intrinsic::
6071 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y;
6073 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
6074 intrinsicID = llvm::Intrinsic::
6075 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z;
6078 return {intrinsicID, args};
6083 llvm::IRBuilderBase &builder) {
6084 auto thisOp = cast<NVVM::PermuteOp>(op);
6085 NVVM::PermuteMode mode = thisOp.getMode();
6087 static constexpr llvm::Intrinsic::ID IDs[] = {
6088 llvm::Intrinsic::nvvm_prmt, llvm::Intrinsic::nvvm_prmt_f4e,
6089 llvm::Intrinsic::nvvm_prmt_b4e, llvm::Intrinsic::nvvm_prmt_rc8,
6090 llvm::Intrinsic::nvvm_prmt_ecl, llvm::Intrinsic::nvvm_prmt_ecr,
6091 llvm::Intrinsic::nvvm_prmt_rc16};
6093 unsigned modeIndex =
static_cast<unsigned>(mode);
6101 args.push_back(mt.
lookupValue(thisOp.getSelector()));
6103 return {IDs[modeIndex], args};
6108 auto thisOp = cast<NVVM::TensormapReplaceOp>(op);
6112 if (thisOp.getOrd())
6113 args.push_back(builder.getInt32(thisOp.getOrd().value()));
6114 if (thisOp.getNewValue())
6115 args.push_back(mt.
lookupValue(thisOp.getNewValue()));
6116 if (
auto attr = thisOp.getNewValueAttr()) {
6119 .Case<TensormapElemtypeAttr, TensormapInterleaveLayoutAttr,
6120 TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
6121 TensormapFillModeAttr>([](
auto attr) {
6122 return static_cast<unsigned>(attr.getValue());
6124 .Default([](
auto attr) {
6125 llvm_unreachable(
"Invalid attribute type");
6128 args.push_back(builder.getInt32(val));
6131 static constexpr llvm::Intrinsic::ID IDs[] = {
6132 llvm::Intrinsic::nvvm_tensormap_replace_global_address,
6133 llvm::Intrinsic::nvvm_tensormap_replace_rank,
6134 llvm::Intrinsic::nvvm_tensormap_replace_box_dim,
6135 llvm::Intrinsic::nvvm_tensormap_replace_global_dim,
6136 llvm::Intrinsic::nvvm_tensormap_replace_global_stride,
6137 llvm::Intrinsic::nvvm_tensormap_replace_element_stride,
6138 llvm::Intrinsic::nvvm_tensormap_replace_elemtype,
6139 llvm::Intrinsic::nvvm_tensormap_replace_interleave_layout,
6140 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_mode,
6141 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_atomicity,
6142 llvm::Intrinsic::nvvm_tensormap_replace_fill_mode,
6145 unsigned fieldIndex =
static_cast<unsigned>(thisOp.getField());
6147 return {IDs[fieldIndex], args};
6154static llvm::nvvm::Tcgen05MMAKind
6157 case NVVM::Tcgen05MMAKind::F16:
6158 return llvm::nvvm::Tcgen05MMAKind::F16;
6159 case NVVM::Tcgen05MMAKind::TF32:
6160 return llvm::nvvm::Tcgen05MMAKind::TF32;
6161 case NVVM::Tcgen05MMAKind::F8F6F4:
6162 return llvm::nvvm::Tcgen05MMAKind::F8F6F4;
6163 case NVVM::Tcgen05MMAKind::I8:
6164 return llvm::nvvm::Tcgen05MMAKind::I8;
6165 case NVVM::Tcgen05MMAKind::TI16:
6166 return llvm::nvvm::Tcgen05MMAKind::TI16;
6167 case NVVM::Tcgen05MMAKind::MXF8F6F4:
6168 case NVVM::Tcgen05MMAKind::MXF4:
6169 case NVVM::Tcgen05MMAKind::MXF4NVF4:
6172 llvm_unreachable(
"Unsupported tcgen05.mma kind");
6178 llvm::IRBuilderBase &builder) {
6180 auto thisOp = cast<NVVM::Tcgen05MMAOp>(op);
6183 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6186 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6189 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6190 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6191 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6193 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6194 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6195 using IsATensorArray = std::array<CtaGroupArray, 2>;
6196 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6197 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6200 static constexpr HasDisableOutputLaneArray tcgen05MMAIDs = {
6206 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic},
6208 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic}}},
6212 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6213 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6217 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6218 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6224 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic},
6226 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic}}},
6230 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6231 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6235 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6236 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6242 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1,
6245 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2,
6250 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1,
6252 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift,
6257 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2,
6259 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift,
6265 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1,
6269 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2,
6274 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1,
6276 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift},
6280 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2,
6282 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift,
6285 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6286 bool hasScaleInputD = ScaleInputD !=
nullptr;
6288 llvm::Value *DisableOutputLane =
6290 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6292 const unsigned ctaGroup =
6295 llvm::Intrinsic::ID ID =
6296 tcgen05MMAIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6297 [ctaGroup - 1][thisOp.getAShift()];
6299 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMAOp.");
6302 args.push_back(ScaleInputD);
6304 if (hasDisableOutputLane)
6305 args.push_back(DisableOutputLane);
6307 args.push_back(builder.getInt32(
6310 if (!hasDisableOutputLane)
6311 args.push_back(builder.getInt32(ctaGroup));
6314 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6317 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6324 NVVM::CTAGroupKind ctaGroup,
bool hasAShift,
6325 NVVM::Tcgen05MMACollectorOp collectorOp,
Location loc) {
6327 if (disableOutputLane) {
6328 mlir::VectorType disableOutputLaneType =
6329 cast<mlir::VectorType>(disableOutputLane.
getType());
6330 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6331 disableOutputLaneType.getNumElements() != 4) ||
6332 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6333 disableOutputLaneType.getNumElements() != 8))
6334 return emitError(loc) <<
"Disable Output Lane of length "
6335 << disableOutputLaneType.getNumElements()
6336 <<
" is incompatible with CtaGroupAttr";
6339 if (hasAShift && !isATensor)
6341 loc,
"A-shift can be applied only when matrix A is in tensor memory");
6343 if (hasAShift ==
true && (collectorOp == Tcgen05MMACollectorOp::FILL ||
6344 collectorOp == Tcgen05MMACollectorOp::USE))
6346 loc,
"Cannot use collector buffer operation fill or use with ashift");
6351LogicalResult Tcgen05MMAOp::verify() {
6353 getDisableOutputLane(), getCtaGroup(), getAShift(),
6354 getCollectorOp(), getLoc());
6364 auto thisOp = cast<NVVM::Tcgen05MMASparseOp>(op);
6367 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6370 bool isATensor = isa<llvm::PointerType>(
A->getType());
6373 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6374 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6375 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6376 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6378 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6379 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6380 using IsATensorArray = std::array<CtaGroupArray, 2>;
6381 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6382 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6385 static constexpr HasDisableOutputLaneArray tcgen05MMASparseIDs = {
6391 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic},
6393 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic}}},
6397 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6398 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6402 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6403 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6409 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6412 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6417 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6418 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6422 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6423 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6430 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1,
6434 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2,
6439 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1,
6441 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift,
6446 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2,
6448 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift,
6454 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1,
6458 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2,
6463 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1,
6465 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift},
6469 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2,
6471 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift,
6474 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6475 bool hasScaleInputD = ScaleInputD !=
nullptr;
6477 llvm::Value *DisableOutputLane =
6479 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6484 llvm::Intrinsic::ID ID =
6485 tcgen05MMASparseIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6486 [ctaGroup - 1][thisOp.getAShift()];
6488 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMASparseOp.");
6491 args.push_back(ScaleInputD);
6493 if (hasDisableOutputLane)
6494 args.push_back(DisableOutputLane);
6496 args.push_back(builder.getInt32(
6499 if (!hasDisableOutputLane)
6500 args.push_back(builder.getInt32(ctaGroup));
6503 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6506 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6511LogicalResult Tcgen05MMASparseOp::verify() {
6513 getDisableOutputLane(), getCtaGroup(), getAShift(),
6514 getCollectorOp(), getLoc());
6524 auto thisOp = cast<NVVM::Tcgen05MMABlockScaleOp>(op);
6527 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6530 bool isATensor = isa<llvm::PointerType>(
A->getType());
6533 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6534 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6535 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6536 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6537 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6538 args.push_back(builder.getInt32(
6541 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6543 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6545 auto kind = thisOp.getKind();
6546 auto blockScale = thisOp.getBlockScale();
6547 llvm::Intrinsic::ID ID = [&]() {
6548 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6549 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6550 return isATensor ? llvm::Intrinsic::
6551 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale
6553 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale;
6554 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6557 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32
6559 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32;
6561 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6562 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6564 ? llvm::Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale
6565 : llvm::Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale;
6566 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6567 return isATensor ? llvm::Intrinsic::
6568 nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32
6570 nvvm_tcgen05_mma_shared_mxf4_block_scale_block32;
6572 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6573 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6576 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32
6578 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32;
6580 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6583 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16
6585 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16;
6588 llvm_unreachable(
"Invalid tcgen05.mma.block_scale attributes");
6595 NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind,
6596 NVVM::Tcgen05MMABlockScale blockScale,
Location loc) {
6597 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT &&
6598 kind == NVVM::Tcgen05MMAKind::MXF4NVF4)
6599 return emitError(loc,
"mxf4nvf4 requires block scale attribute");
6601 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16 &&
6602 kind != NVVM::Tcgen05MMAKind::MXF4NVF4)
6604 llvm::formatv(
"{} kind does not support block16 attribute",
6605 stringifyEnum(kind)));
6610LogicalResult Tcgen05MMABlockScaleOp::verify() {
6612 getBlockScale(), getLoc());
6622 auto thisOp = cast<NVVM::Tcgen05MMASparseBlockScaleOp>(op);
6625 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6628 bool isATensor = isa<llvm::PointerType>(
A->getType());
6631 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6632 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6633 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6634 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6635 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6636 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6637 args.push_back(builder.getInt32(
6640 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6642 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6644 auto kind = thisOp.getKind();
6645 auto blockScale = thisOp.getBlockScale();
6646 llvm::Intrinsic::ID ID = [&]() {
6647 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6648 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6649 return isATensor ? llvm::Intrinsic::
6650 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale
6652 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale;
6653 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6656 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32
6658 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32;
6660 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6661 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6662 return isATensor ? llvm::Intrinsic::
6663 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale
6665 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale;
6666 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6669 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32
6671 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32;
6673 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6674 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6677 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32
6679 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32;
6681 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6684 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16
6686 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16;
6689 llvm_unreachable(
"Invalid tcgen05.mma.sp.block_scale attributes");
6695LogicalResult Tcgen05MMASparseBlockScaleOp::verify() {
6697 getBlockScale(), getLoc());
6707 auto thisOp = cast<NVVM::Tcgen05MMAWsOp>(op);
6710 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6713 bool isATensor = isa<llvm::PointerType>(
A->getType());
6716 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6717 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6718 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6720 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6724 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor_zero_col_mask
6725 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared_zero_col_mask;
6727 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor
6728 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared;
6730 args.push_back(builder.getInt32(
6733 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6735 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6747 auto thisOp = cast<NVVM::Tcgen05MMAWsSparseOp>(op);
6750 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6753 bool isATensor = isa<llvm::PointerType>(
A->getType());
6756 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6757 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6758 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6759 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6761 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6766 ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor_zero_col_mask
6767 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared_zero_col_mask;
6769 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor
6770 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared;
6772 args.push_back(builder.getInt32(
6775 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6777 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6788 auto thisOp = cast<Tcgen05MMADecompressBOp>(op);
6791 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6794 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6797 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6798 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6799 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6800 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6802 llvm::Value *DisableOutputLane =
6804 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6806 NVVM::CTAGroupKind ctaGroup = thisOp.getCtaGroup();
6808 using namespace llvm::Intrinsic;
6809 ID intrinsicID = not_intrinsic;
6811 if (hasDisableOutputLane) {
6812 if (ctaGroup == NVVM::CTAGroupKind::CTA_1) {
6815 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b
6816 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b;
6817 }
else if (ctaGroup == NVVM::CTAGroupKind::CTA_2) {
6820 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b
6821 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b;
6823 llvm_unreachable(
"Unknown ctaGroup for tcgen05.mma.decompress_b");
6826 intrinsicID = isATensor ? nvvm_tcgen05_mma_tensor_f8f6f4_decompress_b
6827 : nvvm_tcgen05_mma_shared_f8f6f4_decompress_b;
6830 assert(intrinsicID != not_intrinsic &&
6831 "Invalid intrinsic for Tcgen05MMADecompressBOp.");
6833 if (hasDisableOutputLane)
6834 args.push_back(DisableOutputLane);
6840 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6842 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6844 return {intrinsicID, args};
6847LogicalResult Tcgen05MMADecompressBOp::verify() {
6848 mlir::Value disableOutputLane = getDisableOutputLane();
6850 if (disableOutputLane) {
6851 NVVM::CTAGroupKind ctaGroup = getCtaGroup();
6853 mlir::VectorType disableOutputLaneType =
6854 cast<mlir::VectorType>(disableOutputLane.
getType());
6855 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6856 disableOutputLaneType.getNumElements() != 4) ||
6857 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6858 disableOutputLaneType.getNumElements() != 8))
6859 return emitOpError() <<
"Disable Output Lane of length "
6860 << disableOutputLaneType.getNumElements()
6861 <<
" is incompatible with CtaGroupAttr";
6873 auto thisOp = cast<Tcgen05MMABlockScaleDecompressBOp>(op);
6876 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6879 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6882 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6883 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6884 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6885 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6886 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6887 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6888 args.push_back(builder.getInt32(
6891 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6893 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6895 llvm::Intrinsic::ID intrinsicID =
6898 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32_decompress_b
6900 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32_decompress_b;
6902 return {intrinsicID, args};
6909#define TCGEN05LDRED(SHAPE, NUM, TYPE) \
6910 llvm::Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_##NUM##_##TYPE
6914 auto thisOp = cast<NVVM::Tcgen05LdRedOp>(op);
6917 mlir::VectorType VecResTy =
6918 cast<mlir::VectorType>(thisOp.getData().getType());
6919 unsigned Num = VecResTy.getNumElements();
6920 bool IsFloat = thisOp.getRedVal().getType().isF32();
6922 llvm::Intrinsic::ID Shape32x32b[][2] = {
6933 llvm::Intrinsic::ID Shape16x32bx2[][2] = {
6944 NVVM::Tcgen05LdStShape
shape = thisOp.getShape();
6945 unsigned ID = [&]() {
6948 unsigned idx = std::log2(Num);
6950 case NVVM::Tcgen05LdStShape::SHAPE_32X32B:
6951 return Shape32x32b[idx][IsFloat];
6952 case NVVM::Tcgen05LdStShape::SHAPE_16X32BX2:
6953 return Shape16x32bx2[idx][IsFloat];
6955 llvm_unreachable(
"unhandled tcgen05.ld lowering");
6961 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2)
6962 args.push_back(mt.
lookupValue(thisOp.getOffset()));
6965 builder.getInt32(thisOp.getOp() == NVVM::ReductionKind::MIN ? 0 : 1));
6968 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getAbs())));
6969 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getNan())));
6974LogicalResult Tcgen05LdRedOp::verify() {
6975 VectorType data = cast<VectorType>(getData().
getType());
6976 Type redVal = getRedVal().getType();
6978 if (data.getElementType() != redVal)
6980 "type of reduction value and element type of vector data should match");
6982 if (getOp() != NVVM::ReductionKind::MIN &&
6983 getOp() != NVVM::ReductionKind::MAX)
6984 return emitError(
"only min and max reduction kinds are supported");
6986 if (redVal.
isInteger() && (getAbs() || getNan())) {
6987 return emitError(
"abs or nan is only applicable for f32 type");
6997struct NVVMInlinerInterface final : DialectInlinerInterface {
6998 using DialectInlinerInterface::DialectInlinerInterface;
6999 bool isLegalToInline(Operation *, Region *,
bool, IRMapping &)
const final {
7006void NVVMDialect::initialize() {
7007 registerNVVMDialectOperations(
this);
7009#define GET_ATTRDEF_LIST
7010#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
7015 allowUnknownOperations();
7016 addInterfaces<NVVMInlinerInterface>();
7017 declarePromisedInterface<ConvertToLLVMPatternInterface, NVVMDialect>();
7018 declarePromisedInterface<gpu::TargetAttrInterface, NVVMTargetAttr>();
7021LogicalResult NVVMDialect::verifyOperationAttribute(
Operation *op,
7023 StringAttr attrName = attr.
getName();
7025 if (attrName == NVVMDialect::getKernelFuncAttrName()) {
7026 if (!isa<LLVM::LLVMFuncOp>(op)) {
7027 return op->
emitError() <<
"'" << NVVMDialect::getKernelFuncAttrName()
7028 <<
"' attribute attached to unexpected op";
7033 if (attrName == NVVMDialect::getMaxntidAttrName() ||
7034 attrName == NVVMDialect::getReqntidAttrName() ||
7035 attrName == NVVMDialect::getClusterDimAttrName()) {
7036 auto values = llvm::dyn_cast<DenseI32ArrayAttr>(attr.
getValue());
7037 if (!values || values.empty() || values.size() > 3) {
7040 <<
"' attribute must be integer array with maximum 3 index";
7045 if (attrName == NVVMDialect::getMinctasmAttrName() ||
7046 attrName == NVVMDialect::getMaxnregAttrName() ||
7047 attrName == NVVMDialect::getClusterMaxBlocksAttrName()) {
7048 if (!llvm::dyn_cast<IntegerAttr>(attr.
getValue())) {
7050 <<
"'" << attrName <<
"' attribute must be integer constant";
7054 if (attrName == NVVMDialect::getBlocksAreClustersAttrName()) {
7058 <<
"'" << attrName <<
"' attribute must be used along with " <<
"'"
7059 << NVVMDialect::getReqntidAttrName() <<
"' and " <<
"'"
7060 << NVVMDialect::getClusterDimAttrName() <<
"'";
7067LogicalResult NVVMDialect::verifyRegionArgAttribute(
Operation *op,
7068 unsigned regionIndex,
7071 auto funcOp = dyn_cast<FunctionOpInterface>(op);
7076 StringAttr attrName = argAttr.
getName();
7077 if (attrName == NVVM::NVVMDialect::getGridConstantAttrName()) {
7081 <<
"' attribute must be present only on kernel arguments";
7083 if (!isa<UnitAttr>(argAttr.
getValue()))
7084 return op->
emitError() <<
"'" << attrName <<
"' must be a unit attribute";
7085 if (!funcOp.getArgAttr(argIndex, LLVM::LLVMDialect::getByValAttrName())) {
7088 <<
"' attribute requires the argument to also have attribute '"
7089 << LLVM::LLVMDialect::getByValAttrName() <<
"'";
7100unsigned NVVMMemorySpaceAttr::getAddressSpace()
const {
7101 return static_cast<unsigned>(getValue());
7104bool NVVMMemorySpaceAttr::isValidLoad(
7105 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
7106 const ::mlir::DataLayout *dataLayout,
7112bool NVVMMemorySpaceAttr::isValidStore(
7113 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
7114 const ::mlir::DataLayout *dataLayout,
7120bool NVVMMemorySpaceAttr::isValidAtomicOp(
7121 ptr::AtomicBinOp op,
Type type, ptr::AtomicOrdering ordering,
7122 std::optional<int64_t> alignment, const ::mlir::DataLayout *dataLayout,
7125 assert(
false &&
"unimplemented, see TODO in the source.");
7129bool NVVMMemorySpaceAttr::isValidAtomicXchg(
7130 Type type, ptr::AtomicOrdering successOrdering,
7131 ptr::AtomicOrdering failureOrdering, std::optional<int64_t> alignment,
7132 const ::mlir::DataLayout *dataLayout,
7135 assert(
false &&
"unimplemented, see TODO in the source.");
7139bool NVVMMemorySpaceAttr::isValidAddrSpaceCast(
7143 assert(
false &&
"unimplemented, see TODO in the source.");
7147bool NVVMMemorySpaceAttr::isValidPtrIntCast(
7152 assert(
false &&
"unimplemented, see TODO in the source.");
7161 int optLevel, StringRef triple, StringRef chip,
7162 StringRef features, DictionaryAttr flags,
7164 if (optLevel < 0 || optLevel > 3) {
7165 emitError() <<
"The optimization level must be a number between 0 and 3.";
7168 if (triple.empty()) {
7169 emitError() <<
"The target triple cannot be empty.";
7173 emitError() <<
"The target chip cannot be empty.";
7176 if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
7177 return mlir::isa_and_nonnull<StringAttr>(attr);
7179 emitError() <<
"All the elements in the `link` array must be strings.";
7185LogicalResult NVVMTargetAttr::verifyTarget(
Operation *gpuModule) {
7186 if (!getVerifyTarget())
7189 auto gpuModuleOp = llvm::dyn_cast<gpu::GPUModuleOp>(gpuModule);
7192 "NVVM target attribute must be attached to a GPU module");
7195 std::optional<unsigned> targetFullSmVersion =
7197 if (!targetFullSmVersion)
7199 <<
"invalid NVVM target chip \"" << getChip()
7200 <<
"\", expected sm_<version>[a|f]";
7204 "Minimum NVVM target SM version is sm_20");
7208 ->
walk([&](Operation *op) {
7209 if (
auto reqOp = llvm::dyn_cast<NVVM::RequiresSMInterface>(op)) {
7210 const NVVMCheckSMVersion requirement =
7211 reqOp.getRequiredMinSMVersion();
7213 op->
emitOpError() <<
"is not supported on " << getChip();
7225#define GET_ATTRDEF_CLASSES
7226#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
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 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 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)
#define GET_F32x2_TO_F8X2_US_ID(rnd, has_satf)
static llvm::nvvm::Tcgen05MMAKind getNVVMTcgen05MMAKind(NVVM::Tcgen05MMAKind kind)
static llvm::Value * getParamCastedAddr(llvm::Value *addr, llvm::IRBuilderBase &builder)
static LogicalResult verifyAddSubFOp(OpType op)
static LogicalResult verifyTcgen05MMABlockScaleOp(NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind, NVVM::Tcgen05MMABlockScale blockScale, Location loc)
static void printMmaUnitProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword)
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)
#define GET_F32x2_TO_F6x2_ID(type, has_relu)
static llvm::Value * getAsPackedI32(llvm::Value *arg, llvm::IRBuilderBase &builder)
static void printMmaEnumProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword, AttrTy value)
#define GET_F16x2_TO_F8X2_ID(type, has_relu)
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)
#define GET_F32x2_TO_F8X2_S_ID(type, has_relu)
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 verifyConvertF32x2ToFP16x2Op(Twine dstType, FPRoundingMode rnd, bool hasRandomBits, Operation *op)
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.
This class represents a diagnostic that is inflight and set to be reported.
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(const ContainerType &container)
Print a comma separated list of operands.
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.
Location getLoc()
The source location the operation was defined or derived from.
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.
static WalkResult advance()
static WalkResult interrupt()
The OpAsmOpInterface, see OpAsmInterface.td for more details.
bool isValidLoadStoreImpl(Type type, ptr::AtomicOrdering ordering, std::optional< int64_t > alignment, const ::mlir::DataLayout *dataLayout, function_ref< InFlightDiagnostic()> emitError)
Checks whether the given type is an LLVM type that can be loaded or stored.
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).
void walk(Operation *op, function_ref< void(Region *)> callback, WalkOrder order)
Walk all of the regions, blocks, or operations nested under (and including) the given operation.
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.
llvm::function_ref< Fn > function_ref
LogicalResult matchAndRewrite(SubFOp op, PatternRewriter &rewriter) const override
static bool isMinimumSMVersion(unsigned fullSmVersion)
static std::optional< unsigned > getTargetFullSmVersionFromStr(StringRef smVersionString)
bool isCompatibleWith(const unsigned &targetFullSmVersion) const
OpRewritePattern(MLIRContext *context, PatternBenefit benefit=1, ArrayRef< StringRef > generatedNames={})
This represents an operation in an abstracted form, suitable for use with the builder APIs.