35#include "llvm/ADT/STLExtras.h"
36#include "llvm/ADT/TypeSwitch.h"
37#include "llvm/IR/IRBuilder.h"
38#include "llvm/IR/NVVMIntrinsicUtils.h"
39#include "llvm/Support/Casting.h"
40#include "llvm/Support/FormatVariadic.h"
41#include "llvm/Support/NVPTXAddrSpace.h"
42#include "llvm/Support/raw_ostream.h"
53#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.cpp.inc"
54#include "mlir/Dialect/LLVMIR/NVVMOpsEnums.cpp.inc"
56static constexpr unsigned notIntrinsic = llvm::Intrinsic::not_intrinsic;
63 auto ptrTy = llvm::cast<LLVM::LLVMPointerType>(
ptr.getType());
64 return ptrTy.getAddressSpace() ==
static_cast<unsigned>(targetAS);
81 NVVMMemorySpace targetAS) {
82 unsigned AS =
static_cast<unsigned>(targetAS);
83 return builder.CreateAddrSpaceCast(
84 ptr, llvm::PointerType::get(builder.getContext(), AS));
88static llvm::nvvm::CTAGroupKind
91 case NVVM::CTAGroupKind::CTA_1:
92 return llvm::nvvm::CTAGroupKind::CG_1;
93 case NVVM::CTAGroupKind::CTA_2:
94 return llvm::nvvm::CTAGroupKind::CG_2;
96 llvm_unreachable(
"unsupported cta_group value");
100 NVVM::CTAGroupKindAttr &groupAttr) {
104 std::optional<NVVM::CTAGroupKind> group =
105 NVVM::symbolizeCTAGroupKind(keyword);
108 groupAttr = NVVM::CTAGroupKindAttr::get(parser.
getContext(), *group);
113 NVVM::CTAGroupKindAttr groupAttr) {
114 printer << NVVM::stringifyCTAGroupKind(groupAttr.getValue());
117template <
typename AttrTy>
124 using EnumTy =
decltype(attr.getValue());
125 std::optional<EnumTy> value = NVVM::symbolizeEnum<EnumTy>(keyword);
127 return parser.
emitError(loc) <<
"unknown enum value '" << keyword <<
"'";
129 attr = AttrTy::get(parser.
getContext(), *value);
142 size_t numIm2ColOffsets,
144 if (tensorDims < 1 || tensorDims > 5)
145 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
153 "to use im2col mode, the tensor has to be at least 3-dimensional");
155 if (numIm2ColOffsets && (tensorDims != (numIm2ColOffsets + 2)))
157 loc,
"im2col offsets must be 2 less than number of coordinates");
166 if (!tensorSize.empty() && coordinates.size() != tensorSize.size()) {
168 emitError(loc,
"Expected coordinates size to be equal to tensor size");
171 if (!lowerStride.empty() && tensorSize.empty()) {
174 "Expected tensor_size to be present when lower_stride is provided");
175 }
else if (!lowerStride.empty() &&
176 lowerStride.size() != tensorSize.size() - 1) {
179 "Expected lower_stride size to be equal to one less than tensor size");
182 if (!lowerStride.empty() !=
static_cast<bool>(upperStride)) {
184 "Expected lower_stride and upper_stride to be either both "
185 "present or both absent");
188 bool isDimStride = tensorSize.size() > 0;
189 if (!isTile && isDimStride) {
191 loc,
"Only tile mode supports override address with dim and stride");
197LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
198 TMAStoreMode mode = getMode();
202 if (getPredicate()) {
203 if (mode != TMAStoreMode::TILE)
204 return emitError(
"Inline-ptx lowering supported only for Tile mode.");
205 if (getL2CacheHint())
206 return emitError(
"Inline-ptx lowering unsupported with L2 cache-hint.");
211 case TMAStoreMode::TILE:
213 case TMAStoreMode::IM2COL:
214 case TMAStoreMode::IM2COL_W:
216 case TMAStoreMode::TILE_SCATTER4:
218 return emitError(
"Scatter4 mode expects 5 coordinates");
223LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::verify() {
224 TMAStoreMode mode = getMode();
226 mode == TMAStoreMode::IM2COL || mode == TMAStoreMode::IM2COL_W;
227 bool isTile = mode == TMAStoreMode::TILE;
233 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
236 if (mode == TMAStoreMode::TILE_SCATTER4 &&
getCoordinates().size() != 5)
237 overrideAddrRes =
emitError(
"Mode tile scatter4 expects 5 coordinates");
242LogicalResult CpAsyncOp::verify() {
243 if (getModifier() != LoadCacheModifierKind::CG &&
244 getModifier() != LoadCacheModifierKind::CA)
245 return emitError(
"Only CG and CA cache modifiers are supported.");
246 if (getSize() != 4 && getSize() != 8 && getSize() != 16)
247 return emitError(
"expected byte size to be either 4, 8 or 16.");
248 if (getModifier() == LoadCacheModifierKind::CG && getSize() != 16)
249 return emitError(
"CG cache modifier is only support for 16 bytes copy.");
256 if (tensorDims < 1 || tensorDims > 5)
257 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
259 auto checkTMALoadParams = [&](TMALoadMode mode,
bool isIm2col,
260 size_t expectedIm2colOff) -> LogicalResult {
261 if (isIm2col && (tensorDims < 3))
264 <<
" mode, the tensor has to be at least 3-dimensional";
266 if (numIm2colOff != expectedIm2colOff)
267 return emitError(loc) <<
" im2col offsets expected " << expectedIm2colOff
268 <<
" (provided " << numIm2colOff <<
")";
274 case TMALoadMode::TILE:
275 return checkTMALoadParams(mode,
false, 0);
276 case TMALoadMode::IM2COL:
277 return checkTMALoadParams(mode,
true, tensorDims - 2);
278 case TMALoadMode::IM2COL_W:
279 case TMALoadMode::IM2COL_W_128:
280 return checkTMALoadParams(mode,
true, 2);
281 case TMALoadMode::TILE_GATHER4:
282 return (tensorDims == 5)
283 ? checkTMALoadParams(mode,
false, 0)
284 :
emitError(loc,
"Gather4 mode expects 5 coordinates");
289LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
291 getMode(), getLoc());
294LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
295 TMALoadMode mode = getMode();
296 bool isCTAOnly = getIsCTAOnly();
297 if (getPredicate()) {
299 return emitError(
"Predicate is supported only for shared::cluster mode.");
300 if (mode != TMALoadMode::TILE && mode != TMALoadMode::IM2COL)
302 "Predicate is supported only for Tile and Im2col modes.");
304 NVVMMemorySpace expectedAS =
305 isCTAOnly ? NVVMMemorySpace::Shared : NVVMMemorySpace::SharedCluster;
306 unsigned AS = llvm::cast<LLVM::LLVMPointerType>(getDstMem().
getType())
308 if (AS != expectedAS)
311 ?
"Shared::cta destination requires address-space 3."
312 :
"Shared::cluster destination requires address-space 7.");
315 if (getMulticastMask())
316 return emitError(
"Multicast is not supported with shared::cta mode.");
318 return emitError(
"CTAGroup is not supported with shared::cta mode.");
323 getMode(), getLoc());
326LogicalResult CpAsyncBulkTensorReduceOp::verify() {
327 TMAStoreMode mode = getMode();
330 case TMAStoreMode::TILE:
332 case TMAStoreMode::IM2COL:
333 case TMAStoreMode::IM2COL_W:
335 case TMAStoreMode::TILE_SCATTER4:
336 return emitError(
"Scatter mode unsupported for CpAsyncBulkTensorReduceOp");
341LogicalResult CpAsyncBulkTensorReduceOverrideAddrOp::verify() {
343 getMode() == TMAStoreMode::IM2COL || getMode() == TMAStoreMode::IM2COL_W;
344 bool isTile = getMode() == TMAStoreMode::TILE;
350 getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
353 if (getMode() == TMAStoreMode::TILE_SCATTER4)
355 "Scatter mode unsupported for CpAsyncBulkTensorReduceOverrideAddrOp");
360LogicalResult CpAsyncBulkGlobalToSharedClusterOp::verify() {
362 if (isSharedCTA && getMulticastMask())
363 return emitError(
"Multicast is not supported with shared::cta mode.");
369 NVVM::MemScopeKind scope,
370 Value retVal =
nullptr) {
371 if (scope != NVVM::MemScopeKind::CTA && scope != NVVM::MemScopeKind::CLUSTER)
372 return op->
emitError(
"mbarrier scope must be either CTA or Cluster");
375 bool hasRetValue =
static_cast<bool>(retVal);
376 if (isSharedCluster && hasRetValue)
378 "mbarrier in shared_cluster space cannot return any value");
383LogicalResult MBarrierArriveOp::verify() {
388LogicalResult MBarrierArriveDropOp::verify() {
393LogicalResult MBarrierArriveExpectTxOp::verify() {
397 if (getPredicate()) {
398 if (getScope() != NVVM::MemScopeKind::CTA)
399 return emitError(
"mbarrier scope must be CTA when using predicate");
402 return emitError(
"mbarrier in shared_cluster space is not supported when "
406 return emitError(
"return-value is not supported when using predicate");
408 if (getRelaxed() ==
true)
409 return emitError(
"mbarrier with relaxed semantics is not supported when "
416LogicalResult MBarrierArriveDropExpectTxOp::verify() {
431 inferredReturnTypes.push_back(IntegerType::get(context, 64));
436MBarrierArriveOp::inferReturnTypes(
MLIRContext *context,
437 std::optional<Location> location,
438 MBarrierArriveOp::Adaptor adaptor,
441 inferredReturnTypes);
444LogicalResult MBarrierArriveDropOp::inferReturnTypes(
445 MLIRContext *context, std::optional<Location> location,
446 MBarrierArriveDropOp::Adaptor adaptor,
449 inferredReturnTypes);
452LogicalResult MBarrierArriveExpectTxOp::inferReturnTypes(
453 MLIRContext *context, std::optional<Location> location,
454 MBarrierArriveExpectTxOp::Adaptor adaptor,
458 if (adaptor.getPredicate())
461 inferredReturnTypes);
464LogicalResult MBarrierArriveDropExpectTxOp::inferReturnTypes(
465 MLIRContext *context, std::optional<Location> location,
466 MBarrierArriveDropExpectTxOp::Adaptor adaptor,
469 inferredReturnTypes);
479 return inferred == actual;
488bool MBarrierArriveExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
492bool MBarrierArriveDropExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
497LogicalResult MBarrierExpectTxOp::verify() {
501LogicalResult MBarrierCompleteTxOp::verify() {
505LogicalResult MBarrierTestWaitOp::verify() {
509LogicalResult MBarrierTryWaitOp::verify() {
513LogicalResult ConvertFloatToTF32Op::verify() {
514 using RndMode = NVVM::FPRoundingMode;
518 return emitError(
"Relu not supported with rna rounding mode.");
525 "Only {rn,rz,rna} rounding modes supported for ConvertFloatToTF32Op.");
530LogicalResult ConvertF32x2ToF6x2Op::verify() {
533 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy())) {
534 return emitOpError(
"Only ")
535 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
536 << mlir::Float6E3M2FNType::get(ctx)
537 <<
" types are supported for conversions from f32x2 to f6x2.";
542LogicalResult ConvertF32x2ToF8x2Op::verify() {
543 using RndMode = NVVM::FPRoundingMode;
544 using SatMode = NVVM::SaturationMode;
546 bool isRoundingModeRN = getRnd() == RndMode::RN;
547 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
548 bool isRoundingModeRP = getRnd() == RndMode::RP;
549 bool isSatFinite = getSat() == SatMode::SATFINITE;
551 bool hasRelu = getRelu();
552 bool hasPzo = getPzo();
557 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
559 if (!(isRoundingModeRN || isRoundingModeRZ)) {
560 return emitOpError(
"Only RN and RZ rounding modes are supported "
561 "for conversions from f32x2 to ")
562 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
563 << mlir::Float8E5M2Type::get(ctx) <<
" types";
566 return emitOpError(
"Only SATFINITE saturation mode is supported "
569 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
570 << mlir::Float8E5M2Type::get(ctx) <<
" types";
574 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
575 if (!(isRoundingModeRZ || isRoundingModeRP)) {
576 return emitOpError(
"Only RZ and RP rounding modes are supported for "
577 "conversions from f32x2 to ")
578 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
580 if (hasRelu || hasPzo) {
581 return emitOpError() << (hasRelu ?
"relu" :
"pzo")
582 <<
" not supported for conversions to "
583 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
588 return emitOpError(
"Only ")
589 << mlir::Float8E4M3FNType::get(ctx) <<
", "
590 << mlir::Float8E5M2Type::get(ctx) <<
", and "
591 << mlir::Float8E8M0FNUType::get(ctx)
593 "supported for conversions from f32x2 to f8x2";
597LogicalResult ConvertF16x2ToF8x2Op::verify() {
600 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy())) {
601 return emitOpError(
"Only ")
602 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
603 << mlir::Float8E5M2Type::get(ctx)
604 <<
" types are supported for conversions from f16x2 to f8x2.";
609LogicalResult ConvertBF16x2ToF8x2Op::verify() {
610 using RndMode = NVVM::FPRoundingMode;
611 using SatMode = NVVM::SaturationMode;
613 bool isRoundingModeRN = getRnd() == RndMode::RN;
614 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
615 bool isRoundingModeRP = getRnd() == RndMode::RP;
616 bool isSatFinite = getSat() == SatMode::SATFINITE;
617 bool hasRelu = getRelu();
618 bool hasPzo = getPzo();
623 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
625 if (!(isRoundingModeRN || isRoundingModeRZ))
626 return emitOpError(
"Only RN and RZ rounding modes are supported "
627 "for conversions from bf16x2 to ")
628 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
629 << mlir::Float8E5M2Type::get(ctx) <<
" types";
631 return emitOpError(
"Only SATFINITE saturation mode is supported "
632 "for conversions from bf16x2 to ")
633 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
634 << mlir::Float8E5M2Type::get(ctx) <<
" types";
637 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
638 if (!(isRoundingModeRZ || isRoundingModeRP))
639 return emitOpError(
"Only RZ and RP rounding modes are supported for "
640 "conversions from bf16x2 to ")
641 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
642 if (hasRelu || hasPzo)
643 return emitOpError() << (hasRelu ?
"relu" :
"pzo")
644 <<
" not supported for conversions to "
645 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
649 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
654LogicalResult ConvertF32x2ToF4x2Op::verify() {
657 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
658 return emitOpError(
"Only ")
659 << mlir::Float4E2M1FNType::get(ctx)
660 <<
" type is supported for conversions from f32x2 to f4x2.";
665LogicalResult ConvertF8x2ToBF16x2Op::verify() {
667 if (llvm::isa<Float8E8M0FNUType>(getSrcType())) {
668 if (getSat() != SaturationMode::NONE)
670 "Only NONE saturation mode is supported for conversions from ")
671 << Float8E8M0FNUType::get(ctx) <<
" type";
672 if (getScaleFactor())
673 return emitOpError(
"scaleFactor not supported for conversions from ")
674 << Float8E8M0FNUType::get(ctx) <<
" type";
676 return emitOpError(
"relu not supported for conversions from ")
677 << Float8E8M0FNUType::get(ctx) <<
" type";
683LogicalResult PermuteOp::verify() {
684 using Mode = NVVM::PermuteMode;
685 bool hasHi =
static_cast<bool>(getHi());
692 return emitError(
"mode '") << getMode() <<
"' requires 'hi' operand.";
700 << getMode() <<
"' does not accept 'hi' operand.";
715 static constexpr FPRoundingMode validRndModes[] = {
716 FPRoundingMode::RN, FPRoundingMode::RZ, FPRoundingMode::RS};
718 if (!llvm::is_contained(validRndModes, rnd)) {
720 "Only RN, RZ, and RS rounding modes are supported for "
721 "conversions from f32x2 to ")
725 if (rnd == FPRoundingMode::RS) {
726 if (!hasRandomBits) {
727 return op->
emitOpError(
"random_bits is required for RS rounding mode.");
732 "random_bits not supported for RN and RZ rounding modes.");
739LogicalResult ConvertF32x2ToF16x2Op::verify() {
741 getRandomBits() ?
true :
false, *
this);
744LogicalResult ConvertF32x2ToBF16x2Op::verify() {
746 getRandomBits() ?
true :
false, *
this);
749LogicalResult ConvertF32x4ToF8x4Op::verify() {
752 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy()))
753 return emitOpError(
"Only ")
754 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
755 << mlir::Float8E5M2Type::get(ctx)
756 <<
" types are supported for conversions from f32x4 to f8x4.";
761LogicalResult ConvertF32x4ToF6x4Op::verify() {
764 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy()))
765 return emitOpError(
"Only ")
766 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
767 << mlir::Float6E3M2FNType::get(ctx)
768 <<
" types are supported for conversions from f32x4 to f6x4.";
773LogicalResult ConvertF32x4ToF4x4Op::verify() {
776 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
777 return emitOpError(
"Only ") << mlir::Float4E2M1FNType::get(ctx)
778 <<
" type is supported for conversions from "
784LogicalResult BulkStoreOp::verify() {
785 if (getInitVal() != 0)
786 return emitOpError(
"only 0 is supported for initVal, got ") << getInitVal();
790LogicalResult AsyncStoreGlobalOp::verify() {
791 NVVM::MemScopeKind scope = getScope();
792 bool isMmio = getMmio();
793 bool isMultimem = getMultimem();
795 if (scope != MemScopeKind::SYS && scope != MemScopeKind::GPU)
796 return emitOpError(
"scope must be either SYS or GPU");
798 if (isMmio && scope != MemScopeKind::SYS)
799 return emitOpError(
"mmio is only supported for SYS scope");
801 if (isMmio && isMultimem)
802 return emitOpError(
"multimem is not supported with mmio");
807LogicalResult PMEventOp::verify() {
808 auto eventId = getEventId();
809 auto maskedEventId = getMaskedEventId();
810 if (!maskedEventId && !eventId) {
811 return emitOpError() <<
"either `id` or `mask` must be set";
814 if (maskedEventId && eventId) {
815 return emitOpError() <<
"`id` and `mask` cannot be set at the same time";
819 if (eventId < 0 || eventId > 15) {
820 return emitOpError() <<
"`id` must be between 0 and 15";
824 return llvm::success();
830std::optional<mlir::NVVM::MMATypes>
831MmaOp::inferOperandMMAType(
Type operandElType,
bool isAccumulator) {
833 VectorType::get(2, Float16Type::get(operandElType.
getContext()));
834 if (operandElType.
isF64())
835 return NVVM::MMATypes::f64;
836 if (operandElType.
isF16() || operandElType == half2Type)
837 return NVVM::MMATypes::f16;
838 if (operandElType.
isF32() && isAccumulator)
839 return NVVM::MMATypes::f32;
840 if (operandElType.
isF32() && !isAccumulator)
841 return NVVM::MMATypes::tf32;
842 if (llvm::isa<IntegerType>(operandElType)) {
844 return NVVM::MMATypes::s32;
848 if (
auto structType = llvm::dyn_cast<LLVM::LLVMStructType>(operandElType)) {
849 if (structType.getBody().empty())
851 return inferOperandMMAType(structType.getBody()[0], isAccumulator);
858 return (type == MMATypes::u4 || type == MMATypes::s4);
862 return (type == MMATypes::u8 || type == MMATypes::s8);
867 type == MMATypes::s32;
870MMATypes MmaOp::accumPtxType() {
871 std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
872 getODSOperands(2).getTypes().front(),
true);
873 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
877MMATypes MmaOp::resultPtxType() {
878 std::optional<mlir::NVVM::MMATypes> val =
879 inferOperandMMAType(getResult().
getType(),
true);
880 assert(val.has_value() &&
"result PTX type should always be inferrable");
884template <
typename AttrTy>
886 StringRef keyword, AttrTy value) {
887 printer << (isFirst ?
" " :
", ") << keyword <<
" = ";
892template <
typename AttrTy>
894 StringRef keyword, AttrTy value) {
895 printer << (isFirst ?
" " :
", ") << keyword <<
" = "
896 << NVVM::stringifyEnum(value.getValue());
902 printer << (isFirst ?
" " :
", ") << keyword;
906template <
typename AttrTy>
910 if (attributes.
get(name))
912 "duplicate property '" + name +
"'");
916 attributes.
append(name, value);
920template <
typename AttrTy>
924 if (attributes.
get(name))
926 "duplicate property '" + name +
"'");
930 attributes.
append(name, value);
935 return llvm::is_contained(
937 "shape",
"b1Op",
"intOverflowBehavior",
"layoutA",
"layoutB",
938 "multiplicandAPtxType",
"multiplicandBPtxType",
"orderedMetadata",
939 "kind",
"scaleVecSize",
"blockScaleFormat",
"operandSegmentSizes"},
951 if (!llvm::is_contained(allowedKeywords, keyword))
953 "unknown MMA property '" + keyword +
"'");
955 ParseResult parseResult =
success();
956 if (keyword ==
"shape")
959 else if (keyword ==
"b1_op")
962 else if (keyword ==
"int_overflow")
964 parser, attributes,
"intOverflowBehavior");
965 else if (keyword ==
"layout_a")
968 else if (keyword ==
"layout_b")
971 else if (keyword ==
"multiplicand_a_ptx_type")
973 parser, attributes,
"multiplicandAPtxType");
974 else if (keyword ==
"multiplicand_b_ptx_type")
976 parser, attributes,
"multiplicandBPtxType");
977 else if (keyword ==
"kind") {
978 if (llvm::is_contained(allowedKeywords,
"block_scale_format"))
980 parser, attributes,
"kind");
984 }
else if (keyword ==
"scale_vec_size")
986 parser, attributes,
"scaleVecSize");
987 else if (keyword ==
"block_scale_format")
989 parser, attributes,
"blockScaleFormat");
990 else if (keyword ==
"ordered_metadata") {
991 if (attributes.
get(
"orderedMetadata"))
993 "duplicate property 'orderedMetadata'");
997 "unknown MMA property '" + keyword +
"'");
999 if (failed(parseResult))
1005 for (StringRef property : requiredProperties) {
1006 if (!attributes.
get(property))
1008 "missing required property '" + property +
"'");
1018 "inherent property '" + attribute.getName().getValue() +
1019 "' must be spelled directly in the operation syntax");
1020 attributes.
append(attribute);
1027 struct MMAOperandFragment {
1028 StringRef operandName;
1029 StringRef ptxTypeAttr;
1030 SmallVector<Value, 4> regs;
1031 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1032 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1035 std::array<MMAOperandFragment, 3> frags{
1036 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1037 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1038 MMAOperandFragment(
"C",
"")};
1040 mlir::NVVM::MmaOp::getOperandSegmentSizeAttr()};
1042 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
1043 auto &frag = frags[fragIdx];
1044 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1045 for (
auto operandIdx = varOperandSpec.first;
1046 operandIdx < varOperandSpec.first + varOperandSpec.second;
1048 frag.regs.push_back(this->getOperand(operandIdx));
1049 if (operandIdx == 0) {
1050 regTypes.push_back(this->getOperand(operandIdx).
getType());
1053 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1054 regTypes.back(), fragIdx >= 2);
1056 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1059 auto printMmaOperand = [&](
const MMAOperandFragment &frag) ->
void {
1060 p <<
" " << frag.operandName;
1066 for (
const auto &frag : frags) {
1067 printMmaOperand(frag);
1070 bool isFirstProperty =
true;
1074 if (getIntOverflowBehaviorAttr())
1076 getIntOverflowBehaviorAttr());
1079 if (getMultiplicandAPtxTypeAttr() &&
1080 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1082 getMultiplicandAPtxTypeAttr());
1083 if (getMultiplicandBPtxTypeAttr() &&
1084 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1086 getMultiplicandBPtxTypeAttr());
1087 llvm::append_range(ignoreAttrNames,
1089 getIntOverflowBehaviorAttrName(),
1090 getLayoutAAttrName(),
1091 getLayoutBAttrName(),
1092 getMultiplicandAPtxTypeAttrName(),
1093 getMultiplicandBPtxTypeAttrName()});
1100 frags[1].regs[0].getType(),
1101 frags[2].regs[0].getType()},
1110 std::optional<MMAIntOverflow> intOverflow,
1111 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
1112 std::optional<std::array<MMALayout, 2>> multiplicandLayouts) {
1114 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1119 result.addOperands(operandA);
1120 result.addOperands(operandB);
1121 result.addOperands(operandC);
1123 if (multiplicandPtxTypes) {
1124 result.addAttribute(
"multiplicandAPtxType",
1125 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1126 result.addAttribute(
"multiplicandBPtxType",
1127 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1129 if (
auto res = inferOperandMMAType(operandA[0].
getType(),
false))
1130 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1131 if (
auto res = inferOperandMMAType(operandB[0].
getType(),
false))
1132 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1135 if (multiplicandLayouts) {
1136 result.addAttribute(
"layoutA",
1137 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[0]));
1138 result.addAttribute(
"layoutB",
1139 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[1]));
1141 result.addAttribute(
"layoutA", MMALayoutAttr::get(ctx, MMALayout::row));
1142 result.addAttribute(
"layoutB", MMALayoutAttr::get(ctx, MMALayout::col));
1145 if (intOverflow.has_value())
1146 result.addAttribute(
"intOverflowBehavior",
1147 MMAIntOverflowAttr::get(ctx, *intOverflow));
1148 if (b1Op.has_value())
1149 result.addAttribute(
"b1Op", MMAB1OpAttr::get(ctx, *b1Op));
1151 result.addTypes(resultType);
1153 MmaOp::getOperandSegmentSizeAttr(),
1155 static_cast<int32_t>(operandB.size()),
1156 static_cast<int32_t>(operandC.size())}));
1165 struct MMAOperandFragment {
1166 std::optional<MMATypes> elemtype;
1167 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1168 SmallVector<Type> regTypes;
1172 std::array<MMAOperandFragment, 4> frags;
1178 MMAOperandFragment &frag) -> LogicalResult {
1197 {
"shape",
"b1_op",
"int_overflow",
"layout_a",
1198 "layout_b",
"multiplicand_a_ptx_type",
1199 "multiplicand_b_ptx_type"},
1200 {
"shape",
"layoutA",
"layoutB"}))
1212 if (operandTypes.size() != 3)
1215 "expected one type for each operand segment but got " +
1216 Twine(operandTypes.size()) +
" types");
1217 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1218 auto &frag = frags[iter.index()];
1219 frag.regTypes.resize(frag.regs.size(), iter.value());
1223 frag.elemtype = inferOperandMMAType(frag.regTypes[0],
1230 frags[3].elemtype = inferOperandMMAType(resultType,
true);
1232 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1233 "multiplicandBPtxType"};
1234 for (
unsigned idx = 0; idx < names.size(); idx++) {
1235 const auto &frag = frags[idx];
1236 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1237 if (!frag.elemtype.has_value() && !attr.has_value()) {
1240 "attribute " + names[idx] +
1241 " is not provided explicitly and cannot be inferred");
1243 if (!attr.has_value())
1245 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1248 result.addTypes(resultType);
1249 if (!namedAttributes.
empty())
1250 result.addAttributes(namedAttributes);
1251 result.addAttribute(MmaOp::getOperandSegmentSizeAttr(),
1253 static_cast<int32_t>(frags[0].regs.size()),
1254 static_cast<int32_t>(frags[1].regs.size()),
1255 static_cast<int32_t>(frags[2].regs.size()),
1260LogicalResult MmaOp::verify() {
1262 auto f16Ty = Float16Type::get(context);
1263 auto i32Ty = IntegerType::get(context, 32);
1264 auto f16x2Ty = VectorType::get(2, f16Ty);
1265 auto f32Ty = Float32Type::get(context);
1266 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1267 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1269 auto s32x4StructTy =
1270 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1271 auto f32x8StructTy =
1273 auto f16x2x2StructTy =
1274 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1275 auto f32x4StructTy =
1276 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1277 auto s32x2StructTy =
1278 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1280 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1281 getShapeAttr().getK()};
1287 AllowedShapes allowedShapes;
1288 AllowedTypes expectedA;
1289 AllowedTypes expectedB;
1290 AllowedTypes expectedC;
1295 if (mmaShape[0] == 16) {
1297 Type multiplicandFragType;
1298 switch (*getMultiplicandAPtxType()) {
1299 case MMATypes::tf32:
1301 multiplicandFragType = i32Ty;
1302 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1303 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1305 case MMATypes::bf16:
1307 multiplicandFragType = i32Ty;
1308 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1309 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1313 multiplicandFragType = f16x2Ty;
1314 expectedResult.push_back(f16x2x2StructTy);
1315 expectedResult.push_back(f32x4StructTy);
1317 case MMATypes::e4m3:
1318 case MMATypes::e5m2:
1322 multiplicandFragType = i32Ty;
1323 expectedResult.push_back(f16x2x2StructTy);
1324 expectedResult.push_back(f32x4StructTy);
1338 return emitError(
"invalid shape or multiplicand type: ")
1339 << getMultiplicandAPtxType().value();
1343 expectedResult.push_back(s32x4StructTy);
1344 expectedC.emplace_back(4, i32Ty);
1345 multiplicandFragType = i32Ty;
1347 expectedC.emplace_back(2, f16x2Ty);
1348 expectedC.emplace_back(4, f32Ty);
1351 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor);
1352 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1353 expectedA.emplace_back(unitA, multiplicandFragType);
1354 expectedB.emplace_back(unitB, multiplicandFragType);
1355 allowedShapes.push_back({16, 8, kFactor});
1356 allowedShapes.push_back({16, 8, kFactor * 2});
1358 if (resultPtxType() != accumPtxType())
1359 return emitOpError(
"ctype does not match dtype");
1363 if (mmaShape[0] == 8) {
1364 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1365 expectedA.emplace_back(2, f16x2Ty);
1366 expectedB.emplace_back(2, f16x2Ty);
1367 expectedResult.push_back(f16x2x4StructTy);
1368 expectedResult.push_back(f32x8StructTy);
1369 expectedC.emplace_back(4, f16x2Ty);
1370 expectedC.emplace_back(8, f32Ty);
1371 allowedShapes.push_back({8, 8, 4});
1373 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1374 Type f64Ty = Float64Type::get(context);
1375 expectedA.emplace_back(1, f64Ty);
1376 expectedB.emplace_back(1, f64Ty);
1377 expectedC.emplace_back(2, f64Ty);
1378 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1380 allowedShapes.push_back({8, 8, 4});
1383 expectedA.push_back({i32Ty});
1384 expectedB.push_back({i32Ty});
1385 expectedC.push_back({i32Ty, i32Ty});
1386 expectedResult.push_back(s32x2StructTy);
1388 allowedShapes.push_back({8, 8, 32});
1390 allowedShapes.push_back({8, 8, 16});
1391 if (getMultiplicandAPtxType().value() == MMATypes::b1)
1392 allowedShapes.push_back({8, 8, 128});
1396 std::string errorMessage;
1397 llvm::raw_string_ostream errorStream(errorMessage);
1400 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1401 !llvm::is_contained(allowedShapes, mmaShape)) {
1402 errorStream <<
"unimplemented variant for MMA shape <";
1403 llvm::interleaveComma(mmaShape, errorStream);
1405 return emitOpError(errorMessage);
1409 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1410 for (
const auto &iter : llvm::enumerate(
1411 std::array<AllowedTypes, 3>{std::move(expectedA),
1412 std::move(expectedB),
1413 std::move(expectedC)})) {
1414 auto spec = this->getODSOperandIndexAndLength(iter.index());
1416 operand_type_begin() + spec.first +
1418 bool match = llvm::is_contained(iter.value(), operandTySeg);
1421 errorStream <<
"Could not match types for the "
1422 << operandNames[iter.index()]
1423 <<
" operands; expected one of ";
1424 for (
const auto &x : iter.value()) {
1425 errorStream << x.size() <<
"x" << x[0] <<
" ";
1427 errorStream <<
"but got ";
1428 llvm::interleaveComma(operandTySeg, errorStream);
1429 return emitOpError(errorMessage);
1434 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1435 return expectedResultType == getResult().getType();
1438 <<
"Could not match allowed types for the result; expected one of ";
1439 llvm::interleaveComma(expectedResult, errorStream);
1440 errorStream <<
" but got " << getResult().getType();
1441 return emitOpError(errorMessage);
1445 if (getMultiplicandAPtxType() == MMATypes::b1 && !getB1Op()) {
1446 return emitOpError(
"op requires " + getB1OpAttrName().strref() +
1454 if (!getIntOverflowBehavior())
1455 return emitOpError(
"op requires " +
1456 getIntOverflowBehaviorAttrName().strref() +
1464 (mmaShape[0] == 8 && mmaShape[1] == 8 && mmaShape[2] == 4 &&
1465 getMultiplicandAPtxType() == MMATypes::f16);
1467 if (!isM8N8K4_F16) {
1469 if (getLayoutA() != MMALayout::row || getLayoutB() != MMALayout::col) {
1470 return emitOpError(
"requires layoutA = #nvvm.mma_layout<row> and "
1471 "layoutB = #nvvm.mma_layout<col> for shape <")
1472 << mmaShape[0] <<
", " << mmaShape[1] <<
", " << mmaShape[2]
1473 <<
"> with element types " << *getMultiplicandAPtxType() <<
" and "
1474 << *getMultiplicandBPtxType()
1475 <<
". Only m8n8k4 with f16 supports other layouts.";
1482MMATypes MmaSpOp::accumPtxType() {
1483 std::optional<mlir::NVVM::MMATypes> val = MmaOp::inferOperandMMAType(
1484 getODSOperands(2).getTypes().front(),
true);
1485 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
1489MMATypes MmaSpOp::resultPtxType() {
1490 std::optional<mlir::NVVM::MMATypes> val =
1491 MmaOp::inferOperandMMAType(getResult().
getType(),
true);
1492 assert(val.has_value() &&
"result PTX type should always be inferrable");
1498 llvm::IRBuilderBase &builder) {
1499 auto thisOp = cast<NVVM::MmaSpOp>(op);
1507 auto intId = MmaSpOp::getIntrinsicID(
1508 thisOp.getShape().getM(), thisOp.getShape().getN(),
1509 thisOp.getShape().getK(), thisOp.getIntOverflowBehavior(),
1510 thisOp.getOrderedMetadata(), thisOp.getKind(),
1511 *thisOp.getMultiplicandAPtxType(), *thisOp.getMultiplicandBPtxType(),
1512 thisOp.accumPtxType(), thisOp.resultPtxType());
1514 return {intId, args};
1519 struct MMAOperandFragment {
1520 StringRef operandName;
1521 StringRef ptxTypeAttr;
1522 SmallVector<Value, 4> regs;
1523 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1524 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1527 std::array<MMAOperandFragment, 5> frags{
1528 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1529 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1530 MMAOperandFragment(
"C",
""), MMAOperandFragment(
"sparseMetadata",
""),
1531 MMAOperandFragment(
"selector",
"")};
1533 mlir::NVVM::MmaSpOp::getOperandSegmentSizeAttr()};
1536 for (
unsigned fragIdx = 0; fragIdx < 3; fragIdx++) {
1537 auto &frag = frags[fragIdx];
1538 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1539 for (
auto operandIdx = varOperandSpec.first;
1540 operandIdx < varOperandSpec.first + varOperandSpec.second;
1542 frag.regs.push_back(this->getOperand(operandIdx));
1543 if (operandIdx == varOperandSpec.first) {
1544 regTypes.push_back(this->getOperand(operandIdx).
getType());
1547 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1548 regTypes.back(), fragIdx >= 2);
1550 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1554 frags[3].regs.push_back(getSparseMetadata());
1555 frags[4].regs.push_back(getSparsitySelector());
1557 auto printMmaSpOperand = [&](
const MMAOperandFragment &frag) ->
void {
1558 p <<
" " << frag.operandName;
1564 for (
const auto &frag : frags)
1565 printMmaSpOperand(frag);
1567 bool isFirstProperty =
true;
1569 if (getIntOverflowBehaviorAttr())
1571 getIntOverflowBehaviorAttr());
1572 if (getMultiplicandAPtxTypeAttr() &&
1573 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
1575 getMultiplicandAPtxTypeAttr());
1576 if (getMultiplicandBPtxTypeAttr() &&
1577 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
1579 getMultiplicandBPtxTypeAttr());
1580 if (getOrderedMetadata())
1587 getMultiplicandAPtxTypeAttrName(),
1588 getMultiplicandBPtxTypeAttrName(),
1589 getOrderedMetadataAttrName(), getKindAttrName()});
1594 for (
int i = 0; i < 3; ++i) {
1599 p <<
") -> " << getResult().getType();
1606 std::optional<MMAIntOverflow> intOverflow,
1607 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
1609 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1614 result.addOperands(operandA);
1615 result.addOperands(operandB);
1616 result.addOperands(operandC);
1617 result.addOperands(sparseMetadata);
1618 result.addOperands(sparsitySelector);
1620 if (multiplicandPtxTypes) {
1621 result.addAttribute(
"multiplicandAPtxType",
1622 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1623 result.addAttribute(
"multiplicandBPtxType",
1624 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1626 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
1627 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1628 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
1629 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1632 if (intOverflow.has_value())
1633 result.addAttribute(
"intOverflowBehavior",
1634 MMAIntOverflowAttr::get(ctx, *intOverflow));
1636 result.addTypes(resultType);
1638 MmaSpOp::getOperandSegmentSizeAttr(),
1640 static_cast<int32_t>(operandB.size()),
1641 static_cast<int32_t>(operandC.size()), 1,
1646 struct MMAOperandFragment {
1647 std::optional<MMATypes> elemtype;
1648 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1649 SmallVector<Type> regTypes;
1653 std::array<MMAOperandFragment, 6> frags;
1658 auto parseMmaSpOperand = [&](StringRef operandName,
1659 MMAOperandFragment &frag) -> LogicalResult {
1670 if (parseMmaSpOperand(
"A", frags[0]).
failed())
1672 if (parseMmaSpOperand(
"B", frags[1]).
failed())
1674 if (parseMmaSpOperand(
"C", frags[2]).
failed())
1676 if (parseMmaSpOperand(
"sparseMetadata", frags[3]).
failed())
1678 if (parseMmaSpOperand(
"selector", frags[4]).
failed())
1682 {
"shape",
"int_overflow",
"multiplicand_a_ptx_type",
1683 "multiplicand_b_ptx_type",
"ordered_metadata",
1698 if (operandTypes.size() != 3)
1701 "expected one type for each operand segment but got " +
1702 Twine(operandTypes.size()) +
" types");
1703 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1704 auto &frag = frags[iter.index()];
1705 frag.regTypes.resize(frag.regs.size(), iter.value());
1710 MmaOp::inferOperandMMAType(frag.regTypes[0],
1718 MmaOp::inferOperandMMAType(resultType,
true);
1733 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1734 "multiplicandBPtxType"};
1735 for (
unsigned idx = 0; idx < names.size(); idx++) {
1736 const auto &frag = frags[idx];
1737 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1738 if (!frag.elemtype.has_value() && !attr.has_value()) {
1741 "attribute " + names[idx] +
1742 " is not provided explicitly and cannot be inferred");
1744 if (!attr.has_value())
1746 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1749 result.addTypes(resultType);
1750 if (!namedAttributes.
empty())
1751 result.addAttributes(namedAttributes);
1752 result.addAttribute(MmaSpOp::getOperandSegmentSizeAttr(),
1754 static_cast<int32_t>(frags[0].regs.size()),
1755 static_cast<int32_t>(frags[1].regs.size()),
1756 static_cast<int32_t>(frags[2].regs.size()),
1763LogicalResult MmaSpOp::verify() {
1765 auto f16Ty = Float16Type::get(context);
1766 auto i32Ty = IntegerType::get(context, 32);
1767 auto f16x2Ty = VectorType::get(2, f16Ty);
1768 auto f32Ty = Float32Type::get(context);
1769 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1770 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1772 auto s32x4StructTy =
1773 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1774 auto f32x8StructTy =
1776 auto f16x2x2StructTy =
1777 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1778 auto f32x4StructTy =
1779 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1780 auto s32x2StructTy =
1781 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1783 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1784 getShapeAttr().getK()};
1790 AllowedShapes allowedShapes;
1791 AllowedTypes expectedA;
1792 AllowedTypes expectedB;
1793 AllowedTypes expectedC;
1798 if (mmaShape[0] == 16) {
1800 Type multiplicandFragType;
1801 switch (*getMultiplicandAPtxType()) {
1802 case MMATypes::tf32:
1804 multiplicandFragType = i32Ty;
1805 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1806 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1808 allowedShapes.push_back({16, 8, 8});
1809 allowedShapes.push_back({16, 8, 16});
1811 case MMATypes::bf16:
1813 multiplicandFragType = i32Ty;
1814 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1815 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1817 allowedShapes.push_back({16, 8, 16});
1818 allowedShapes.push_back({16, 8, 32});
1822 multiplicandFragType = f16x2Ty;
1823 expectedResult.push_back(f16x2x2StructTy);
1824 expectedResult.push_back(f32x4StructTy);
1826 allowedShapes.push_back({16, 8, 16});
1827 allowedShapes.push_back({16, 8, 32});
1833 allowedShapes.push_back({16, 8, 64});
1834 allowedShapes.push_back({16, 8, 128});
1840 allowedShapes.push_back({16, 8, 32});
1841 allowedShapes.push_back({16, 8, 64});
1843 case MMATypes::e4m3:
1844 case MMATypes::e5m2:
1845 case MMATypes::e3m2:
1846 case MMATypes::e2m3:
1847 case MMATypes::e2m1:
1849 multiplicandFragType = i32Ty;
1850 expectedResult.push_back(f16x2x2StructTy);
1851 expectedResult.push_back(f32x4StructTy);
1853 allowedShapes.push_back({16, 8, 64});
1856 return emitError(
"invalid shape or multiplicand type: ")
1857 << getMultiplicandAPtxType().value();
1861 expectedResult.push_back(s32x4StructTy);
1862 expectedC.emplace_back(4, i32Ty);
1863 multiplicandFragType = i32Ty;
1864 }
else if (*getMultiplicandAPtxType() >= MMATypes::e4m3 &&
1865 *getMultiplicandAPtxType() <= MMATypes::e2m1) {
1867 expectedC.emplace_back(2, f16x2Ty);
1868 expectedC.emplace_back(4, f32Ty);
1870 expectedC.emplace_back(2, f16x2Ty);
1871 expectedC.emplace_back(4, f32Ty);
1876 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor) / 2;
1877 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1878 expectedA.emplace_back(unitA, multiplicandFragType);
1879 expectedB.emplace_back(unitB, multiplicandFragType);
1881 if (resultPtxType() != accumPtxType())
1882 return emitOpError(
"ctype does not match dtype");
1886 if (mmaShape[0] == 8) {
1887 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1888 expectedA.emplace_back(2, f16x2Ty);
1889 expectedB.emplace_back(2, f16x2Ty);
1890 expectedResult.push_back(f16x2x4StructTy);
1891 expectedResult.push_back(f32x8StructTy);
1892 expectedC.emplace_back(4, f16x2Ty);
1893 expectedC.emplace_back(8, f32Ty);
1894 allowedShapes.push_back({8, 8, 4});
1896 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1897 Type f64Ty = Float64Type::get(context);
1898 expectedA.emplace_back(1, f64Ty);
1899 expectedB.emplace_back(1, f64Ty);
1900 expectedC.emplace_back(2, f64Ty);
1901 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1903 allowedShapes.push_back({8, 8, 4});
1906 expectedA.push_back({i32Ty});
1907 expectedB.push_back({i32Ty});
1908 expectedC.push_back({i32Ty, i32Ty});
1909 expectedResult.push_back(s32x2StructTy);
1911 allowedShapes.push_back({8, 8, 32});
1913 allowedShapes.push_back({8, 8, 16});
1917 std::string errorMessage;
1918 llvm::raw_string_ostream errorStream(errorMessage);
1921 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1922 !llvm::is_contained(allowedShapes, mmaShape)) {
1923 errorStream <<
"unimplemented variant for MMA shape <";
1924 llvm::interleaveComma(mmaShape, errorStream);
1926 return emitOpError(errorMessage);
1930 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1931 for (
const auto &iter : llvm::enumerate(
1932 std::array<AllowedTypes, 3>{std::move(expectedA),
1933 std::move(expectedB),
1934 std::move(expectedC)})) {
1935 auto spec = this->getODSOperandIndexAndLength(iter.index());
1937 operand_type_begin() + spec.first +
1939 bool match = llvm::is_contained(iter.value(), operandTySeg);
1942 errorStream <<
"Could not match types for the "
1943 << operandNames[iter.index()]
1944 <<
" operands; expected one of ";
1945 for (
const auto &x : iter.value()) {
1946 errorStream << x.size() <<
"x" << x[0] <<
" ";
1948 errorStream <<
"but got ";
1949 llvm::interleaveComma(operandTySeg, errorStream);
1950 return emitOpError(errorMessage);
1955 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1956 return expectedResultType == getResult().getType();
1959 <<
"Could not match allowed types for the result; expected one of ";
1960 llvm::interleaveComma(expectedResult, errorStream);
1961 errorStream <<
" but got " << getResult().getType();
1962 return emitOpError(errorMessage);
1969 if (!getIntOverflowBehavior())
1970 return emitOpError(
"op requires " +
1971 getIntOverflowBehaviorAttrName().strref() +
1976 if (!getSparseMetadata().
getType().isInteger(32)) {
1977 return emitOpError() <<
"sparse metadata must be i32 type";
1981 if (!getSparsitySelector().
getType().isInteger(32)) {
1982 return emitOpError() <<
"sparsity selector must be i32 type";
1994struct MMAOperandFragment {
1995 StringRef operandName;
1996 StringRef ptxTypeAttr;
1997 SmallVector<Value, 4> regs;
1998 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1999 : operandName(name), ptxTypeAttr(ptxTypeName) {}
2006 p <<
" " << name <<
"[";
2025template <
typename Op>
2030 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
2031 auto &frag = frags[fragIdx];
2032 auto varOperandSpec = op.getODSOperandIndexAndLength(fragIdx);
2033 for (
auto operandIdx = varOperandSpec.first;
2034 operandIdx < varOperandSpec.first + varOperandSpec.second;
2036 frag.regs.push_back(op.getOperand(operandIdx));
2037 if (fragIdx == 0 && operandIdx == varOperandSpec.first) {
2038 regTypes.push_back(op.getOperand(operandIdx).getType());
2042 regTypes.push_back(frag.regs[0].getType());
2044 std::optional<MMATypes> inferredType =
2045 MmaOp::inferOperandMMAType(regTypes.back(),
2048 ignoreAttrNames.push_back(frag.ptxTypeAttr);
2059 auto typeParser = [&]() {
2063 operandTypes.push_back(ty);
2069 if (operandTypes.size() != 3)
2071 "expected exactly 3 types");
2080 if (!attrs.
get(
"multiplicandAPtxType")) {
2081 if (
auto inferredType =
2082 MmaOp::inferOperandMMAType(operandTypes[0],
false)) {
2083 attrs.
set(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *inferredType));
2086 if (!attrs.
get(
"multiplicandBPtxType")) {
2087 if (
auto inferredType =
2088 MmaOp::inferOperandMMAType(operandTypes[1],
false)) {
2089 attrs.
set(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *inferredType));
2095template <
typename OpType>
2098 ScaleVecSize scaleVecSize,
2099 BlockScaleFormat blockScaleFormat,
2100 MMABlockScaleKind kind) {
2102 auto &properties =
result.getOrAddProperties<
typename OpType::Properties>();
2103 properties.setShape(
2105 properties.setScaleVecSize(ScaleVecSizeAttr::get(ctx, scaleVecSize));
2106 properties.setBlockScaleFormat(
2107 BlockScaleFormatAttr::get(ctx, blockScaleFormat));
2108 properties.setKind(MMABlockScaleKindAttr::get(ctx, kind));
2115 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
2116 if (multiplicandPtxTypes) {
2117 result.addAttribute(
"multiplicandAPtxType",
2118 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
2119 result.addAttribute(
"multiplicandBPtxType",
2120 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
2122 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
2123 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
2124 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
2125 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
2130template <
typename OpTy>
2132 return *MmaOp::inferOperandMMAType(
2133 cast<LLVM::LLVMStructType>(op.getRes().getType()).getBody()[0],
2143 std::array<MMAOperandFragment, 3> frags{
2144 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2145 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2146 MMAOperandFragment(
"C",
"")};
2148 mlir::NVVM::MmaBlockScaleOp::getOperandSegmentSizeAttr()};
2153 for (
const auto &frag : frags)
2158 {getScaleAData(), getByteIdA(), getThreadIdA()});
2160 {getScaleBData(), getByteIdB(), getThreadIdB()});
2162 bool isFirstProperty =
true;
2164 if (getMultiplicandAPtxTypeAttr() &&
2165 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2167 getMultiplicandAPtxTypeAttr());
2168 if (getMultiplicandBPtxTypeAttr() &&
2169 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2171 getMultiplicandBPtxTypeAttr());
2173 getScaleVecSizeAttr());
2175 getBlockScaleFormatAttr());
2180 getMultiplicandBPtxTypeAttrName(),
2181 getScaleVecSizeAttrName(),
2182 getBlockScaleFormatAttrName(), getKindAttrName()});
2189 frags[1].regs[0].getType(),
2190 frags[2].regs[0].getType()},
2196ParseResult MmaBlockScaleOp::parse(
OpAsmParser &parser,
2198 struct LocalOperandFragment {
2199 std::optional<MMATypes> elemtype;
2200 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2204 std::array<LocalOperandFragment, 3> frags;
2220 {
"shape",
"multiplicand_a_ptx_type",
2221 "multiplicand_b_ptx_type",
"scale_vec_size",
2222 "block_scale_format",
"kind"},
2223 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2237 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2238 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2241 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2251 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2261 result.addAttributes(namedAttributes);
2265 result.addTypes(resultTypes);
2266 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2268 static_cast<int32_t>(frags[0].regs.size()),
2269 static_cast<int32_t>(frags[1].regs.size()),
2270 static_cast<int32_t>(frags[2].regs.size()),
2281void MmaBlockScaleOp::build(
2286 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2287 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2288 MMABlockScaleKind kind) {
2289 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2292 blockScaleFormat, kind);
2294 result.addOperands(operandA);
2295 result.addOperands(operandB);
2296 result.addOperands(operandC);
2298 {scaleAData, byteIdA, threadIdA, scaleBData, byteIdB, threadIdB});
2301 multiplicandPtxTypes);
2303 result.addTypes(resultType);
2304 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
2306 static_cast<int32_t>(operandA.size()),
2307 static_cast<int32_t>(operandB.size()),
2308 static_cast<int32_t>(operandC.size()),
2320 auto curOp = cast<NVVM::MmaBlockScaleOp>(op);
2324 for (
Value operand : curOp.getOperandA())
2326 for (
Value operand : curOp.getOperandB())
2328 for (
Value operand : curOp.getOperandC())
2332 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2333 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2334 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2335 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2336 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2337 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2339 unsigned intId = MmaBlockScaleOp::getIntrinsicID(
2340 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2341 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2343 curOp.getBlockScaleFormat(), curOp.getKind());
2345 return {intId, args};
2348LogicalResult MmaBlockScaleOp::verify() {
2354 if (m == 16 && n == 8 && k == 64) {
2355 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2356 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2358 "unsupported MMATypes attribute for mma.m16n8k64.(mxf4nvf4|mxf4)");
2359 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2360 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2362 "unsupported ScaleVecSize attribute for mma.m16n8k64.mxf4");
2363 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2365 "unsupported BlockScaleFormat attribute for mma.m16n8k64.mxf4");
2366 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2367 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2368 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2369 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2370 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2371 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2372 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2373 "attributes for mma.m16n8k64.mxf4nvf4");
2375 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k64");
2377 }
else if (m == 16 && n == 8 && k == 32) {
2378 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2379 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2380 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2382 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2383 "attributes for mma.m16n8k32");
2385 result = emitOpError(
"unsupported Geom for mma with block scaling");
2396 std::array<MMAOperandFragment, 3> frags{
2397 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2398 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2399 MMAOperandFragment(
"C",
"")};
2401 mlir::NVVM::MmaSpBlockScaleOp::getOperandSegmentSizeAttr()};
2406 for (
const auto &frag : frags)
2415 {getScaleAData(), getByteIdA(), getThreadIdA()});
2417 {getScaleBData(), getByteIdB(), getThreadIdB()});
2419 bool isFirstProperty =
true;
2421 if (getMultiplicandAPtxTypeAttr() &&
2422 !llvm::is_contained(ignoreAttrNames, getMultiplicandAPtxTypeAttrName()))
2424 getMultiplicandAPtxTypeAttr());
2425 if (getMultiplicandBPtxTypeAttr() &&
2426 !llvm::is_contained(ignoreAttrNames, getMultiplicandBPtxTypeAttrName()))
2428 getMultiplicandBPtxTypeAttr());
2431 getScaleVecSizeAttr());
2433 getBlockScaleFormatAttr());
2438 getMultiplicandBPtxTypeAttrName(),
2439 getOrderedMetadataAttrName(),
2440 getScaleVecSizeAttrName(),
2441 getBlockScaleFormatAttrName(), getKindAttrName()});
2448 frags[1].regs[0].getType(),
2449 frags[2].regs[0].getType()},
2455ParseResult MmaSpBlockScaleOp::parse(
OpAsmParser &parser,
2457 struct LocalOperandFragment {
2458 std::optional<MMATypes> elemtype;
2459 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2463 std::array<LocalOperandFragment, 3> frags;
2486 {
"shape",
"multiplicand_a_ptx_type",
2487 "multiplicand_b_ptx_type",
"ordered_metadata",
2488 "scale_vec_size",
"block_scale_format",
"kind"},
2489 {
"shape",
"scaleVecSize",
"blockScaleFormat",
"kind"}))
2503 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2504 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2507 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2516 .resolveOperands(metadataOperands, i32Type, parser.
getNameLoc(),
2529 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2539 result.addAttributes(namedAttributes);
2544 if (!
result.attributes.get(
"orderedMetadata"))
2547 result.addTypes(resultTypes);
2548 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2550 static_cast<int32_t>(frags[0].regs.size()),
2551 static_cast<int32_t>(frags[1].regs.size()),
2552 static_cast<int32_t>(frags[2].regs.size()),
2565void MmaSpBlockScaleOp::build(
2571 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2572 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2573 MMABlockScaleKind kind) {
2574 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2577 builder,
result,
shape, scaleVecSize, blockScaleFormat, kind);
2580 result.addOperands(operandA);
2581 result.addOperands(operandB);
2582 result.addOperands(operandC);
2583 result.addOperands({sparseMetadata, sparsitySelector, scaleAData, byteIdA,
2584 threadIdA, scaleBData, byteIdB, threadIdB});
2587 multiplicandPtxTypes);
2589 result.addTypes(resultType);
2590 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2592 static_cast<int32_t>(operandA.size()),
2593 static_cast<int32_t>(operandB.size()),
2594 static_cast<int32_t>(operandC.size()),
2608 auto curOp = cast<NVVM::MmaSpBlockScaleOp>(op);
2612 for (
Value operand : curOp.getOperandA())
2614 for (
Value operand : curOp.getOperandB())
2616 for (
Value operand : curOp.getOperandC())
2620 args.push_back(mt.
lookupValue(curOp.getSparseMetadata()));
2621 args.push_back(mt.
lookupValue(curOp.getSparsitySelector()));
2624 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2625 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2626 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2627 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2628 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2629 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2631 unsigned intId = MmaSpBlockScaleOp::getIntrinsicID(
2632 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2633 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2635 curOp.getBlockScaleFormat(), curOp.getKind());
2637 return {intId, args};
2640LogicalResult MmaSpBlockScaleOp::verify() {
2642 if (!getOrderedMetadata()) {
2643 return emitOpError(
"'orderedMetadata' attribute is mandatory");
2651 if (m == 16 && n == 8 && k == 128) {
2652 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2653 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2655 "unsupported MMATypes attribute for mma.m16n8k128.(mxf4nvf4|mxf4)");
2656 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2657 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2659 "unsupported ScaleVecSize attribute for mma.m16n8k128.mxf4");
2660 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2662 "unsupported BlockScaleFormat attribute for mma.m16n8k128.mxf4");
2663 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2664 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2665 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2666 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2667 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2668 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2669 result = emitOpError(
"unsupported ScaleVecSize and BlockScaleFormat "
2670 "attributes for mma.m16n8k128.mxf4nvf4");
2672 result = emitOpError(
"unsupported Kind attribute for mma.m16n8k128");
2674 }
else if (m == 16 && n == 8 && k == 64) {
2675 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2676 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2677 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2679 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2680 "attributes for mma.m16n8k64");
2682 result = emitOpError(
"unsupported Geom for sparse mma with block scaling");
2687LogicalResult ShflOp::verify() {
2688 auto returnStructType = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
2690 auto verifyTypeError = [&](Twine desc,
Type expectedType,
2691 Type actualType) -> LogicalResult {
2692 return emitOpError(
"expected " + desc +
" to be of type ")
2693 << expectedType <<
" but got " << actualType <<
" instead";
2696 if (returnStructType) {
2697 if (!getReturnValueAndIsValid())
2698 return emitOpError(
"\"return_value_and_is_valid\" attribute must be "
2699 "specified when the return type is a struct type");
2701 if (returnStructType.getBody().size() != 2)
2702 return emitOpError(
"expected return type to be a two-element struct");
2705 auto resultType = returnStruct[0];
2706 if (resultType != getVal().
getType())
2707 return verifyTypeError(
"first element in the returned struct",
2708 getVal().
getType(), resultType);
2710 auto predicateType = returnStruct[1];
2711 if (!predicateType.isInteger(1))
2712 return verifyTypeError(
"second element in the returned struct",
2716 if (getReturnValueAndIsValid())
2717 return emitOpError(
"expected return type to be a two-element struct");
2720 return verifyTypeError(
"return type", getVal().
getType(),
getType());
2726ShflOp::inferReturnTypes(
MLIRContext *context, std::optional<Location> location,
2727 ShflOp::Adaptor adaptor,
2729 Type valType = adaptor.getVal().getType();
2730 if (adaptor.getReturnValueAndIsValid())
2731 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2732 context, {valType, IntegerType::get(context, 1)}));
2734 inferredReturnTypes.push_back(valType);
2739 NVVM::MMAFrag frag,
int nRow,
2742 unsigned numberElements = 0;
2745 Type f16x2 = VectorType::get(2, builder.getF16Type());
2746 if (type == NVVM::MMATypes::f16) {
2747 elementType = f16x2;
2748 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2752 }
else if (type == NVVM::MMATypes::f32) {
2753 elementType = builder.getF32Type();
2755 }
else if (type == NVVM::MMATypes::f64) {
2756 elementType = builder.getF64Type();
2757 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2761 }
else if (type == NVVM::MMATypes::tf32) {
2762 elementType = builder.getI32Type();
2764 }
else if (type == NVVM::MMATypes::s8 || type == NVVM::MMATypes::u8) {
2765 elementType = builder.getI32Type();
2766 int parallelSize = 0;
2767 if (frag == NVVM::MMAFrag::a)
2768 parallelSize = nRow;
2769 if (frag == NVVM::MMAFrag::b)
2770 parallelSize = nCol;
2773 if (parallelSize == 16)
2776 else if (parallelSize == 8)
2778 else if (parallelSize == 32)
2780 }
else if (type == NVVM::MMATypes::s32) {
2781 elementType = builder.getI32Type();
2784 assert(numberElements != 0 && elementType !=
nullptr);
2785 return std::make_pair(elementType, numberElements);
2788static std::pair<mlir::Type, unsigned>
2792 if (frag == NVVM::MMAFrag::a) {
2795 }
else if (frag == NVVM::MMAFrag::b) {
2802 assert(nRow && nCol);
2806LogicalResult NVVM::WMMALoadOp::verify() {
2807 unsigned addressSpace =
2808 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2809 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2810 addressSpace != NVVMMemorySpace::Shared)
2811 return emitOpError(
"expected source pointer in memory "
2814 if (NVVM::WMMALoadOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2815 getEltype(), getFrag()) == 0)
2816 return emitOpError() <<
"invalid attribute combination";
2821 if (typeInfo.first == f64Ty && typeInfo.second == 1) {
2823 return emitOpError(
"expected destination type to be f64");
2827 Type dstType = LLVM::LLVMStructType::getLiteral(
2830 return emitOpError(
"expected destination type is a structure of ")
2831 << typeInfo.second <<
" elements of type " << typeInfo.first;
2835LogicalResult NVVM::WMMAStoreOp::verify() {
2836 unsigned addressSpace =
2837 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2838 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2839 addressSpace != NVVMMemorySpace::Shared)
2840 return emitOpError(
"expected operands to be a source pointer in memory "
2843 if (NVVM::WMMAStoreOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2845 return emitOpError() <<
"invalid attribute combination";
2848 if (getArgs().size() != typeInfo.second)
2849 return emitOpError() <<
"expected " << typeInfo.second <<
" data operands";
2850 if (llvm::any_of(getArgs(), [&typeInfo](
Value operands) {
2851 return operands.
getType() != typeInfo.first;
2853 return emitOpError() <<
"expected data operands of type " << typeInfo.first;
2857LogicalResult NVVM::WMMAMmaOp::verify() {
2858 if (NVVM::WMMAMmaOp::getIntrinsicID(
getM(),
getN(), getK(), getLayoutA(),
2859 getLayoutB(), getEltypeA(),
2861 return emitOpError() <<
"invalid attribute combination";
2869 arguments.append(typeInfoA.second, typeInfoA.first);
2870 arguments.append(typeInfoB.second, typeInfoB.first);
2871 arguments.append(typeInfoC.second, typeInfoC.first);
2872 unsigned numArgs = arguments.size();
2873 if (getArgs().size() != numArgs)
2874 return emitOpError() <<
"expected " << numArgs <<
" arguments";
2875 for (
unsigned i = 0; i < numArgs; i++) {
2876 if (getArgs()[i].
getType() != arguments[i])
2877 return emitOpError() <<
"expected argument " << i <<
" to be of type "
2880 Type dstType = LLVM::LLVMStructType::getLiteral(
2883 return emitOpError(
"expected destination type is a structure of ")
2884 << typeInfoC.second <<
" elements of type " << typeInfoC.first;
2888LogicalResult NVVM::LdMatrixOp::verify() {
2890 if (m == 8 && n == 8) {
2891 if (num != 1 && num != 2 && num != 4) {
2892 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x8 "
2895 if (getEltType() != LdStMatrixEltType::B16) {
2896 return emitOpError(
"expected element type to be b16 for 8x8 matrix");
2898 }
else if (m == 8 && n == 16) {
2899 if (num != 1 && num != 2 && num != 4) {
2900 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x16 "
2903 if (getLayout() != MMALayout::row) {
2904 return emitOpError(
"expected layout to be row for 8x16 matrix");
2906 if (getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2907 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2908 return emitOpError(
"expected element type to be b8x16.b4x16_p64 or "
2909 "b8x16.b6x16_p32 for 8x16 matrix");
2911 }
else if (m == 16 && n == 16) {
2912 if (num != 1 && num != 2) {
2913 return emitOpError(
"expected num attribute to be 1 or 2 for 16x16 "
2916 if (getLayout() != MMALayout::col) {
2917 return emitOpError(
"expected layout to be col for 16x16 matrix");
2919 if (getEltType() != LdStMatrixEltType::B8 &&
2920 getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2921 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2922 return emitOpError(
"expected element type to be b8, b8x16.b4x16_p64 or "
2923 "b8x16.b6x16_p32 for 16x16 matrix");
2926 return emitOpError(
"expected shape to be 8x8, 8x16 or 16x16");
2930 uint32_t numElements = (m == 16 && n == 16 ? num * 2 : num);
2931 if (numElements == 1 &&
getType() != i32)
2932 return emitOpError(
"expected destination type is i32");
2933 if (numElements == 2 || numElements == 4) {
2934 Type dstType = LLVM::LLVMStructType::getLiteral(
2937 return emitOpError(
"expected destination type is a structure of ")
2938 << numElements <<
" elements of type i32";
2944LogicalResult LdMatrixOp::inferReturnTypes(
2945 MLIRContext *context, std::optional<Location> location,
2947 uint32_t num = adaptor.getNum();
2948 uint32_t m = adaptor.getShape().getM();
2949 uint32_t n = adaptor.getShape().getN();
2950 uint32_t numElements = (m == 16 && n == 16) ? num * 2 : num;
2952 Type i32 = IntegerType::get(context, 32);
2953 if (numElements == 1)
2954 inferredReturnTypes.push_back(i32);
2956 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2961LogicalResult NVVM::StMatrixOp::verify() {
2962 int numMatrix = getSources().size();
2963 if (numMatrix != 1 && numMatrix != 2 && numMatrix != 4)
2964 return emitOpError(
"expected num attribute to be 1, 2 or 4");
2967 if (m == 8 && n == 8) {
2968 if (getEltType() != NVVM::LdStMatrixEltType::B16) {
2969 return emitOpError(
"expected element type to be B16 for 8x8 matrix");
2971 }
else if (m == 16 && n == 8) {
2972 if (getEltType() != NVVM::LdStMatrixEltType::B8) {
2973 return emitOpError(
"expected element type to be B8 for 16x8 matrix");
2975 if (getLayout() != NVVM::MMALayout::col) {
2976 return emitOpError(
"expected layout to be col for 16x8 matrix");
2979 return emitOpError(
"expected shape to be 8x8 or 16x8");
2985LogicalResult NVVM::MovMatrixOp::verify() {
2987 if (m != 8 || n != 8)
2988 return emitOpError(
"expected shape to be 8x8");
2989 if (getLayout() != NVVM::MMALayout::col)
2990 return emitOpError(
"expected layout to be col");
2991 if (getEltType() != NVVM::LdStMatrixEltType::B16)
2992 return emitOpError(
"expected element type to be b16");
2997 if (typeA == NVVM::WGMMATypes::tf32)
2999 if (typeA == NVVM::WGMMATypes::f16 || typeA == NVVM::WGMMATypes::bf16)
3001 if (typeA == NVVM::WGMMATypes::s8 || typeA == NVVM::WGMMATypes::u8)
3003 if (typeA == NVVM::WGMMATypes::e4m3 || typeA == NVVM::WGMMATypes::e5m2)
3005 if (typeA == NVVM::WGMMATypes::b1)
3011 NVVM::WGMMATypes typeA,
3012 NVVM::WGMMATypes typeB) {
3014 case NVVM::WGMMATypes::f16:
3015 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3016 typeB == NVVM::WGMMATypes::f16)
3019 case NVVM::WGMMATypes::tf32:
3020 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::tf32)
3023 case NVVM::WGMMATypes::u8:
3024 case NVVM::WGMMATypes::s8:
3025 if (typeD == NVVM::WGMMATypes::s32 &&
3026 (typeB == NVVM::WGMMATypes::u8 || typeB == NVVM::WGMMATypes::s8))
3029 case NVVM::WGMMATypes::b1:
3030 if (typeD == NVVM::WGMMATypes::s32 && typeB == NVVM::WGMMATypes::b1)
3033 case NVVM::WGMMATypes::bf16:
3034 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::bf16)
3037 case NVVM::WGMMATypes::e4m3:
3038 case NVVM::WGMMATypes::e5m2:
3039 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
3040 (typeB == NVVM::WGMMATypes::e5m2 || typeB == NVVM::WGMMATypes::e4m3))
3043 case WGMMATypes::f32:
3044 case WGMMATypes::s32:
3045 llvm_unreachable(
"unsupported input types");
3053 72, 80, 88, 96, 104, 112, 120, 128,
3054 136, 144, 152, 160, 168, 176, 184, 192,
3055 200, 208, 216, 224, 232, 240, 248, 256};
3057 80, 96, 112, 128, 144, 160,
3058 176, 192, 208, 224, 240, 256};
3060 case WGMMATypes::f16:
3061 case WGMMATypes::tf32:
3062 case WGMMATypes::bf16:
3063 case WGMMATypes::e4m3:
3064 case WGMMATypes::e5m2:
3065 if (llvm::is_contained(allowedN, sizeN))
3068 case WGMMATypes::u8:
3069 case WGMMATypes::s8:
3070 case WGMMATypes::b1:
3071 if (llvm::is_contained(allowedNshort, sizeN))
3074 case WGMMATypes::f32:
3075 case WGMMATypes::s32:
3076 llvm_unreachable(
"unsupported input types");
3082LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
3083 Value outValue = getResults();
3084 auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.
getType());
3086 return emitOpError() <<
"expected results to be struct";
3087 int outputSize = stype.getBody().size();
3088 WGMMATypes typeD = getTypeD();
3089 WGMMATypes typeA = getTypeA();
3090 WGMMATypes typeB = getTypeB();
3092 for (
Type t : stype.getBody()) {
3093 if (t != stype.getBody().front())
3094 return emitOpError()
3095 <<
"all elements in struct must be same type but there is " << t;
3098 if (typeD != WGMMATypes::f32 && typeD != WGMMATypes::f16 &&
3099 typeD != WGMMATypes::s32) {
3100 return emitOpError() <<
"does not support the given output type " << typeD;
3102 if (typeD == WGMMATypes::s32 &&
3103 (getScaleA() == WGMMAScaleIn::neg || getScaleB() == WGMMAScaleIn::neg)) {
3104 return emitOpError() <<
"has s32 output, scaleA and scaleB cannot be neg";
3108 return emitOpError() << typeD <<
" += " << typeA <<
" * " << typeB
3109 <<
", it is not supported.";
3114 return emitOpError() <<
"shape 'm' must be 64";
3119 return emitOpError() <<
"shape 'k' must be " << allowedK.value()
3120 <<
" for input type " << typeA;
3124 return emitOpError() <<
"has input type " << typeA <<
" n is set to "
3125 <<
getShape().getN() <<
", it is not supported.";
3132 if ((typeA != WGMMATypes::f16 && typeA != WGMMATypes::bf16) &&
3133 (getLayoutA() == mlir::NVVM::MMALayout::col ||
3134 getLayoutB() == mlir::NVVM::MMALayout::row)) {
3135 return emitOpError()
3136 <<
"given layouts layout_a = " << getLayoutA()
3137 <<
" and layout_b = " << getLayoutB() <<
" for input types " << typeA
3139 <<
" requires transpose. However, this is only supported for: "
3140 << MMATypes::f16 <<
" and " << MMATypes::bf16;
3144 int expectedOutput = 0;
3145 if (typeD == WGMMATypes::f32 || typeD == WGMMATypes::s32)
3146 expectedOutput =
getShape().getN() / 2;
3147 if (typeD == WGMMATypes::f16)
3148 expectedOutput =
getShape().getN() / 4;
3149 if (outputSize != expectedOutput) {
3150 return emitOpError() <<
"results " << expectedOutput
3151 <<
", however output struct has " << outputSize
3155 if (typeD != WGMMATypes::s32 &&
3156 getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3157 NVVM::MMAIntOverflow::satfinite) {
3158 return emitOpError()
3159 <<
" `satfinite` can be only used with s32 accumulator, however "
3160 "the current accumulator is "
3167std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
3170 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3172 StringRef outputTypeName = stringifyWGMMATypes(getTypeD());
3174 int expectedOutputRegisters = 0;
3175 if (getTypeD() == WGMMATypes::f16)
3176 expectedOutputRegisters =
getShape().getN() / 4;
3178 expectedOutputRegisters =
getShape().getN() / 2;
3181 llvm::raw_string_ostream ss(ptx);
3186 << ((expectedOutputRegisters * 2) + 2)
3188 "wgmma.mma_async.sync.aligned.m"
3189 << m <<
"n" << n <<
"k" << k <<
"." << outputTypeName <<
"." << getTypeA()
3190 <<
"." << getTypeB();
3191 if (getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
3192 NVVM::MMAIntOverflow::satfinite)
3196 for (; regCnt < expectedOutputRegisters; ++regCnt) {
3197 ss <<
"$" << regCnt;
3198 if (regCnt != expectedOutputRegisters - 1)
3204 regCnt = (regCnt * 2);
3205 ss <<
" $" << (regCnt) <<
"," <<
" $" << (regCnt + 1) <<
"," <<
" p";
3206 if (getTypeD() != WGMMATypes::s32) {
3207 ss <<
", $" << (regCnt + 3) <<
", $" << (regCnt + 4);
3211 ss <<
", $" << (regCnt + 5) <<
", $" << (regCnt + 6);
3218bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
3222 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
3229 asmValues.push_back({makeConstantI32(rewriter,
static_cast<int>(getScaleD())),
3231 if (getTypeD() != WGMMATypes::s32) {
3232 asmValues.push_back(
3233 {makeConstantI32(rewriter,
3234 getScaleA() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3236 asmValues.push_back(
3237 {makeConstantI32(rewriter,
3238 getScaleB() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
3242 asmValues.push_back(
3243 {makeConstantI32(rewriter,
static_cast<int>(getLayoutA())),
3245 asmValues.push_back(
3246 {makeConstantI32(rewriter, 1 -
static_cast<int>(getLayoutB())),
3252LogicalResult NVVM::FenceProxyOp::verify() {
3253 if (getKind() == NVVM::ProxyKind::async_shared && !getSpace().has_value()) {
3254 return emitOpError() <<
"async_shared fence requires space attribute";
3256 if (getKind() != NVVM::ProxyKind::async_shared && getSpace().has_value()) {
3257 return emitOpError() <<
"only async_shared fence can have space attribute";
3262LogicalResult NVVM::FenceProxyAcquireOp::verify() {
3263 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3264 return emitOpError(
"uni-directional proxies only support generic for "
3265 "from_proxy attribute");
3267 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3268 return emitOpError(
"uni-directional proxies only support tensormap "
3269 "for to_proxy attribute");
3273LogicalResult NVVM::FenceProxyReleaseOp::verify() {
3274 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3275 return emitOpError(
"uni-directional proxies only support generic for "
3276 "from_proxy attribute");
3278 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
3279 return emitOpError(
"uni-directional proxies only support tensormap "
3280 "for to_proxy attribute");
3284LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
3285 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
3286 return emitOpError(
"only generic is support for from_proxy attribute");
3288 if (getToProxy() != NVVM::ProxyKind::async)
3289 return emitOpError(
"only async is supported for to_proxy attribute");
3293LogicalResult NVVM::SetMaxRegisterOp::verify() {
3294 if (getRegCount() % 8)
3295 return emitOpError(
"new register size must be multiple of 8");
3296 if (getRegCount() < 24 || getRegCount() > 256)
3297 return emitOpError(
"new register size must be in between 24 to 256");
3301LogicalResult NVVM::Tcgen05CpOp::verify() {
3302 auto mc = getMulticast();
3304 using SH = Tcgen05CpShape;
3305 using MC = Tcgen05CpMulticast;
3307 case SH::SHAPE_128x256b:
3308 case SH::SHAPE_128x128b:
3309 case SH::SHAPE_4x256b:
3311 return emitError(
"Invalid multicast type for tcgen05.cp Op");
3313 case SH::SHAPE_64x128b:
3314 if (mc != MC::WARPX2_01_23 && mc != MC::WARPX2_02_13)
3315 return emitError(
"Shape 64x128b requires multicast warpx2_01_23 or "
3316 "warpx2_02_13 for tcgen05.cp Op");
3318 case SH::SHAPE_32x128b:
3319 if (mc != MC::WARPX4)
3321 "Shape 32x128b requires multicast warpx4 for tcgen05.cp Op");
3327LogicalResult NVVM::MatchSyncOp::verify() {
3328 if (getKind() == NVVM::MatchSyncKind::all) {
3329 auto type = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
3330 if (!type || type.getBody().size() != 2 ||
3331 !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
3332 return emitOpError(
"match.sync 'all' returns a two element struct with "
3333 "first element as i32 and second element as i1");
3336 if (!
getType().isInteger(32)) {
3337 return emitOpError(
"match.sync 'any' returns an i32");
3343LogicalResult MatchSyncOp::inferReturnTypes(
3344 MLIRContext *context, std::optional<Location> location,
3346 if (adaptor.getKind() == NVVM::MatchSyncKind::all)
3347 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
3349 {IntegerType::get(context, 32), IntegerType::get(context, 1)}));
3351 inferredReturnTypes.push_back(IntegerType::get(context, 32));
3355LogicalResult NVVM::VoteSyncOp::verify() {
3356 if (getKind() == NVVM::VoteSyncKind::ballot) {
3357 if (!
getType().isInteger(32)) {
3358 return emitOpError(
"vote.sync 'ballot' returns an i32");
3361 if (!
getType().isInteger(1)) {
3362 return emitOpError(
"vote.sync 'any', 'all' and 'uni' returns an i1");
3368LogicalResult VoteSyncOp::inferReturnTypes(
3369 MLIRContext *context, std::optional<Location> location,
3371 unsigned width = adaptor.getKind() == NVVM::VoteSyncKind::ballot ? 32 : 1;
3372 inferredReturnTypes.push_back(IntegerType::get(context, width));
3376LogicalResult NVVM::PrefetchOp::verify() {
3377 using MemSpace = NVVM::NVVMMemorySpace;
3378 using CacheLevel = NVVM::PrefetchCacheLevel;
3380 unsigned addressSpace =
3381 llvm::cast<LLVM::LLVMPointerType>(getAddr().
getType()).getAddressSpace();
3382 std::optional<NVVM::CacheEvictionPriority> evictPriority = getEvictPriority();
3383 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = getCacheLevel();
3385 if (getTensormap() && cacheLevel)
3386 return emitOpError(
"cannot specify both tensormap and cache level");
3388 if (getTensormap()) {
3389 if (addressSpace != MemSpace::Generic &&
3390 addressSpace != MemSpace::Constant) {
3392 "prefetch tensormap requires a generic or constant pointer");
3395 if (evictPriority) {
3397 "prefetch tensormap does not support eviction priority");
3400 if (getInParamSpace() && addressSpace != MemSpace::Generic) {
3402 "in_param_space can only be specified for a generic pointer");
3405 }
else if (cacheLevel) {
3406 if (addressSpace != MemSpace::Generic && addressSpace != MemSpace::Global &&
3407 addressSpace != MemSpace::Local) {
3408 return emitOpError(
"prefetch to cache level requires a generic, global, "
3409 "or local pointer");
3413 if (*cacheLevel != CacheLevel::L1) {
3415 "unsupported cache level, the only supported uniform "
3416 "cache level is L1");
3419 if (addressSpace != MemSpace::Generic) {
3421 "prefetch to uniform cache requires a generic pointer");
3425 if (evictPriority) {
3426 if (*cacheLevel != CacheLevel::L2)
3428 "cache eviction priority supported only for cache level L2");
3430 if (addressSpace != MemSpace::Global)
3431 return emitOpError(
"cache eviction priority requires a global pointer");
3433 if (*evictPriority != NVVM::CacheEvictionPriority::EvictNormal &&
3434 *evictPriority != NVVM::CacheEvictionPriority::EvictLast)
3436 "unsupported cache eviction priority, only evict_last and "
3437 "evict_normal are supported");
3441 return emitOpError(
"predicate supported only on prefetch tensormap");
3445 "requires specification of either cache level or tensormap");
3451LogicalResult NVVM::ClusterLaunchControlQueryCancelOp::verify() {
3452 switch (getQueryType()) {
3453 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
3455 return emitOpError(
"is_canceled query type returns an i1");
3457 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
3458 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
3459 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
3460 if (!
getType().isInteger(32)) {
3461 return emitOpError(
"get_first_cta_id_x, get_first_cta_id_y, "
3462 "get_first_cta_id_z query types return an i32");
3469LogicalResult ClusterLaunchControlQueryCancelOp::inferReturnTypes(
3470 MLIRContext *context, std::optional<Location> location,
3471 ClusterLaunchControlQueryCancelOp::Adaptor adaptor,
3474 adaptor.getQueryType() == NVVM::ClusterLaunchControlQueryType::IS_CANCELED
3477 inferredReturnTypes.push_back(IntegerType::get(context, width));
3481LogicalResult NVVM::ReduxOp::verify() {
3484 if (!reduxType.
isF32()) {
3486 return emitOpError(
"abs attribute is supported only for f32 type");
3488 return emitOpError(
"nan attribute is supported only for f32 type");
3491 NVVM::ReductionKind kind = getKind();
3493 case NVVM::ReductionKind::ADD:
3494 case NVVM::ReductionKind::AND:
3495 case NVVM::ReductionKind::OR:
3496 case NVVM::ReductionKind::XOR:
3497 case NVVM::ReductionKind::MAX:
3498 case NVVM::ReductionKind::MIN:
3499 case NVVM::ReductionKind::UMAX:
3500 case NVVM::ReductionKind::UMIN:
3502 return emitOpError(
"'")
3503 << kind <<
"' reduction kind unsupported with " << reduxType
3504 <<
" type. Only supported type is 'i32'.";
3506 case NVVM::ReductionKind::FMIN:
3507 case NVVM::ReductionKind::FMAX:
3508 if (!reduxType.isF32())
3509 return emitOpError(
"'")
3510 << kind <<
"' reduction kind unsupported with " << reduxType
3511 <<
" type. Only supported type is 'f32'.";
3518LogicalResult NVVM::TensormapReplaceOp::verify() {
3519 auto ord = getOrd();
3520 Value newVal = getNewValue();
3521 auto newValAttr = getNewValueAttr();
3522 auto fieldName = stringifyEnum(getField());
3524 if (ord && !llvm::is_contained({NVVM::TensormapField::BOX_DIM,
3525 NVVM::TensormapField::GLOBAL_DIM,
3526 NVVM::TensormapField::GLOBAL_STRIDE,
3527 NVVM::TensormapField::ELEMENT_STRIDE},
3529 return emitOpError(
"ordinal is not supported for ")
3530 << fieldName <<
" field";
3532 auto invalidNewVal = [&](llvm::Twine type) -> std::string {
3533 return llvm::Twine(
"new_value must be specified and must be an " + type +
3534 " for " + llvm::Twine(fieldName) +
" field")
3538 auto invalidNewValAttr = [&]() -> std::string {
3539 return (llvm::Twine(
3540 "new_value_attr must be specified and must be a valid ") +
3541 llvm::Twine(fieldName) +
" attribute for " + fieldName +
" field")
3545 switch (getField()) {
3546 case NVVM::TensormapField::GLOBAL_ADDRESS:
3548 return emitOpError(invalidNewVal(
"i64"));
3550 case NVVM::TensormapField::RANK:
3552 return emitOpError(invalidNewVal(
"i32"));
3554 case NVVM::TensormapField::GLOBAL_STRIDE:
3556 return emitOpError(
"ordinal is required for global_stride field");
3558 return emitOpError(invalidNewVal(
"i64"));
3560 case NVVM::TensormapField::BOX_DIM:
3561 case NVVM::TensormapField::GLOBAL_DIM:
3562 case NVVM::TensormapField::ELEMENT_STRIDE:
3564 return emitOpError(
"ordinal is required for ")
3565 << stringifyEnum(getField()) <<
" field";
3567 return emitOpError(invalidNewVal(
"i32"));
3569 case NVVM::TensormapField::ELEMTYPE:
3570 if (!(newValAttr && llvm::isa<TensormapElemtypeAttr>(*newValAttr)))
3571 return emitOpError(invalidNewValAttr());
3573 case NVVM::TensormapField::INTERLEAVE_LAYOUT:
3574 if (!(newValAttr && llvm::isa<TensormapInterleaveLayoutAttr>(*newValAttr)))
3575 return emitOpError(invalidNewValAttr());
3577 case NVVM::TensormapField::SWIZZLE_MODE:
3578 if (!(newValAttr && llvm::isa<TensormapSwizzleModeAttr>(*newValAttr)))
3579 return emitOpError(invalidNewValAttr());
3581 case NVVM::TensormapField::SWIZZLE_ATOMICITY:
3582 if (!(newValAttr && llvm::isa<TensormapSwizzleAtomicityAttr>(*newValAttr)))
3583 return emitOpError(invalidNewValAttr());
3585 case NVVM::TensormapField::FILL_MODE:
3586 if (!(newValAttr && llvm::isa<TensormapFillModeAttr>(*newValAttr)))
3587 return emitOpError(invalidNewValAttr());
3594template <
typename OpType>
3596 mlir::NVVM::FPRoundingMode rndMode = op.getRnd();
3597 mlir::NVVM::SaturationMode satMode = op.getSat();
3598 bool isFTZ = op.getFtz();
3601 mlir::Type opBaseType = isa<VectorType>(opType)
3602 ? cast<VectorType>(opType).getElementType()
3605 if (opBaseType.
isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3606 return op.emitOpError(
"FTZ and saturation are not supported for "
3607 "additions/subtractions involving f64 type");
3609 if (opBaseType.
isF16() && !(rndMode == NVVM::FPRoundingMode::RN ||
3610 rndMode == NVVM::FPRoundingMode::NONE))
3611 return op.emitOpError(
"only RN rounding mode is supported for f16 and "
3612 "vector<2xf16> additions/subtractions");
3614 if (opBaseType.
isBF16()) {
3615 if (rndMode != NVVM::FPRoundingMode::RN &&
3616 rndMode != NVVM::FPRoundingMode::NONE)
3617 return op.emitOpError(
"only RN rounding mode is supported for bf16 and "
3618 "vector<2xbf16> additions/subtractions");
3619 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3620 return op.emitOpError(
"FTZ and saturation are not supported for bf16 and "
3621 "vector<2xbf16> additions/subtractions");
3631LogicalResult NVVM::FmaOp::verify() {
3632 auto opType = getRes().getType();
3633 mlir::NVVM::FPRoundingMode rndMode = getRnd();
3634 mlir::NVVM::SaturationMode satMode = getSat();
3635 bool isFTZ = getFtz();
3636 bool isRelu = getRelu();
3637 bool hasOOB = getOob();
3639 auto getBaseFType = [](
Type type) ->
Type {
3640 if (isa<VectorType>(type))
3641 return cast<VectorType>(type).getElementType();
3645 auto opBaseType = getBaseFType(opType);
3647 if (rndMode == NVVM::FPRoundingMode::NONE)
3648 return emitOpError(
"rounding mode must be specified");
3650 if (isRelu && satMode == NVVM::SaturationMode::SAT)
3651 return emitOpError(
"relu and saturation are not supported together");
3653 if (hasOOB && (satMode == NVVM::SaturationMode::SAT || isFTZ))
3654 return emitOpError(
"oob is not supported with saturation or FTZ");
3656 if (!(opBaseType.isF16() || opBaseType.isBF16()) && (isRelu || hasOOB))
3657 return emitOpError(
"relu and oob are only supported for f16 and bf16");
3659 if (opBaseType.isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3660 return emitOpError(
"FTZ and saturation are not supported for f64 type");
3662 if (opBaseType.isF16() && rndMode != NVVM::FPRoundingMode::RN)
3664 "only RN rounding mode is supported for f16 and vector<2xf16>");
3666 if (opBaseType.isBF16()) {
3667 if (rndMode != NVVM::FPRoundingMode::RN)
3669 "only RN rounding mode is supported for bf16 and vector<2xbf16>");
3670 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3672 "FTZ and saturation are not supported for bf16 and vector<2xbf16>");
3678LogicalResult NVVM::SqrtOp::verify() {
3679 if (getRnd() == NVVM::FPRoundingMode::NONE)
3680 return emitOpError(
"rounding mode cannot be None");
3682 if (getRes().
getType().isF64() && getFtz())
3683 return emitOpError(
"FTZ is not supported for f64");
3688LogicalResult NVVM::DivFOp::verify() {
3689 bool isApprox = getApprox();
3690 bool isFull = getFull();
3691 bool isF64 = getRes().getType().isF64();
3692 bool isFtz = getFtz();
3693 NVVM::FPRoundingMode rndMode = getRnd();
3695 if (isApprox && isFull)
3696 return emitOpError(
"'approx' and 'full' are mutually exclusive");
3698 if (isApprox || isFull) {
3700 return emitOpError(
"'approx' and 'full' forms are f32-only");
3701 if (rndMode != NVVM::FPRoundingMode::NONE)
3703 "'approx' and 'full' forms do not accept a rounding mode");
3708 if (rndMode == NVVM::FPRoundingMode::NONE)
3709 return emitOpError(
"rounding mode cannot be None for the rounded divide");
3711 return emitOpError(
"FTZ is not supported for f64");
3722 unsigned sizeInBits,
3724 field = builder.CreateZExtOrBitCast(field, builder.getInt32Ty());
3726 unsigned mask = (sizeInBits < 32 ? ((1u << sizeInBits) - 1) : 0xffffffffu);
3727 if (mask != 0xffffffffu)
3728 field = builder.CreateAnd(field, builder.getInt32(mask));
3730 field = builder.CreateZExtOrBitCast(field, builder.getInt64Ty());
3731 field = builder.CreateShl(field, start);
3733 return builder.CreateOr(
result, field);
3736void Tcgen05MmaSmemDescOp::createSmemDescriptor(
Operation &op,
3738 llvm::IRBuilderBase &builder) {
3739 auto thisOp = cast<NVVM::Tcgen05MmaSmemDescOp>(op);
3740 llvm::Value *smemDesc = builder.getInt64(0);
3745 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimOffset()), 14, 16);
3747 builder, smemDesc, mt.
lookupValue(thisOp.getStrideDimOffset()), 14, 32);
3753 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimMode()), 1, 52);
3757 mt.
mapValue(thisOp.getRes()) = smemDesc;
3764std::string NVVM::MBarrierInitOp::getPtx() {
3767 std::string layout =
3768 getLayout() == 1 ? std::string(
".layout::v1") : std::string();
3770 return llvm::formatv(
"mbarrier.init{0}{1}.b64 [%0], %1;", layout, space)
3774std::string NVVM::MBarrierArriveExpectTxOp::getPtx() {
3777 ? std::string(
"mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;")
3778 : std::string(
"mbarrier.arrive.expect_tx.b64 _, [%0], %1;");
3781std::string NVVM::MBarrierTryWaitParityOp::getPtx() {
3783 llvm::StringRef space = isShared ?
".shared" :
"";
3785 return llvm::formatv(
"{\n\t"
3786 ".reg .pred P1; \n\t"
3788 "mbarrier.try_wait.parity{0}.b64 P1, [%0], %1, %2; \n\t"
3789 "@P1 bra.uni DONE; \n\t"
3790 "bra.uni LAB_WAIT; \n\t"
3807 LLVM::FNegOp::create(rewriter, loc, op.getRhs().getType(), op.getRhs());
3810 op.getRnd(), op.getSat(), op.getFtz());
3829 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count
3830 : llvm::Intrinsic::nvvm_barrier_cta_sync_count;
3832 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all
3833 : llvm::Intrinsic::nvvm_barrier_cta_sync_all;
3838static llvm::Intrinsic::ID
3841 case NVVM::BarrierReduction::AND:
3842 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_and_aligned_all
3843 : llvm::Intrinsic::nvvm_barrier_cta_red_and_all;
3844 case NVVM::BarrierReduction::OR:
3845 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_or_aligned_all
3846 : llvm::Intrinsic::nvvm_barrier_cta_red_or_all;
3847 case NVVM::BarrierReduction::POPC:
3848 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_popc_aligned_all
3849 : llvm::Intrinsic::nvvm_barrier_cta_red_popc_all;
3851 llvm_unreachable(
"unknown BarrierReduction kind");
3856 auto thisOp = cast<NVVM::BarrierOp>(op);
3857 llvm::Value *barrierId = thisOp.getBarrierId()
3859 : builder.getInt32(0);
3860 bool hasCount =
static_cast<bool>(thisOp.getNumberOfThreads());
3861 llvm::Intrinsic::ID
id =
3865 args.push_back(mt.
lookupValue(thisOp.getNumberOfThreads()));
3866 return {id, std::move(args)};
3871 auto thisOp = cast<NVVM::BarrierArriveOp>(op);
3872 llvm::Value *barrierId = thisOp.getBarrierId()
3874 : builder.getInt32(0);
3875 llvm::Value *numThreads = mt.
lookupValue(thisOp.getNumberOfThreads());
3876 llvm::Intrinsic::ID
id =
3878 ? llvm::Intrinsic::nvvm_barrier_cta_arrive_aligned_count
3879 : llvm::Intrinsic::nvvm_barrier_cta_arrive_count;
3880 return {id, {barrierId, numThreads}};
3885 auto thisOp = cast<NVVM::BarrierReductionOp>(op);
3887 thisOp.getAligned(), thisOp.getReductionOp());
3888 llvm::Value *barrierId = thisOp.getBarrierId()
3890 : builder.getInt32(0);
3893 builder.CreateICmpNE(mt.
lookupValue(thisOp.getReductionPredicate()),
3894 builder.getInt32(0))};
3895 return {id, std::move(args)};
3900 llvm::IRBuilderBase &builder) {
3901 auto thisOp = cast<NVVM::CosOp>(op);
3902 llvm::Intrinsic::ID
id = thisOp.getFtz()
3903 ? llvm::Intrinsic::nvvm_cos_approx_ftz_f
3904 : llvm::Intrinsic::nvvm_cos_approx_f;
3910 llvm::IRBuilderBase &builder) {
3911 auto thisOp = cast<NVVM::SinOp>(op);
3912 llvm::Intrinsic::ID
id = thisOp.getFtz()
3913 ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
3914 : llvm::Intrinsic::nvvm_sin_approx_f;
3920 llvm::IRBuilderBase &builder) {
3921 auto thisOp = cast<NVVM::Log2Op>(op);
3922 llvm::Intrinsic::ID
id = thisOp.getFtz()
3923 ? llvm::Intrinsic::nvvm_lg2_approx_ftz_f
3924 : llvm::Intrinsic::nvvm_lg2_approx_f;
3930 llvm::IRBuilderBase &builder) {
3931 auto thisOp = cast<NVVM::Ex2Op>(op);
3932 llvm::Intrinsic::ID
id = thisOp.getFtz()
3933 ? llvm::Intrinsic::nvvm_ex2_approx_ftz
3934 : llvm::Intrinsic::nvvm_ex2_approx;
3938LogicalResult NVVM::Ex2Op::verify() {
3939 auto vectorType = dyn_cast<VectorType>(getSrc().
getType());
3943 if (vectorType.getElementType().isF16() && getFtz())
3944 return emitOpError(
"FTZ is not supported for vector<2xf16>");
3945 if (vectorType.getElementType().isBF16() && !getFtz())
3946 return emitOpError(
"FTZ is required for vector<2xbf16>");
3952 llvm::IRBuilderBase &builder) {
3953 auto thisOp = cast<NVVM::RsqrtOp>(op);
3954 Type t = thisOp.getRes().getType();
3955 bool isFtz = thisOp.getFtz();
3957 llvm::Intrinsic::ID
id = [&] {
3959 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_f
3960 : llvm::Intrinsic::nvvm_rsqrt_approx_f;
3963 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_d
3964 : llvm::Intrinsic::nvvm_rsqrt_approx_d;
3972 llvm::IRBuilderBase &builder) {
3973 auto thisOp = cast<NVVM::SqrtOp>(op);
3974 Type t = thisOp.getRes().getType();
3975 NVVM::FPRoundingMode rndMode = thisOp.getRnd();
3976 bool isFtz = thisOp.getFtz();
3980 unsigned rndIndex =
static_cast<unsigned>(rndMode) - 1;
3982 static constexpr llvm::Intrinsic::ID f32IDs[] = {
3983 llvm::Intrinsic::nvvm_sqrt_rn_f,
3984 llvm::Intrinsic::nvvm_sqrt_rm_f,
3985 llvm::Intrinsic::nvvm_sqrt_rp_f,
3986 llvm::Intrinsic::nvvm_sqrt_rz_f,
3988 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
3989 llvm::Intrinsic::nvvm_sqrt_rn_ftz_f,
3990 llvm::Intrinsic::nvvm_sqrt_rm_ftz_f,
3991 llvm::Intrinsic::nvvm_sqrt_rp_ftz_f,
3992 llvm::Intrinsic::nvvm_sqrt_rz_ftz_f,
3994 static constexpr llvm::Intrinsic::ID f64IDs[] = {
3995 llvm::Intrinsic::nvvm_sqrt_rn_d,
3996 llvm::Intrinsic::nvvm_sqrt_rm_d,
3997 llvm::Intrinsic::nvvm_sqrt_rp_d,
3998 llvm::Intrinsic::nvvm_sqrt_rz_d,
4001 llvm::Intrinsic::ID
id =
4002 t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4010 llvm::IRBuilderBase &builder) {
4011 auto thisOp = cast<NVVM::SqrtApproxOp>(op);
4012 llvm::Intrinsic::ID
id = thisOp.getFtz()
4013 ? llvm::Intrinsic::nvvm_sqrt_approx_ftz_f
4014 : llvm::Intrinsic::nvvm_sqrt_approx_f;
4020 llvm::IRBuilderBase &builder) {
4021 auto thisOp = cast<NVVM::DivFOp>(op);
4022 bool isFtz = thisOp.getFtz();
4024 llvm::Intrinsic::ID id;
4026 if (thisOp.getApprox()) {
4027 id = isFtz ? llvm::Intrinsic::nvvm_div_approx_ftz_f
4028 : llvm::Intrinsic::nvvm_div_approx_f;
4029 }
else if (thisOp.getFull()) {
4032 id = isFtz ? llvm::Intrinsic::nvvm_div_full_ftz
4033 : llvm::Intrinsic::nvvm_div_full;
4036 unsigned rndIndex =
static_cast<unsigned>(thisOp.getRnd()) - 1;
4038 static constexpr llvm::Intrinsic::ID f32IDs[] = {
4039 llvm::Intrinsic::nvvm_div_rn_f,
4040 llvm::Intrinsic::nvvm_div_rm_f,
4041 llvm::Intrinsic::nvvm_div_rp_f,
4042 llvm::Intrinsic::nvvm_div_rz_f,
4044 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
4045 llvm::Intrinsic::nvvm_div_rn_ftz_f,
4046 llvm::Intrinsic::nvvm_div_rm_ftz_f,
4047 llvm::Intrinsic::nvvm_div_rp_ftz_f,
4048 llvm::Intrinsic::nvvm_div_rz_ftz_f,
4050 static constexpr llvm::Intrinsic::ID f64IDs[] = {
4051 llvm::Intrinsic::nvvm_div_rn_d,
4052 llvm::Intrinsic::nvvm_div_rm_d,
4053 llvm::Intrinsic::nvvm_div_rp_d,
4054 llvm::Intrinsic::nvvm_div_rz_d,
4056 Type t = thisOp.getRes().getType();
4057 id = t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
4068 auto thisOp = cast<NVVM::AsyncStoreGlobalOp>(op);
4069 mlir::NVVM::MemScopeKind scope = thisOp.getScope();
4070 bool isMmio = thisOp.getMmio();
4072 llvm::Value *addr = mt.
lookupValue(thisOp.getAddr());
4073 llvm::Value *value = mt.
lookupValue(thisOp.getValue());
4074 llvm::Value *isMultimem = builder.getInt1(thisOp.getMultimem());
4076 if (scope == MemScopeKind::SYS) {
4077 return isMmio ?
IDArgPair(llvm::Intrinsic::nvvm_st_async_mmio_sys,
4080 {addr, value, isMultimem});
4081 }
else if (scope == MemScopeKind::GPU) {
4082 return IDArgPair(llvm::Intrinsic::nvvm_st_async_gpu,
4083 {addr, value, isMultimem});
4085 llvm_unreachable(
"unsupported scope for AsyncStoreGlobalOp");
4090 llvm::IRBuilderBase &builder) {
4091 auto thisOp = cast<NVVM::PMEventOp>(op);
4095 llvm::Value *maskVal;
4096 if (
auto eventAttr = thisOp.getEventIdAttr()) {
4097 uint16_t mask =
static_cast<uint16_t
>(1u << eventAttr.getInt());
4098 maskVal = llvm::ConstantInt::get(i16Ty, mask);
4101 llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
4104 return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
4107bool MBarrierInitOp::getAsmValues(
4114 for (
auto val : getOperands())
4122 auto thisOp = cast<NVVM::MBarrierInitOp>(op);
4126 return {llvm::Intrinsic::nvvm_mbarrier_init,
4128 builder.getInt32(thisOp.getLayout())}};
4133 auto thisOp = cast<NVVM::MBarrierInvalOp>(op);
4135 llvm::Intrinsic::ID
id = isShared
4136 ? llvm::Intrinsic::nvvm_mbarrier_inval_shared
4137 : llvm::Intrinsic::nvvm_mbarrier_inval;
4144 auto thisOp = cast<NVVM::MBarrierCheckLayoutOp>(op);
4147 llvm::Intrinsic::nvvm_mbarrier_check_layout,
4148 {mt.
lookupValue(thisOp.getAddr()), builder.getInt32(thisOp.getLayout())}};
4153 auto thisOp = cast<NVVM::MBarrierExpectTxOp>(op);
4156 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4159 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4161 static constexpr llvm::Intrinsic::ID IDs[] = {
4162 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cta,
4163 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cluster,
4164 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cta,
4165 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cluster};
4170 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4172 return {IDs[
index], std::move(args)};
4177 auto thisOp = cast<NVVM::MBarrierCompleteTxOp>(op);
4180 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4183 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4185 static constexpr llvm::Intrinsic::ID IDs[] = {
4186 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cta,
4187 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cluster,
4188 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cta,
4189 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cluster};
4194 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
4196 return {IDs[
index], std::move(args)};
4201 auto thisOp = cast<NVVM::MBarrierArriveOp>(op);
4204 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4207 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4209 static constexpr llvm::Intrinsic::ID IDs[] = {
4210 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta,
4211 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cluster,
4212 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cta,
4213 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cluster};
4214 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4215 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cta,
4216 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cluster,
4217 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cta,
4219 nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cluster};
4220 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4224 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4231 bool hasCount =
static_cast<bool>(thisOp.getCount());
4233 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta))
4234 return {llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {mbar}};
4238 llvm::Value *count =
4240 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4241 return {id, {mbar, count}};
4246 auto thisOp = cast<NVVM::MBarrierArriveDropOp>(op);
4249 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4252 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4254 static constexpr llvm::Intrinsic::ID IDs[] = {
4255 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta,
4256 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cluster,
4257 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cta,
4258 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cluster};
4259 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4260 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cta,
4262 nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cluster,
4264 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cta,
4266 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cluster};
4267 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4271 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4278 bool hasCount =
static_cast<bool>(thisOp.getCount());
4280 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta))
4281 return {llvm::Intrinsic::nvvm_mbarrier_arrive_drop_shared, {mbar}};
4285 llvm::Value *count =
4287 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
4288 return {id, {mbar, count}};
4291bool MBarrierArriveExpectTxOp::getAsmValues(
4298 for (
auto val : getOperands())
4306 auto thisOp = cast<NVVM::MBarrierArriveExpectTxOp>(op);
4309 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4312 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4315 static constexpr llvm::Intrinsic::ID IDs[] = {
4316 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cta,
4317 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cluster,
4318 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cta,
4319 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cluster};
4320 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4321 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cta,
4322 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cluster,
4323 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cta,
4324 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cluster};
4326 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4329 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4330 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4335 return {id, {mbar, txcount}};
4340 auto thisOp = cast<NVVM::MBarrierArriveDropExpectTxOp>(op);
4343 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4346 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
4349 static constexpr llvm::Intrinsic::ID IDs[] = {
4350 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cta,
4351 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cluster,
4352 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cta,
4353 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cluster};
4354 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4355 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cta,
4356 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cluster,
4357 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cta,
4358 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cluster};
4360 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4363 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
4364 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4369 return {id, {mbar, txcount}};
4374 auto thisOp = cast<NVVM::MBarrierArriveNocompleteOp>(op);
4376 llvm::Intrinsic::ID
id =
4377 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared
4378 : llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete;
4382 args.push_back(mt.
lookupValue(thisOp.getCount()));
4384 return {id, std::move(args)};
4389 auto thisOp = cast<NVVM::MBarrierArriveDropNocompleteOp>(op);
4391 llvm::Intrinsic::ID
id =
4392 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete_shared
4393 : llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete;
4397 args.push_back(mt.
lookupValue(thisOp.getCount()));
4399 return {id, std::move(args)};
4404 auto thisOp = cast<NVVM::MBarrierTestWaitOp>(op);
4405 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4406 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4409 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isPhaseParity ? 1 : 0);
4412 static constexpr llvm::Intrinsic::ID IDs[] = {
4413 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cta_space_cta,
4414 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cta_space_cta,
4415 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cluster_space_cta,
4416 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cluster_space_cta};
4417 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4418 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cta_space_cta,
4419 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cta_space_cta,
4420 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cluster_space_cta,
4421 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cluster_space_cta};
4423 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4426 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4427 llvm::Value *input = mt.
lookupValue(thisOp.getStateOrPhase());
4432 return {id, {mbar, input}};
4437 auto thisOp = cast<NVVM::MBarrierTryWaitOp>(op);
4438 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
4439 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
4440 bool hasTicks =
static_cast<bool>(thisOp.getTicks());
4444 size_t index = ((hasTicks ? 1 : 0) << 2) | ((isClusterScope ? 1 : 0) << 1) |
4445 (isPhaseParity ? 1 : 0);
4448 static constexpr llvm::Intrinsic::ID IDs[] = {
4449 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cta_space_cta,
4450 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cta_space_cta,
4451 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cluster_space_cta,
4452 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cluster_space_cta,
4453 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cta_space_cta,
4454 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cta_space_cta,
4455 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cluster_space_cta,
4456 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cluster_space_cta};
4457 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4458 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cta_space_cta,
4459 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cta_space_cta,
4460 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cluster_space_cta,
4461 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cluster_space_cta,
4462 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cta_space_cta,
4463 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cta_space_cta,
4464 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cluster_space_cta,
4465 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cluster_space_cta};
4467 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4470 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4477 args.push_back(mbar);
4478 args.push_back(mt.
lookupValue(thisOp.getStateOrPhase()));
4480 args.push_back(mt.
lookupValue(thisOp.getTicks()));
4482 return {id, std::move(args)};
4487 auto thisOp = cast<NVVM::CpAsyncMBarrierArriveOp>(op);
4490 llvm::Intrinsic::ID id;
4491 if (thisOp.getNoinc()) {
4492 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc_shared
4493 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc;
4495 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_shared
4496 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive;
4504 llvm::IRBuilderBase &builder) {
4505 auto thisOp = cast<NVVM::MovMatrixOp>(op);
4506 return {llvm::Intrinsic::nvvm_movmatrix_sync_aligned_m8n8_trans_b16,
4510#define CP_ASYNC_ID_IMPL(mod, size, suffix) \
4511 llvm::Intrinsic::nvvm_cp_async_##mod##_shared_global_##size##suffix
4513#define GET_CP_ASYNC_ID(mod, size, has_cpsize) \
4514 has_cpsize ? CP_ASYNC_ID_IMPL(mod, size, _s) : CP_ASYNC_ID_IMPL(mod, size, )
4519 llvm::Intrinsic::ID id;
4521 auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
4522 bool hasCpSize =
static_cast<bool>(cpAsyncOp.getCpSize());
4523 switch (cpAsyncOp.getSize()) {
4531 id = (cpAsyncOp.getModifier() == NVVM::LoadCacheModifierKind::CG)
4536 llvm_unreachable(
"Invalid copy size in CpAsyncOp.");
4540 args.push_back(mt.
lookupValue(cpAsyncOp.getDst()));
4541 args.push_back(mt.
lookupValue(cpAsyncOp.getSrc()));
4543 args.push_back(mt.
lookupValue(cpAsyncOp.getCpSize()));
4550 auto thisOp = cast<NVVM::CpAsyncBulkPrefetchOp>(op);
4552 llvm::Intrinsic::ID
id = llvm::Intrinsic::nvvm_cp_async_bulk_prefetch_L2;
4555 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4559 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4560 llvm::Value *i64Unused =
4561 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4562 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4563 args.push_back(builder.getInt1(hasCacheHint));
4565 return {id, std::move(args)};
4570 auto thisOp = cast<NVVM::CpAsyncBulkGlobalToSharedClusterOp>(op);
4574 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4576 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4580 mlir::Value multicastMask = thisOp.getMulticastMask();
4581 const bool hasMulticastMask =
static_cast<bool>(multicastMask);
4584 llvm::Value *i16Unused = llvm::ConstantInt::get(builder.getInt16Ty(), 0);
4585 args.push_back(hasMulticastMask ? mt.
lookupValue(multicastMask)
4588 args.push_back(builder.getInt32(0));
4589 args.push_back(builder.getInt32(0));
4594 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4595 llvm::Value *i64Unused = llvm::ConstantInt::get(builder.getInt64Ty(), 0);
4596 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4600 args.push_back(builder.getInt1(hasCacheHint));
4601 args.push_back(builder.getInt1(
false));
4603 args.push_back(builder.getInt1(hasMulticastMask));
4604 args.push_back(builder.getInt1(hasCacheHint));
4608 args.push_back(builder.getInt32(0));
4610 llvm::Intrinsic::ID
id =
4612 ? llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta
4613 : llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
4615 return {id, std::move(args)};
4620 auto thisOp = cast<NVVM::CpAsyncBulkSharedCTAToGlobalOp>(op);
4622 llvm::Intrinsic::ID
id =
4623 llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global;
4626 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4627 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4631 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4632 llvm::Value *i64Unused =
4633 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4634 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4635 args.push_back(builder.getInt1(hasCacheHint));
4638 if (
mlir::Value byteMask = thisOp.getByteMask()) {
4640 id = llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global_bytemask;
4643 return {id, std::move(args)};
4646bool CpAsyncBulkTensorGlobalToSharedClusterOp::getAsmValues(
4653 for (
auto val : getOperands())
4660CpAsyncBulkTensorGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
4662 auto thisOp = cast<NVVM::CpAsyncBulkTensorGlobalToSharedClusterOp>(op);
4663 const bool isCTAOnly = thisOp.getIsCTAOnly();
4667 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4669 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4679 const bool hasMC =
static_cast<bool>(mcMask);
4680 llvm::Value *i16Zero =
4681 llvm::ConstantInt::get(llvm::Type::getInt16Ty(mt.
getLLVMContext()), 0);
4685 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4686 llvm::Value *i64Zero =
4687 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4693 thisOp.getGroup() ? (
static_cast<int32_t
>(*thisOp.getGroup()) + 1) : 0;
4695 llvm::ConstantInt::get(llvm::Type::getInt32Ty(mt.
getLLVMContext()), val);
4698 llvm::Value *flagValidPattern = builder.getInt32(0);
4702 args.push_back(hasMC ? mt.
lookupValue(mcMask) : i16Zero);
4703 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4704 args.push_back(builder.getInt1(hasMC));
4705 args.push_back(builder.getInt1(hasCacheHint));
4707 args.push_back(flagValidPattern);
4710 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4711 args.push_back(builder.getInt1(hasCacheHint));
4712 args.push_back(flagValidPattern);
4715 constexpr size_t numDims = 5;
4716 constexpr size_t numModes = 5;
4717 using rowTy = std::array<llvm::Intrinsic::ID, numDims + 1>;
4718 using TableTy = std::array<rowTy, numModes>;
4719 static constexpr TableTy IDTable{
4720 {{
notIntrinsic, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d,
4721 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d,
4722 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d,
4723 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d,
4724 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d},
4726 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d,
4727 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d,
4728 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d},
4730 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_3d,
4731 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_4d,
4732 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_5d},
4734 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_3d,
4735 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_4d,
4736 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_5d},
4738 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_gather4_2d}}};
4740 static constexpr TableTy IDTableCTA{
4742 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_1d,
4743 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_2d,
4744 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_3d,
4745 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_4d,
4746 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_5d},
4748 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_3d,
4749 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_4d,
4750 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_5d},
4752 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_3d,
4753 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_4d,
4754 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_5d},
4756 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_3d,
4757 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_4d,
4758 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_5d},
4760 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_gather4_2d}}};
4763 (getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1) &&
4764 (getMaxEnumValForTMALoadMode() == std::size(IDTableCTA) - 1),
4765 "TMALoadModes must match number of rows in IDTable and IDTableCTA");
4766 size_t mode =
static_cast<size_t>(thisOp.getMode());
4767 size_t dim = thisOp.getCoordinates().size();
4768 auto id = isCTAOnly ? IDTableCTA[mode][dim] : IDTable[mode][dim];
4770 "Invalid intrinsic for CpAsyncBulkTensorGlobalToSharedClusterOp.");
4772 return {id, std::move(args)};
4777 auto thisOp = cast<NVVM::CpAsyncBulkTensorPrefetchOp>(op);
4781 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4783 for (
auto v : thisOp.getCoordinates())
4785 for (
auto v : thisOp.getIm2colOffsets())
4789 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4790 llvm::Value *i64Unused =
4791 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4792 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4793 args.push_back(builder.getInt1(hasCacheHint));
4795 const unsigned NI = llvm::Intrinsic::not_intrinsic;
4796 static constexpr llvm::Intrinsic::ID IDTable[][6] = {
4797 {NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_1d,
4798 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_2d,
4799 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_3d,
4800 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_4d,
4801 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_5d},
4803 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_3d,
4804 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_4d,
4805 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_5d},
4807 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_3d,
4808 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_4d,
4809 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_5d},
4811 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_3d,
4812 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_4d,
4813 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_5d},
4814 {NI, NI, NI, NI, NI,
4815 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_gather4_2d}};
4817 static_assert(getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1,
4818 "TMALoadModes must match number of rows in IDTable");
4819 size_t mode =
static_cast<size_t>(thisOp.getMode());
4820 size_t dim = thisOp.getCoordinates().size();
4821 llvm::Intrinsic::ID
id = IDTable[mode][dim];
4822 if (
id == llvm::Intrinsic::not_intrinsic)
4823 llvm_unreachable(
"Invalid intrinsic for CpAsyncBulkTensorPrefetchOp.");
4825 return {id, std::move(args)};
4829CpAsyncBulkTensorSharedCTAToGlobalOp::getIntrinsicIDAndArgs(
4831 auto thisOp = cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp>(op);
4835 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4836 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4838 for (
auto v : thisOp.getCoordinates())
4842 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4843 llvm::Value *i64Unused =
4844 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4845 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4846 args.push_back(builder.getInt1(hasCacheHint));
4848 using namespace llvm::Intrinsic;
4849 const unsigned NI = not_intrinsic;
4850 static constexpr ID IDTable[][6] = {
4851 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_1d,
4852 nvvm_cp_async_bulk_tensor_s2g_tile_2d,
4853 nvvm_cp_async_bulk_tensor_s2g_tile_3d,
4854 nvvm_cp_async_bulk_tensor_s2g_tile_4d,
4855 nvvm_cp_async_bulk_tensor_s2g_tile_5d},
4856 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_3d,
4857 nvvm_cp_async_bulk_tensor_s2g_im2col_4d,
4858 nvvm_cp_async_bulk_tensor_s2g_im2col_5d},
4859 {NI, NI, NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_2d},
4860 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_3d,
4861 nvvm_cp_async_bulk_tensor_s2g_im2col_w_4d,
4862 nvvm_cp_async_bulk_tensor_s2g_im2col_w_5d}};
4864 static_assert(getMaxEnumValForTMAStoreMode() == std::size(IDTable) - 1,
4865 "TMAStoreModes must match number of rows in IDTable");
4866 size_t mode =
static_cast<size_t>(thisOp.getMode());
4867 size_t dim = thisOp.getCoordinates().size();
4868 ID
id = IDTable[mode][dim];
4869 if (
id == llvm::Intrinsic::not_intrinsic)
4871 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOp.");
4873 return {id, std::move(args)};
4877CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
4880 cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp>(op);
4883 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4884 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4885 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
4886 for (
Value v : thisOp.getTensorSize())
4888 for (
Value v : thisOp.getLowerStride())
4890 if (thisOp.getUpperStride())
4891 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
4892 for (
Value v : thisOp.getCoordinates())
4896 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4897 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4898 : builder.getInt64(0));
4899 args.push_back(builder.getInt1(hasCacheHint));
4901 using namespace llvm::Intrinsic;
4902 const unsigned NI = not_intrinsic;
4905 static constexpr ID IDTable[][6] = {
4906 {NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_1d,
4907 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_2d,
4908 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_3d,
4909 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_4d,
4910 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_5d},
4911 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_3d,
4912 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_4d,
4913 nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_5d},
4914 {NI, NI, NI, NI, NI,
4915 nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_override_addr_2d},
4916 {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_3d,
4917 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_4d,
4918 nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_5d}};
4922 static constexpr ID dimStrideIDTable[] = {
4923 NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_1d,
4924 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_2d,
4925 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_3d,
4926 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_4d,
4927 nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_5d};
4930 size_t mode =
static_cast<size_t>(thisOp.getMode());
4931 size_t dim = thisOp.getCoordinates().size();
4932 bool isDimStride = !thisOp.getTensorSize().empty();
4934 assert(mode < std::size(IDTable) &&
4935 "Invalid mode for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4936 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
4937 "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4939 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
4941 intrinsicID != NI &&
4942 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
4943 return {intrinsicID, std::move(args)};
4948 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOp>(op);
4951 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4952 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4953 for (
Value v : thisOp.getCoordinates())
4957 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4958 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
4959 : builder.getInt64(0));
4960 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
4961 args.push_back(builder.getInt1(hasCacheHint));
4963 using namespace llvm::Intrinsic;
4964 const unsigned NI = not_intrinsic;
4965 static constexpr ID IDTable[][6] = {
4966 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_1d,
4967 nvvm_cp_async_bulk_tensor_reduce_tile_2d,
4968 nvvm_cp_async_bulk_tensor_reduce_tile_3d,
4969 nvvm_cp_async_bulk_tensor_reduce_tile_4d,
4970 nvvm_cp_async_bulk_tensor_reduce_tile_5d},
4971 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_3d,
4972 nvvm_cp_async_bulk_tensor_reduce_im2col_4d,
4973 nvvm_cp_async_bulk_tensor_reduce_im2col_5d},
4974 {NI, NI, NI, NI, NI, NI},
4975 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_3d,
4976 nvvm_cp_async_bulk_tensor_reduce_im2col_w_4d,
4977 nvvm_cp_async_bulk_tensor_reduce_im2col_w_5d}};
4979 size_t mode =
static_cast<size_t>(thisOp.getMode());
4980 size_t dim = thisOp.getCoordinates().size();
4981 assert(mode < std::size(IDTable) &&
4982 "Invalid mode for CpAsyncBulkTensorReduceOp");
4983 assert(dim < std::size(IDTable[mode]) &&
4984 "Invalid dim for CpAsyncBulkTensorReduceOp");
4986 ID intrinsicID = IDTable[mode][dim];
4987 assert(intrinsicID != NI &&
4988 "Invalid intrinsic for CpAsyncBulkTensorReduceOp");
4989 return {intrinsicID, std::move(args)};
4992NVVM::IDArgPair CpAsyncBulkTensorReduceOverrideAddrOp::getIntrinsicIDAndArgs(
4994 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOverrideAddrOp>(op);
4997 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4998 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4999 args.push_back(mt.
lookupValue(thisOp.getOverrideAddr()));
5001 for (
Value v : thisOp.getTensorSize())
5003 for (
Value v : thisOp.getLowerStride())
5005 if (thisOp.getUpperStride())
5006 args.push_back(mt.
lookupValue(thisOp.getUpperStride()));
5007 for (
Value v : thisOp.getCoordinates())
5011 const bool hasCacheHint =
static_cast<bool>(cacheHint);
5012 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint)
5013 : builder.getInt64(0));
5014 args.push_back(builder.getInt32(
static_cast<uint32_t
>(thisOp.getRedKind())));
5015 args.push_back(builder.getInt1(hasCacheHint));
5017 using namespace llvm::Intrinsic;
5018 const unsigned NI = not_intrinsic;
5021static constexpr ID IDTable[][6] = {
5022 {NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_1d,
5023 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_2d,
5024 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_3d,
5025 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_4d,
5026 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_5d},
5027 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_3d,
5028 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_4d,
5029 nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_5d},
5030 {NI, NI, NI, NI, NI, NI},
5031 {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_3d,
5032 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_4d,
5033 nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_5d}};
5037static constexpr ID dimStrideIDTable[] = {
5038 NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_1d,
5039 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_2d,
5040 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_3d,
5041 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_4d,
5042 nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_5d};
5045 size_t mode =
static_cast<size_t>(thisOp.getMode());
5046 size_t dim = thisOp.getCoordinates().size();
5047 bool isDimStride = !thisOp.getTensorSize().empty();
5049 assert(mode < std::size(IDTable) &&
5050 "Invalid mode for CpAsyncBulkTensorReduceOverrideAddrOp");
5051 assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
5052 "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
5054 ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
5055 assert(intrinsicID != NI &&
5056 "Invalid intrinsic for CpAsyncBulkTensorReduceOverrideAddrOp");
5057 return {intrinsicID, std::move(args)};
5062#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5063 hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
5064 : llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
5066#define GET_CVT_F2TF32_ID(rnd, relu, sf) \
5067 hasSatFinite ? CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
5068 : CVT_F2TF32_ID_IMPL(rnd, relu, )
5071ConvertFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
5072 NVVM::SaturationMode sat,
bool hasRelu) {
5073 using RndMode = NVVM::FPRoundingMode;
5074 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5083 llvm_unreachable(
"Invalid RoundingMode for CvtFloatToTF32Op");
5088ConvertF32x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF4x2Op op,
5090 llvm::IRBuilderBase &builder) {
5094 args.push_back(builder.getInt1(op.getPzo()));
5097 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5098 static_cast<unsigned>(op.getRelu());
5099 static constexpr llvm::Intrinsic::ID ids[] = {
5100 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_satfinite,
5101 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_relu_satfinite,
5102 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rz_satfinite,
5103 llvm::Intrinsic::nvvm_ff_to_e2m1x2_rz_relu_satfinite,
5106 return {ids[idx], std::move(args)};
5110ConvertF32x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5112 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5113 static_cast<unsigned>(hasRelu);
5114 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5115 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rn_satfinite,
5116 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rn_relu_satfinite,
5117 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rz_satfinite,
5118 llvm::Intrinsic::nvvm_ff_to_e2m3x2_rz_relu_satfinite,
5120 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5121 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rn_satfinite,
5122 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rn_relu_satfinite,
5123 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rz_satfinite,
5124 llvm::Intrinsic::nvvm_ff_to_e3m2x2_rz_relu_satfinite,
5128 .Case([&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5129 .Case([&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5131 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF6x2Op");
5132 return llvm::Intrinsic::not_intrinsic;
5137ConvertF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF16x2ToF4x2Op &op,
5139 llvm::IRBuilderBase &builder) {
5143 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5144 static_cast<unsigned>(op.getRelu());
5145 static constexpr llvm::Intrinsic::ID ids[] = {
5146 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_satfinite,
5147 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_relu_satfinite,
5148 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rz_satfinite,
5149 llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rz_relu_satfinite,
5152 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5153 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5158 args.push_back(builder.getInt1(op.getPzo()));
5160 return {intId, std::move(args)};
5164ConvertBF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertBF16x2ToF4x2Op &op,
5166 llvm::IRBuilderBase &builder) {
5170 (
static_cast<unsigned>(op.getRnd() == NVVM::FPRoundingMode::RZ) << 1) |
5171 static_cast<unsigned>(op.getRelu());
5172 static constexpr llvm::Intrinsic::ID ids[] = {
5173 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_satfinite,
5174 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite,
5175 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rz_satfinite,
5176 llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rz_relu_satfinite,
5179 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
5180 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
5185 args.push_back(builder.getInt1(op.getPzo()));
5187 return {intId, std::move(args)};
5191ConvertF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5193 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5194 static_cast<unsigned>(hasRelu);
5195 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5196 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_satfinite,
5197 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_relu_satfinite,
5198 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rz_satfinite,
5199 llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rz_relu_satfinite,
5201 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5202 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_satfinite,
5203 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_relu_satfinite,
5204 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rz_satfinite,
5205 llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rz_relu_satfinite,
5209 .Case<mlir::Float6E2M3FNType>(
5210 [&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5211 .Case<mlir::Float6E3M2FNType>(
5212 [&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5214 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF6x2Op");
5215 return llvm::Intrinsic::not_intrinsic;
5220ConvertBF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
5221 NVVM::FPRoundingMode rnd,
bool hasRelu) {
5222 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5223 static_cast<unsigned>(hasRelu);
5224 static constexpr llvm::Intrinsic::ID e2m3x2IDs[] = {
5225 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_satfinite,
5226 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite,
5227 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rz_satfinite,
5228 llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite,
5230 static constexpr llvm::Intrinsic::ID e3m2x2IDs[] = {
5231 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_satfinite,
5232 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite,
5233 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rz_satfinite,
5234 llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite,
5238 .Case<mlir::Float6E2M3FNType>(
5239 [&](mlir::Float6E2M3FNType) {
return e2m3x2IDs[idx]; })
5240 .Case<mlir::Float6E3M2FNType>(
5241 [&](mlir::Float6E3M2FNType) {
return e3m2x2IDs[idx]; })
5243 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF6x2Op");
5244 return llvm::Intrinsic::not_intrinsic;
5249ConvertF32x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5250 NVVM::SaturationMode sat,
bool hasRelu) {
5251 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5252 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5253 static_cast<unsigned>(hasRelu);
5254 unsigned ue8m0Idx = (
static_cast<unsigned>(hasSatFinite) << 1) |
5255 static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RP);
5257 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5258 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rn,
5259 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rn_relu,
5260 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rz,
5261 llvm::Intrinsic::nvvm_ff_to_e4m3x2_rz_relu,
5263 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5264 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rn,
5265 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rn_relu,
5266 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rz,
5267 llvm::Intrinsic::nvvm_ff_to_e5m2x2_rz_relu,
5269 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
5270 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rz,
5271 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rp,
5272 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rz_satfinite,
5273 llvm::Intrinsic::nvvm_ff_to_ue8m0x2_rp_satfinite,
5277 .Case([&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5278 .Case([&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5279 .Case([&](mlir::Float8E8M0FNUType) {
5280 if (rnd != NVVM::FPRoundingMode::RZ && rnd != NVVM::FPRoundingMode::RP)
5281 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5282 return ue8m0x2IDs[ue8m0Idx];
5285 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
5286 return llvm::Intrinsic::not_intrinsic;
5291ConvertF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
5293 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5294 static_cast<unsigned>(hasRelu);
5295 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5296 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rn,
5297 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rn_relu,
5298 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rz,
5299 llvm::Intrinsic::nvvm_f16x2_to_e4m3x2_rz_relu,
5301 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5302 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rn,
5303 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rn_relu,
5304 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rz,
5305 llvm::Intrinsic::nvvm_f16x2_to_e5m2x2_rz_relu,
5309 .Case([&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5310 .Case([&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5312 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF8x2Op");
5313 return llvm::Intrinsic::not_intrinsic;
5318ConvertBF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
5319 NVVM::FPRoundingMode rnd,
5320 NVVM::SaturationMode sat,
bool hasRelu) {
5321 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
5322 unsigned idx = (
static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RZ) << 1) |
5323 static_cast<unsigned>(hasRelu);
5324 unsigned ue8m0Idx = (
static_cast<unsigned>(hasSatFinite) << 1) |
5325 static_cast<unsigned>(rnd == NVVM::FPRoundingMode::RP);
5327 static constexpr llvm::Intrinsic::ID e4m3x2IDs[] = {
5328 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_satfinite,
5329 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite,
5330 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rz_satfinite,
5331 llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite,
5333 static constexpr llvm::Intrinsic::ID e5m2x2IDs[] = {
5334 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_satfinite,
5335 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite,
5336 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rz_satfinite,
5337 llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite,
5339 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
5340 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz,
5341 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp,
5342 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz_satfinite,
5343 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp_satfinite,
5347 .Case<mlir::Float8E4M3FNType>(
5348 [&](mlir::Float8E4M3FNType) {
return e4m3x2IDs[idx]; })
5349 .Case<mlir::Float8E5M2Type>(
5350 [&](mlir::Float8E5M2Type) {
return e5m2x2IDs[idx]; })
5351 .Case<mlir::Float8E8M0FNUType>(
5352 [&](mlir::Float8E8M0FNUType) {
return ue8m0x2IDs[ue8m0Idx]; })
5354 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
5355 return llvm::Intrinsic::not_intrinsic;
5361 auto curOp = cast<NVVM::ConvertF8x2ToF16x2Op>(op);
5363 bool hasRelu = curOp.getRelu();
5365 llvm::Intrinsic::ID intId =
5367 .Case([&](Float8E4M3FNType type) {
5368 return hasRelu ? llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn_relu
5369 : llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn;
5371 .Case([&](Float8E5M2Type type) {
5372 return hasRelu ? llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn_relu
5373 : llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn;
5376 llvm_unreachable(
"Invalid type for ConvertF8x2ToF16x2Op");
5377 return llvm::Intrinsic::not_intrinsic;
5380 llvm::Value *packedI16 =
5381 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5382 llvm::Type::getInt16Ty(builder.getContext()));
5384 return {intId, {packedI16}};
5389 auto curOp = cast<NVVM::ConvertF8x2ToBF16x2Op>(op);
5390 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5391 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5392 bool hasRelu = curOp.getRelu();
5394 static constexpr llvm::Intrinsic::ID E4M3Ids[] = {
5395 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5396 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5397 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5398 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5401 static constexpr llvm::Intrinsic::ID E5M2Ids[] = {
5402 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5403 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5404 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5405 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5408 llvm::Intrinsic::ID intId =
5410 .Case([&](Float8E8M0FNUType type) {
5411 return llvm::Intrinsic::nvvm_ue8m0x2_to_bf16x2;
5413 .Case([&](Float8E4M3FNType type) {
5414 return E4M3Ids[hasSatfinite << 1 | hasRelu];
5416 .Case([&](Float8E5M2Type type) {
5417 return E5M2Ids[hasSatfinite << 1 | hasRelu];
5420 llvm_unreachable(
"Invalid type for ConvertF8x2ToBF16x2Op");
5421 return llvm::Intrinsic::not_intrinsic;
5423 llvm::Value *packedI16 =
5424 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5425 llvm::Type::getInt16Ty(builder.getContext()));
5428 args.push_back(packedI16);
5429 if (!isa<Float8E8M0FNUType>(curOp.getSrcType()))
5432 : builder.getInt16(0x7f7f));
5435 return {intId, std::move(args)};
5440 auto curOp = cast<NVVM::ConvertF6x2ToF16x2Op>(op);
5442 bool hasRelu = curOp.getRelu();
5444 llvm::Intrinsic::ID intId =
5446 .Case([&](Float6E2M3FNType type) {
5447 return hasRelu ? llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn_relu
5448 : llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn;
5450 .Case([&](Float6E3M2FNType type) {
5451 return hasRelu ? llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn_relu
5452 : llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn;
5455 llvm_unreachable(
"Invalid type for ConvertF6x2ToF16x2Op");
5456 return llvm::Intrinsic::not_intrinsic;
5459 llvm::Value *packedI16 =
5460 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5461 llvm::Type::getInt16Ty(builder.getContext()));
5463 return {intId, {packedI16}};
5468 auto curOp = cast<NVVM::ConvertF6x2ToBF16x2Op>(op);
5469 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5470 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5471 bool hasRelu = curOp.getRelu();
5473 static constexpr llvm::Intrinsic::ID E2M3Ids[] = {
5474 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_scale_n2_ue8m0,
5475 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5476 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5477 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5480 static constexpr llvm::Intrinsic::ID E3M2Ids[] = {
5481 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_scale_n2_ue8m0,
5482 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5483 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5484 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5487 unsigned idx = (hasSatfinite << 1) | hasRelu;
5488 llvm::Intrinsic::ID intId =
5490 .Case([&](Float6E2M3FNType type) {
return E2M3Ids[idx]; })
5491 .Case([&](Float6E3M2FNType type) {
return E3M2Ids[idx]; })
5493 llvm_unreachable(
"Invalid type for ConvertF6x2ToBF16x2Op");
5494 return llvm::Intrinsic::not_intrinsic;
5497 llvm::Value *packedI16 =
5498 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
5499 llvm::Type::getInt16Ty(builder.getContext()));
5502 args.push_back(packedI16);
5509 return {intId, std::move(args)};
5514 auto curOp = cast<NVVM::ConvertF4x2ToF16x2Op>(op);
5516 bool hasRelu = curOp.getRelu();
5518 llvm::Intrinsic::ID intId =
5520 .Case([&](Float4E2M1FNType type) {
5521 return hasRelu ? llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn_relu
5522 : llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn;
5525 llvm_unreachable(
"Invalid type for ConvertF4x2ToF16x2Op");
5526 return llvm::Intrinsic::not_intrinsic;
5529 llvm::Value *extendedI16 =
5530 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5531 llvm::Type::getInt16Ty(builder.getContext()));
5533 return {intId, {extendedI16}};
5538 auto curOp = cast<NVVM::ConvertF4x2ToBF16x2Op>(op);
5539 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
5540 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
5541 bool hasRelu = curOp.getRelu();
5543 static constexpr llvm::Intrinsic::ID E2M1Ids[] = {
5544 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_scale_n2_ue8m0,
5545 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5546 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5547 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5550 unsigned idx = (hasSatfinite << 1) | hasRelu;
5551 llvm::Intrinsic::ID intId =
5553 .Case([&](Float4E2M1FNType type) {
return E2M1Ids[idx]; })
5555 llvm_unreachable(
"Invalid type for ConvertF4x2ToBF16x2Op");
5556 return llvm::Intrinsic::not_intrinsic;
5559 llvm::Value *extendedI16 =
5560 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
5561 llvm::Type::getInt16Ty(builder.getContext()));
5564 args.push_back(extendedI16);
5571 return {intId, std::move(args)};
5576 auto thisOp = cast<NVVM::ConvertF32x2ToS2F6x2Op>(op);
5577 bool hasRelu = thisOp.getRelu();
5578 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5580 llvm::Intrinsic::ID
id =
5582 ? llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5583 : llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5589 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5590 : builder.getInt16(0x7f7f));
5591 return {id, std::move(args)};
5596 auto thisOp = cast<NVVM::ConvertBF16x2ToS2F6x2Op>(op);
5597 bool hasRelu = thisOp.getRelu();
5598 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5600 llvm::Intrinsic::ID
id =
5603 nvvm_bf16x2_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5604 : llvm::Intrinsic::nvvm_bf16x2_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5609 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5610 : builder.getInt16(0x7f7f));
5611 return {id, std::move(args)};
5616 auto thisOp = cast<NVVM::ConvertS2F6x2ToBF16x2Op>(op);
5617 bool hasRelu = thisOp.getRelu();
5618 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5619 bool hasSat = thisOp.getSat() == NVVM::SaturationMode::SATFINITE;
5621 static constexpr llvm::Intrinsic::ID ids[] = {
5622 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_scale_n2_ue8m0,
5623 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5624 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5625 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5628 unsigned idx = (hasSat << 1) | hasRelu;
5632 llvm::Value *packedI16 =
5633 builder.CreateBitCast(mt.
lookupValue(thisOp.getSrc()),
5634 llvm::Type::getInt16Ty(builder.getContext()));
5635 args.push_back(packedI16);
5636 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5637 : builder.getInt16(0x7f7f));
5639 return {ids[idx], std::move(args)};
5644 auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
5645 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5647 llvm::Intrinsic::ID
id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
5648 : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
5652 builder.getInt1(curOp.getIsExclusive())}};
5657 auto curOp = cast<NVVM::Tcgen05DeallocOp>(op);
5658 llvm::Intrinsic::ID
id = (curOp.getGroup() == CTAGroupKind::CTA_1)
5659 ? llvm::Intrinsic::nvvm_tcgen05_dealloc_cg1
5660 : llvm::Intrinsic::nvvm_tcgen05_dealloc_cg2;
5664 builder.getInt1(curOp.getIsExclusive())}};
5668Tcgen05CommitOp::getIntrinsicIDAndArgs(
Operation &op,
5671 auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
5672 bool hasMulticast =
static_cast<bool>(curOp.getMulticastMask());
5673 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5674 bool hasSmemARead = curOp.getSmemARead();
5675 unsigned index = (
static_cast<unsigned>(hasSmemARead) << 1) |
5676 static_cast<unsigned>(is2CTAMode);
5678 using namespace llvm::Intrinsic;
5679 static constexpr ID IDs[] = {
5680 nvvm_tcgen05_commit_cg1,
5681 nvvm_tcgen05_commit_cg2,
5682 nvvm_tcgen05_commit_smem_a_read_cg1,
5683 nvvm_tcgen05_commit_smem_a_read_cg2,
5686 static constexpr ID multicastIDs[] = {
5687 nvvm_tcgen05_commit_mc_cg1,
5688 nvvm_tcgen05_commit_mc_cg2,
5689 nvvm_tcgen05_commit_smem_a_read_mc_cg1,
5690 nvvm_tcgen05_commit_smem_a_read_mc_cg2,
5693 ID
id = hasMulticast ? multicastIDs[
index] : IDs[
index];
5697 args.push_back(mt.
lookupValue(curOp.getMulticastMask()));
5702#define TCGEN05_CP_IMPL(shape_mc, src_fmt, cg) \
5703 llvm::Intrinsic::nvvm_tcgen05_cp##shape_mc##src_fmt##cg
5705#define TCGEN05_CP_2CTA(shape_mc, src_fmt, is_2cta) \
5706 is_2cta ? TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg2) \
5707 : TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
5709#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
5711 if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
5712 return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
5713 if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
5714 return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
5715 return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
5719ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
5721 llvm::IRBuilderBase &builder) {
5722 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5723 llvm::Intrinsic::nvvm_ff2f16x2_rn,
5724 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu,
5725 llvm::Intrinsic::nvvm_ff2f16x2_rn_satfinite,
5726 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu_satfinite,
5728 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5729 llvm::Intrinsic::nvvm_ff2f16x2_rz,
5730 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu,
5731 llvm::Intrinsic::nvvm_ff2f16x2_rz_satfinite,
5732 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu_satfinite,
5734 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5735 llvm::Intrinsic::nvvm_ff2f16x2_rs,
5736 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu,
5737 llvm::Intrinsic::nvvm_ff2f16x2_rs_satfinite,
5738 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu_satfinite,
5741 unsigned hasRelu = op.getRelu() ? 1 : 0;
5742 unsigned hasSatFinite =
5743 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5746 unsigned idx = (hasSatFinite << 1) | hasRelu;
5751 if (op.getRandomBits())
5752 args.push_back(mt.
lookupValue(op.getRandomBits()));
5756 if (op.getRnd() != FPRoundingMode::RS)
5757 args.push_back(builder.getInt1(
false));
5759 switch (op.getRnd()) {
5760 case FPRoundingMode::RN:
5761 return {rndRNIds[idx], std::move(args)};
5762 case FPRoundingMode::RZ:
5763 return {rndRZIds[idx], std::move(args)};
5764 case FPRoundingMode::RS:
5765 return {rndRSIds[idx], std::move(args)};
5767 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToF16x2Op");
5772ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
5774 llvm::IRBuilderBase &builder) {
5775 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5776 llvm::Intrinsic::nvvm_ff2bf16x2_rn,
5777 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu,
5778 llvm::Intrinsic::nvvm_ff2bf16x2_rn_satfinite,
5779 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu_satfinite,
5781 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5782 llvm::Intrinsic::nvvm_ff2bf16x2_rz,
5783 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu,
5784 llvm::Intrinsic::nvvm_ff2bf16x2_rz_satfinite,
5785 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu_satfinite,
5787 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5788 llvm::Intrinsic::nvvm_ff2bf16x2_rs,
5789 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu,
5790 llvm::Intrinsic::nvvm_ff2bf16x2_rs_satfinite,
5791 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu_satfinite,
5794 unsigned hasRelu = op.getRelu() ? 1 : 0;
5795 unsigned hasSatFinite =
5796 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5799 unsigned idx = (hasSatFinite << 1) | hasRelu;
5804 if (op.getRandomBits())
5805 args.push_back(mt.
lookupValue(op.getRandomBits()));
5808 if (op.getRnd() != FPRoundingMode::RS)
5809 args.push_back(builder.getInt1(
false));
5811 switch (op.getRnd()) {
5812 case FPRoundingMode::RN:
5813 return {rndRNIds[idx], std::move(args)};
5814 case FPRoundingMode::RZ:
5815 return {rndRZIds[idx], std::move(args)};
5816 case FPRoundingMode::RS:
5817 return {rndRSIds[idx], std::move(args)};
5819 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToBF16x2Op");
5823llvm::Intrinsic::ID ConvertF32x4ToF8x4Op::getIntrinsicID() {
5825 bool hasRelu = getRelu();
5828 .Case([&](mlir::Float8E4M3FNType) {
5829 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_relu_satfinite
5830 : llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_satfinite;
5832 .Case([&](mlir::Float8E5M2Type) {
5833 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_relu_satfinite
5834 : llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_satfinite;
5837 llvm_unreachable(
"Invalid F8 type in ConvertF32x4ToF8x4Op");
5838 return llvm::Intrinsic::not_intrinsic;
5842llvm::Intrinsic::ID ConvertF32x4ToF6x4Op::getIntrinsicID() {
5844 bool hasRelu = getRelu();
5847 .Case([&](mlir::Float6E2M3FNType) {
5848 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_relu_satfinite
5849 : llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_satfinite;
5851 .Case([&](mlir::Float6E3M2FNType) {
5852 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_relu_satfinite
5853 : llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_satfinite;
5856 llvm_unreachable(
"Invalid F6 type in ConvertF32x4ToF6x4Op");
5857 return llvm::Intrinsic::not_intrinsic;
5861llvm::Intrinsic::ID ConvertF32x4ToF4x4Op::getIntrinsicID() {
5863 bool hasRelu = getRelu();
5866 .Case([&](mlir::Float4E2M1FNType) {
5867 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_relu_satfinite
5868 : llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_satfinite;
5871 llvm_unreachable(
"Invalid F4 type in ConvertF32x4ToF4x4Op");
5872 return llvm::Intrinsic::not_intrinsic;
5876llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(
Operation &op) {
5877 auto curOp = cast<NVVM::Tcgen05CpOp>(op);
5878 bool is2CTA = curOp.getGroup() == CTAGroupKind::CTA_2;
5879 auto srcFmt = curOp.getSrcFormat();
5880 auto mc = curOp.getMulticast();
5882 switch (curOp.getShape()) {
5883 case Tcgen05CpShape::SHAPE_128x256b:
5885 case Tcgen05CpShape::SHAPE_128x128b:
5887 case Tcgen05CpShape::SHAPE_4x256b:
5889 case Tcgen05CpShape::SHAPE_32x128b:
5891 case Tcgen05CpShape::SHAPE_64x128b:
5892 return (mc == Tcgen05CpMulticast::WARPX2_01_23)
5896 llvm_unreachable(
"Invalid shape in tcgen05 cp Op");
5903 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
5905 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
5910LogicalResult Tcgen05LdOp::verify() {
5912 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5915 if (
getShape() != NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && getOffset())
5916 result =
emitError(
"offset argument is only supported for shape 16x32bx2");
5918 auto resTy = getRes().getType();
5919 unsigned resLen = isa<VectorType>(resTy)
5920 ? llvm::cast<VectorType>(resTy).getNumElements()
5923 result =
emitError(llvm::formatv(
"invalid result type length {0} for shape "
5924 "{1} in tcgen05.ld Op",
5925 resLen, stringifyEnum(
getShape())));
5930LogicalResult Tcgen05StOp::verify() {
5932 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5935 auto valTy = getVal().getType();
5936 unsigned valLen = isa<VectorType>(valTy)
5937 ? llvm::cast<VectorType>(valTy).getNumElements()
5940 result =
emitError(llvm::formatv(
"invalid input length {0} for shape "
5941 "{1} in tcgen05.st Op",
5942 valLen, stringifyEnum(
getShape())));
5950 std::optional<LLVM::ConstantRangeAttr> range,
Value result,
5954 setResultRanges(
result, {range->getLower(), range->getUpper(),
5955 range->getLower(), range->getUpper()});
5964 Operation *op, std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
5968 const llvm::APInt &lower = rangeAttr->getLower();
5969 const llvm::APInt &upper = rangeAttr->getUpper();
5972 if (lower == upper && !lower.isMaxValue() && !lower.isMinValue()) {
5973 unsigned bitWidth = lower.getBitWidth();
5974 llvm::APInt minVal = llvm::APInt::getMinValue(bitWidth);
5975 llvm::APInt maxVal = llvm::APInt::getMaxValue(bitWidth);
5977 "invalid range attribute: Lower == Upper, but they aren't min (")
5978 << llvm::toString(minVal, 10,
false) <<
") or max ("
5979 << llvm::toString(maxVal, 10,
false)
5980 <<
") value! This is an invalid constant range.";
5987 llvm::IRBuilderBase &builder) {
5988 return builder.CreateBitCast(arg,
5989 llvm::Type::getInt32Ty(builder.getContext()));
5994 auto curOp = cast<NVVM::DotAccumulate4WayOp>(op);
6001 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
6002 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
6003 unsigned type = (isASigned << 1) | isBSigned;
6004 const llvm::Intrinsic::ID ids[] = {
6005 llvm::Intrinsic::nvvm_idp4a_u_u,
6006 llvm::Intrinsic::nvvm_idp4a_u_s,
6007 llvm::Intrinsic::nvvm_idp4a_s_u,
6008 llvm::Intrinsic::nvvm_idp4a_s_s,
6010 return {ids[type], args};
6015 auto curOp = cast<NVVM::DotAccumulate2WayOp>(op);
6020 args.push_back(builder.getInt1(curOp.getBHi()));
6023 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
6024 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
6025 unsigned type = (isASigned << 1) | isBSigned;
6026 const llvm::Intrinsic::ID ids[] = {
6027 llvm::Intrinsic::nvvm_idp2a_u_u,
6028 llvm::Intrinsic::nvvm_idp2a_u_s,
6029 llvm::Intrinsic::nvvm_idp2a_s_u,
6030 llvm::Intrinsic::nvvm_idp2a_s_s,
6032 return {ids[type], args};
6036 llvm::IRBuilderBase &builder) {
6037 return builder.CreateAddrSpaceCast(
6038 addr, builder.getPtrTy(llvm::NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM));
6042PrefetchOp::getIntrinsicIDAndArgs(NVVM::PrefetchOp &op,
6044 llvm::IRBuilderBase &builder) {
6045 using MemSpace = NVVM::NVVMMemorySpace;
6046 using CacheLevel = NVVM::PrefetchCacheLevel;
6048 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = op.getCacheLevel();
6049 std::optional<NVVM::CacheEvictionPriority> evictPriority =
6050 op.getEvictPriority();
6051 unsigned addressSpace =
6052 llvm::cast<LLVM::LLVMPointerType>(op.getAddr().getType())
6060 if (op.getTensormap())
6061 return {llvm::Intrinsic::nvvm_prefetch_tensormap, args};
6063 assert(cacheLevel &&
"expected cache level for non-tensormap prefetch");
6065 if (op.getUniform() && *cacheLevel == CacheLevel::L1)
6066 return {llvm::Intrinsic::nvvm_prefetchu_L1, args};
6068 if (evictPriority && *cacheLevel == CacheLevel::L2) {
6069 switch (*evictPriority) {
6070 case NVVM::CacheEvictionPriority::EvictLast:
6071 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_last, args};
6072 case NVVM::CacheEvictionPriority::EvictNormal:
6073 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_normal, args};
6075 llvm_unreachable(
"Invalid cache eviction priority");
6079 switch (
static_cast<MemSpace
>(addressSpace)) {
6080 case MemSpace::Generic:
6081 return *cacheLevel == CacheLevel::L1
6083 :
NVVM::
IDArgPair({llvm::Intrinsic::nvvm_prefetch_L2, args});
6084 case MemSpace::Global:
6085 return *cacheLevel == CacheLevel::L1
6087 {llvm::Intrinsic::nvvm_prefetch_global_L1, args})
6089 {llvm::Intrinsic::nvvm_prefetch_global_L2, args});
6090 case MemSpace::Local:
6091 return *cacheLevel == CacheLevel::L1
6093 {llvm::Intrinsic::nvvm_prefetch_local_L1, args})
6095 {llvm::Intrinsic::nvvm_prefetch_local_L2, args});
6097 llvm_unreachable(
"Invalid pointer address space");
6101bool NVVM::InlinePtxOp::getAsmValues(
6105 for (
auto arg : getReadWriteArgs())
6107 for (
auto arg : getResults())
6109 for (
auto arg : getReadOnlyArgs())
6116NVVM::IDArgPair ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
6118 auto curOp = cast<NVVM::ClusterLaunchControlTryCancelOp>(op);
6120 args.push_back(mt.
lookupValue(curOp.getSmemAddress()));
6121 args.push_back(mt.
lookupValue(curOp.getMbarrier()));
6123 llvm::Intrinsic::ID intrinsicID =
6124 curOp.getMulticast()
6126 nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
6127 : llvm::Intrinsic::nvvm_clusterlaunchcontrol_try_cancel_async_shared;
6129 return {intrinsicID, args};
6132NVVM::IDArgPair ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
6134 auto curOp = cast<NVVM::ClusterLaunchControlQueryCancelOp>(op);
6136 args.push_back(mt.
lookupValue(curOp.getTryCancelResponse()));
6138 llvm::Intrinsic::ID intrinsicID;
6140 switch (curOp.getQueryType()) {
6141 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
6143 llvm::Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled;
6145 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
6146 intrinsicID = llvm::Intrinsic::
6147 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x;
6149 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
6150 intrinsicID = llvm::Intrinsic::
6151 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y;
6153 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
6154 intrinsicID = llvm::Intrinsic::
6155 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z;
6158 return {intrinsicID, args};
6163 llvm::IRBuilderBase &builder) {
6164 auto thisOp = cast<NVVM::PermuteOp>(op);
6165 NVVM::PermuteMode mode = thisOp.getMode();
6167 static constexpr llvm::Intrinsic::ID IDs[] = {
6168 llvm::Intrinsic::nvvm_prmt, llvm::Intrinsic::nvvm_prmt_f4e,
6169 llvm::Intrinsic::nvvm_prmt_b4e, llvm::Intrinsic::nvvm_prmt_rc8,
6170 llvm::Intrinsic::nvvm_prmt_ecl, llvm::Intrinsic::nvvm_prmt_ecr,
6171 llvm::Intrinsic::nvvm_prmt_rc16};
6173 unsigned modeIndex =
static_cast<unsigned>(mode);
6181 args.push_back(mt.
lookupValue(thisOp.getSelector()));
6183 return {IDs[modeIndex], args};
6188 auto thisOp = cast<NVVM::TensormapReplaceOp>(op);
6192 if (thisOp.getOrd())
6193 args.push_back(builder.getInt32(thisOp.getOrd().value()));
6194 if (thisOp.getNewValue())
6195 args.push_back(mt.
lookupValue(thisOp.getNewValue()));
6196 if (
auto attr = thisOp.getNewValueAttr()) {
6199 .Case<TensormapElemtypeAttr, TensormapInterleaveLayoutAttr,
6200 TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
6201 TensormapFillModeAttr>([](
auto attr) {
6202 return static_cast<unsigned>(attr.getValue());
6204 .Default([](
auto attr) {
6205 llvm_unreachable(
"Invalid attribute type");
6208 args.push_back(builder.getInt32(val));
6211 static constexpr llvm::Intrinsic::ID IDs[] = {
6212 llvm::Intrinsic::nvvm_tensormap_replace_global_address,
6213 llvm::Intrinsic::nvvm_tensormap_replace_rank,
6214 llvm::Intrinsic::nvvm_tensormap_replace_box_dim,
6215 llvm::Intrinsic::nvvm_tensormap_replace_global_dim,
6216 llvm::Intrinsic::nvvm_tensormap_replace_global_stride,
6217 llvm::Intrinsic::nvvm_tensormap_replace_element_stride,
6218 llvm::Intrinsic::nvvm_tensormap_replace_elemtype,
6219 llvm::Intrinsic::nvvm_tensormap_replace_interleave_layout,
6220 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_mode,
6221 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_atomicity,
6222 llvm::Intrinsic::nvvm_tensormap_replace_fill_mode,
6225 unsigned fieldIndex =
static_cast<unsigned>(thisOp.getField());
6227 return {IDs[fieldIndex], args};
6234static llvm::nvvm::Tcgen05MMAKind
6237 case NVVM::Tcgen05MMAKind::F16:
6238 return llvm::nvvm::Tcgen05MMAKind::F16;
6239 case NVVM::Tcgen05MMAKind::TF32:
6240 return llvm::nvvm::Tcgen05MMAKind::TF32;
6241 case NVVM::Tcgen05MMAKind::F8F6F4:
6242 return llvm::nvvm::Tcgen05MMAKind::F8F6F4;
6243 case NVVM::Tcgen05MMAKind::I8:
6244 return llvm::nvvm::Tcgen05MMAKind::I8;
6245 case NVVM::Tcgen05MMAKind::TI16:
6246 return llvm::nvvm::Tcgen05MMAKind::TI16;
6247 case NVVM::Tcgen05MMAKind::MXF8F6F4:
6248 case NVVM::Tcgen05MMAKind::MXF4:
6249 case NVVM::Tcgen05MMAKind::MXF4NVF4:
6252 llvm_unreachable(
"Unsupported tcgen05.mma kind");
6258 llvm::IRBuilderBase &builder) {
6260 auto thisOp = cast<NVVM::Tcgen05MMAOp>(op);
6263 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6266 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6269 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6270 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6271 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6273 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6274 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6275 using IsATensorArray = std::array<CtaGroupArray, 2>;
6276 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6277 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6280 static constexpr HasDisableOutputLaneArray tcgen05MMAIDs = {
6286 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic},
6288 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic}}},
6292 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6293 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6297 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
6298 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
6304 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic},
6306 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic}}},
6310 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6311 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6315 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
6316 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
6322 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1,
6325 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2,
6330 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1,
6332 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift,
6337 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2,
6339 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift,
6345 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1,
6349 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2,
6354 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1,
6356 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift},
6360 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2,
6362 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift,
6365 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6366 bool hasScaleInputD = ScaleInputD !=
nullptr;
6368 llvm::Value *DisableOutputLane =
6370 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6372 const unsigned ctaGroup =
6375 llvm::Intrinsic::ID ID =
6376 tcgen05MMAIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6377 [ctaGroup - 1][thisOp.getAShift()];
6379 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMAOp.");
6382 args.push_back(ScaleInputD);
6384 if (hasDisableOutputLane)
6385 args.push_back(DisableOutputLane);
6387 args.push_back(builder.getInt32(
6390 if (!hasDisableOutputLane)
6391 args.push_back(builder.getInt32(ctaGroup));
6394 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6397 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6404 NVVM::CTAGroupKind ctaGroup,
bool hasAShift,
6405 NVVM::Tcgen05MMACollectorOp collectorOp,
Location loc) {
6407 if (disableOutputLane) {
6408 mlir::VectorType disableOutputLaneType =
6409 cast<mlir::VectorType>(disableOutputLane.
getType());
6410 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6411 disableOutputLaneType.getNumElements() != 4) ||
6412 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6413 disableOutputLaneType.getNumElements() != 8))
6414 return emitError(loc) <<
"Disable Output Lane of length "
6415 << disableOutputLaneType.getNumElements()
6416 <<
" is incompatible with CtaGroupAttr";
6419 if (hasAShift && !isATensor)
6421 loc,
"A-shift can be applied only when matrix A is in tensor memory");
6423 if (hasAShift ==
true && (collectorOp == Tcgen05MMACollectorOp::FILL ||
6424 collectorOp == Tcgen05MMACollectorOp::USE))
6426 loc,
"Cannot use collector buffer operation fill or use with ashift");
6431LogicalResult Tcgen05MMAOp::verify() {
6433 getDisableOutputLane(), getCtaGroup(), getAShift(),
6434 getCollectorOp(), getLoc());
6444 auto thisOp = cast<NVVM::Tcgen05MMASparseOp>(op);
6447 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6450 bool isATensor = isa<llvm::PointerType>(
A->getType());
6453 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6454 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6455 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6456 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6458 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
6459 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
6460 using IsATensorArray = std::array<CtaGroupArray, 2>;
6461 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
6462 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
6465 static constexpr HasDisableOutputLaneArray tcgen05MMASparseIDs = {
6471 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic},
6473 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic}}},
6477 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6478 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6482 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
6483 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
6489 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6492 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
6497 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6498 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6502 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
6503 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
6510 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1,
6514 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2,
6519 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1,
6521 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift,
6526 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2,
6528 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift,
6534 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1,
6538 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2,
6543 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1,
6545 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift},
6549 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2,
6551 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift,
6554 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
6555 bool hasScaleInputD = ScaleInputD !=
nullptr;
6557 llvm::Value *DisableOutputLane =
6559 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6564 llvm::Intrinsic::ID ID =
6565 tcgen05MMASparseIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
6566 [ctaGroup - 1][thisOp.getAShift()];
6568 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMASparseOp.");
6571 args.push_back(ScaleInputD);
6573 if (hasDisableOutputLane)
6574 args.push_back(DisableOutputLane);
6576 args.push_back(builder.getInt32(
6579 if (!hasDisableOutputLane)
6580 args.push_back(builder.getInt32(ctaGroup));
6583 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6586 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6591LogicalResult Tcgen05MMASparseOp::verify() {
6593 getDisableOutputLane(), getCtaGroup(), getAShift(),
6594 getCollectorOp(), getLoc());
6604 auto thisOp = cast<NVVM::Tcgen05MMABlockScaleOp>(op);
6607 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6610 bool isATensor = isa<llvm::PointerType>(
A->getType());
6613 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6614 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6615 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6616 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6617 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6618 args.push_back(builder.getInt32(
6621 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6623 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6625 auto kind = thisOp.getKind();
6626 auto blockScale = thisOp.getBlockScale();
6627 llvm::Intrinsic::ID ID = [&]() {
6628 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6629 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6630 return isATensor ? llvm::Intrinsic::
6631 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale
6633 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale;
6634 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6637 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32
6639 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32;
6641 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6642 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6644 ? llvm::Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale
6645 : llvm::Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale;
6646 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6647 return isATensor ? llvm::Intrinsic::
6648 nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32
6650 nvvm_tcgen05_mma_shared_mxf4_block_scale_block32;
6652 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6653 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6656 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32
6658 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32;
6660 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6663 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16
6665 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16;
6668 llvm_unreachable(
"Invalid tcgen05.mma.block_scale attributes");
6675 NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind,
6676 NVVM::Tcgen05MMABlockScale blockScale,
Location loc) {
6677 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT &&
6678 kind == NVVM::Tcgen05MMAKind::MXF4NVF4)
6679 return emitError(loc,
"mxf4nvf4 requires block scale attribute");
6681 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16 &&
6682 kind != NVVM::Tcgen05MMAKind::MXF4NVF4)
6684 llvm::formatv(
"{} kind does not support block16 attribute",
6685 stringifyEnum(kind)));
6690LogicalResult Tcgen05MMABlockScaleOp::verify() {
6692 getBlockScale(), getLoc());
6702 auto thisOp = cast<NVVM::Tcgen05MMASparseBlockScaleOp>(op);
6705 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6708 bool isATensor = isa<llvm::PointerType>(
A->getType());
6711 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6712 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6713 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6714 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6715 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6716 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6717 args.push_back(builder.getInt32(
6720 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6722 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6724 auto kind = thisOp.getKind();
6725 auto blockScale = thisOp.getBlockScale();
6726 llvm::Intrinsic::ID ID = [&]() {
6727 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6728 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6729 return isATensor ? llvm::Intrinsic::
6730 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale
6732 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale;
6733 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6736 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32
6738 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32;
6740 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6741 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6742 return isATensor ? llvm::Intrinsic::
6743 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale
6745 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale;
6746 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6749 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32
6751 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32;
6753 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6754 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6757 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32
6759 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32;
6761 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6764 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16
6766 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16;
6769 llvm_unreachable(
"Invalid tcgen05.mma.sp.block_scale attributes");
6775LogicalResult Tcgen05MMASparseBlockScaleOp::verify() {
6777 getBlockScale(), getLoc());
6787 auto thisOp = cast<NVVM::Tcgen05MMAWsOp>(op);
6790 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6793 bool isATensor = isa<llvm::PointerType>(
A->getType());
6796 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6797 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6798 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6800 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6804 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor_zero_col_mask
6805 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared_zero_col_mask;
6807 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor
6808 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared;
6810 args.push_back(builder.getInt32(
6813 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6815 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6827 auto thisOp = cast<NVVM::Tcgen05MMAWsSparseOp>(op);
6830 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6833 bool isATensor = isa<llvm::PointerType>(
A->getType());
6836 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6837 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6838 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6839 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6841 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6846 ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor_zero_col_mask
6847 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared_zero_col_mask;
6849 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor
6850 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared;
6852 args.push_back(builder.getInt32(
6855 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6857 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6868 auto thisOp = cast<Tcgen05MMADecompressBOp>(op);
6871 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6874 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6877 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6878 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6879 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6880 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6882 llvm::Value *DisableOutputLane =
6884 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
6886 NVVM::CTAGroupKind ctaGroup = thisOp.getCtaGroup();
6888 using namespace llvm::Intrinsic;
6889 ID intrinsicID = not_intrinsic;
6891 if (hasDisableOutputLane) {
6892 if (ctaGroup == NVVM::CTAGroupKind::CTA_1) {
6895 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg1_decompress_b
6896 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg1_decompress_b;
6897 }
else if (ctaGroup == NVVM::CTAGroupKind::CTA_2) {
6900 ? nvvm_tcgen05_mma_tensor_f8f6f4_disable_output_lane_cg2_decompress_b
6901 : nvvm_tcgen05_mma_shared_f8f6f4_disable_output_lane_cg2_decompress_b;
6903 llvm_unreachable(
"Unknown ctaGroup for tcgen05.mma.decompress_b");
6906 intrinsicID = isATensor ? nvvm_tcgen05_mma_tensor_f8f6f4_decompress_b
6907 : nvvm_tcgen05_mma_shared_f8f6f4_decompress_b;
6910 assert(intrinsicID != not_intrinsic &&
6911 "Invalid intrinsic for Tcgen05MMADecompressBOp.");
6913 if (hasDisableOutputLane)
6914 args.push_back(DisableOutputLane);
6920 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6922 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6924 return {intrinsicID, args};
6927LogicalResult Tcgen05MMADecompressBOp::verify() {
6928 mlir::Value disableOutputLane = getDisableOutputLane();
6930 if (disableOutputLane) {
6931 NVVM::CTAGroupKind ctaGroup = getCtaGroup();
6933 mlir::VectorType disableOutputLaneType =
6934 cast<mlir::VectorType>(disableOutputLane.
getType());
6935 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
6936 disableOutputLaneType.getNumElements() != 4) ||
6937 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
6938 disableOutputLaneType.getNumElements() != 8))
6939 return emitOpError() <<
"Disable Output Lane of length "
6940 << disableOutputLaneType.getNumElements()
6941 <<
" is incompatible with CtaGroupAttr";
6953 auto thisOp = cast<Tcgen05MMABlockScaleDecompressBOp>(op);
6956 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6959 const bool isATensor = isa<llvm::PointerType>(
A->getType());
6962 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6963 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6964 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6965 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6966 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6967 args.push_back(mt.
lookupValue(thisOp.getDecompressBMetadata()));
6968 args.push_back(builder.getInt32(
6971 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpA())));
6973 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOpB())));
6975 llvm::Intrinsic::ID intrinsicID =
6978 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32_decompress_b
6980 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32_decompress_b;
6982 return {intrinsicID, args};
6989#define TCGEN05LDRED(SHAPE, NUM, TYPE) \
6990 llvm::Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_##NUM##_##TYPE
6994 auto thisOp = cast<NVVM::Tcgen05LdRedOp>(op);
6997 mlir::VectorType VecResTy =
6998 cast<mlir::VectorType>(thisOp.getData().getType());
6999 unsigned Num = VecResTy.getNumElements();
7000 bool IsFloat = thisOp.getRedVal().getType().isF32();
7002 llvm::Intrinsic::ID Shape32x32b[][2] = {
7013 llvm::Intrinsic::ID Shape16x32bx2[][2] = {
7024 NVVM::Tcgen05LdStShape
shape = thisOp.getShape();
7025 unsigned ID = [&]() {
7028 unsigned idx = std::log2(Num);
7030 case NVVM::Tcgen05LdStShape::SHAPE_32X32B:
7031 return Shape32x32b[idx][IsFloat];
7032 case NVVM::Tcgen05LdStShape::SHAPE_16X32BX2:
7033 return Shape16x32bx2[idx][IsFloat];
7035 llvm_unreachable(
"unhandled tcgen05.ld lowering");
7041 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2)
7042 args.push_back(mt.
lookupValue(thisOp.getOffset()));
7045 builder.getInt32(thisOp.getOp() == NVVM::ReductionKind::MIN ? 0 : 1));
7048 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getAbs())));
7049 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getNan())));
7054LogicalResult Tcgen05LdRedOp::verify() {
7055 VectorType data = cast<VectorType>(getData().
getType());
7056 Type redVal = getRedVal().getType();
7058 if (data.getElementType() != redVal)
7060 "type of reduction value and element type of vector data should match");
7062 if (getOp() != NVVM::ReductionKind::MIN &&
7063 getOp() != NVVM::ReductionKind::MAX)
7064 return emitError(
"only min and max reduction kinds are supported");
7066 if (redVal.
isInteger() && (getAbs() || getNan())) {
7067 return emitError(
"abs or nan is only applicable for f32 type");
7077struct NVVMInlinerInterface final : DialectInlinerInterface {
7078 using DialectInlinerInterface::DialectInlinerInterface;
7079 bool isLegalToInline(Operation *, Region *,
bool, IRMapping &)
const final {
7086void NVVMDialect::initialize() {
7087 registerNVVMDialectOperations(
this);
7089#define GET_ATTRDEF_LIST
7090#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
7095 allowUnknownOperations();
7096 addInterfaces<NVVMInlinerInterface>();
7097 declarePromisedInterface<ConvertToLLVMPatternInterface, NVVMDialect>();
7098 declarePromisedInterface<gpu::TargetAttrInterface, NVVMTargetAttr>();
7101LogicalResult NVVMDialect::verifyOperationAttribute(
Operation *op,
7103 StringAttr attrName = attr.
getName();
7105 if (attrName == NVVMDialect::getKernelFuncAttrName()) {
7106 if (!isa<LLVM::LLVMFuncOp>(op)) {
7107 return op->
emitError() <<
"'" << NVVMDialect::getKernelFuncAttrName()
7108 <<
"' attribute attached to unexpected op";
7113 if (attrName == NVVMDialect::getMaxntidAttrName() ||
7114 attrName == NVVMDialect::getReqntidAttrName() ||
7115 attrName == NVVMDialect::getClusterDimAttrName()) {
7116 auto values = llvm::dyn_cast<DenseI32ArrayAttr>(attr.
getValue());
7117 if (!values || values.empty() || values.size() > 3) {
7120 <<
"' attribute must be integer array with maximum 3 index";
7125 if (attrName == NVVMDialect::getMinctasmAttrName() ||
7126 attrName == NVVMDialect::getMaxnregAttrName() ||
7127 attrName == NVVMDialect::getClusterMaxBlocksAttrName()) {
7128 if (!llvm::dyn_cast<IntegerAttr>(attr.
getValue())) {
7130 <<
"'" << attrName <<
"' attribute must be integer constant";
7134 if (attrName == NVVMDialect::getBlocksAreClustersAttrName()) {
7138 <<
"'" << attrName <<
"' attribute must be used along with " <<
"'"
7139 << NVVMDialect::getReqntidAttrName() <<
"' and " <<
"'"
7140 << NVVMDialect::getClusterDimAttrName() <<
"'";
7147LogicalResult NVVMDialect::verifyRegionArgAttribute(
Operation *op,
7148 unsigned regionIndex,
7151 auto funcOp = dyn_cast<FunctionOpInterface>(op);
7156 StringAttr attrName = argAttr.
getName();
7157 if (attrName == NVVM::NVVMDialect::getGridConstantAttrName()) {
7161 <<
"' attribute must be present only on kernel arguments";
7163 if (!isa<UnitAttr>(argAttr.
getValue()))
7164 return op->
emitError() <<
"'" << attrName <<
"' must be a unit attribute";
7165 if (!funcOp.getArgAttr(argIndex, LLVM::LLVMDialect::getByValAttrName())) {
7168 <<
"' attribute requires the argument to also have attribute '"
7169 << LLVM::LLVMDialect::getByValAttrName() <<
"'";
7180unsigned NVVMMemorySpaceAttr::getAddressSpace()
const {
7181 return static_cast<unsigned>(getValue());
7184bool NVVMMemorySpaceAttr::isValidLoad(
7185 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
7186 const ::mlir::DataLayout *dataLayout,
7192bool NVVMMemorySpaceAttr::isValidStore(
7193 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
7194 const ::mlir::DataLayout *dataLayout,
7200bool NVVMMemorySpaceAttr::isValidAtomicOp(
7201 ptr::AtomicBinOp op,
Type type, ptr::AtomicOrdering ordering,
7202 std::optional<int64_t> alignment, const ::mlir::DataLayout *dataLayout,
7205 assert(
false &&
"unimplemented, see TODO in the source.");
7209bool NVVMMemorySpaceAttr::isValidAtomicXchg(
7210 Type type, ptr::AtomicOrdering successOrdering,
7211 ptr::AtomicOrdering failureOrdering, std::optional<int64_t> alignment,
7212 const ::mlir::DataLayout *dataLayout,
7215 assert(
false &&
"unimplemented, see TODO in the source.");
7219bool NVVMMemorySpaceAttr::isValidAddrSpaceCast(
7223 assert(
false &&
"unimplemented, see TODO in the source.");
7227bool NVVMMemorySpaceAttr::isValidPtrIntCast(
7232 assert(
false &&
"unimplemented, see TODO in the source.");
7241 int optLevel, StringRef triple, StringRef chip,
7242 StringRef features, DictionaryAttr flags,
7244 if (optLevel < 0 || optLevel > 3) {
7245 emitError() <<
"The optimization level must be a number between 0 and 3.";
7248 if (triple.empty()) {
7249 emitError() <<
"The target triple cannot be empty.";
7253 emitError() <<
"The target chip cannot be empty.";
7256 if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
7257 return mlir::isa_and_nonnull<StringAttr>(attr);
7259 emitError() <<
"All the elements in the `link` array must be strings.";
7265LogicalResult NVVMTargetAttr::verifyTarget(
Operation *gpuModule) {
7266 if (!getVerifyTarget())
7269 auto gpuModuleOp = llvm::dyn_cast<gpu::GPUModuleOp>(gpuModule);
7272 "NVVM target attribute must be attached to a GPU module");
7275 std::optional<unsigned> targetFullSmVersion =
7277 if (!targetFullSmVersion)
7279 <<
"invalid NVVM target chip \"" << getChip()
7280 <<
"\", expected sm_<version>[a|f]";
7284 "Minimum NVVM target SM version is sm_20");
7288 ->
walk([&](Operation *op) {
7289 if (
auto reqOp = llvm::dyn_cast<NVVM::RequiresSMInterface>(op)) {
7290 const NVVMCheckSMVersion requirement =
7291 reqOp.getRequiredMinSMVersion();
7293 op->
emitOpError() <<
"is not supported on " << getChip();
7305#define GET_ATTRDEF_CLASSES
7306#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)
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)
static llvm::Value * getAsPackedI32(llvm::Value *arg, llvm::IRBuilderBase &builder)
static void printMmaEnumProperty(OpAsmPrinter &printer, bool &isFirst, StringRef keyword, AttrTy value)
static LogicalResult verifyMBarrierArriveLikeOp(Operation *op, Value addr, NVVM::MemScopeKind scope, Value retVal=nullptr)
static llvm::Value * castPtrToAddrSpace(llvm::IRBuilderBase &builder, llvm::Value *ptr, NVVMMemorySpace targetAS)
static LogicalResult isAllowedWGMMADataType(NVVM::WGMMATypes typeD, NVVM::WGMMATypes typeA, NVVM::WGMMATypes typeB)
static llvm::Intrinsic::ID getBarrierReductionIntrinsic(bool aligned, NVVM::BarrierReduction kind)
Maps the (aligned, kind) pair to the @llvm.nvvm.barrier.cta.red.
static ParseResult parseMmaProperties(OpAsmParser &parser, NamedAttrList &attributes, ArrayRef< StringRef > allowedKeywords, ArrayRef< StringRef > requiredProperties)
static void inferAndSetMultiplicandTypes(MLIRContext *ctx, NamedAttrList &attrs, const SmallVectorImpl< Type > &operandTypes)
static LogicalResult parseMmaOperand(OpAsmParser &parser, StringRef operandName, SmallVectorImpl< OpAsmParser::UnresolvedOperand > ®s)
static std::pair< mlir::Type, unsigned > inferMMATypeFromMNK(NVVM::MMATypes type, NVVM::MMAFrag frag, int m, int n, int k, MLIRContext *context)
static bool isInt8PtxType(MMATypes type)
#define TCGEN05LDRED(SHAPE, NUM, TYPE)
static bool isInt4PtxType(MMATypes type)
static bool isIntegerPtxType(MMATypes type)
static MMATypes inferPtxTypeFromResult(OpTy op)
static LogicalResult parseMmaTypeSignature(OpAsmParser &parser, SmallVectorImpl< Type > &operandTypes)
static FailureOr< int > getAllowedSizeK(NVVM::WGMMATypes typeA)
static bool isPtrInSharedClusterSpace(mlir::Value ptr)
LogicalResult CpAsyncBulkTensorOverrideAddrCommonVerifier(OperandRange coordinates, OperandRange tensorSize, OperandRange lowerStride, Value upperStride, bool isTile, Location loc)
static ParseResult parseMmaEnumPropertyValue(OpAsmParser &parser, NamedAttrList &attributes, StringRef name)
#define GET_CP_ASYNC_ID(mod, size, has_cpsize)
static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape shape, unsigned vecLen)
static LogicalResult 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(OperandRange operands)
Print a comma separated range of operation operands out of line to avoid instantiating the range iter...
virtual void printOptionalAttrDict(ArrayRef< NamedAttribute > attrs, ArrayRef< StringRef > elidedAttrs={})=0
If the specified operation has attributes, print out an attribute dictionary with their values.
This class helps build Operations.
This provides public APIs that all operations should have.
This class implements the operand iterators for the Operation class.
Operation is the basic unit of execution within MLIR.
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.