32#include "llvm/ADT/STLExtras.h"
33#include "llvm/ADT/TypeSwitch.h"
34#include "llvm/IR/IRBuilder.h"
35#include "llvm/IR/NVVMIntrinsicUtils.h"
36#include "llvm/Support/Casting.h"
37#include "llvm/Support/FormatVariadic.h"
38#include "llvm/Support/NVPTXAddrSpace.h"
39#include "llvm/Support/raw_ostream.h"
47#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.cpp.inc"
48#include "mlir/Dialect/LLVMIR/NVVMOpsEnums.cpp.inc"
50static constexpr unsigned notIntrinsic = llvm::Intrinsic::not_intrinsic;
57 auto ptrTy = llvm::cast<LLVM::LLVMPointerType>(
ptr.getType());
58 return ptrTy.getAddressSpace() ==
static_cast<unsigned>(targetAS);
75 NVVMMemorySpace targetAS) {
76 unsigned AS =
static_cast<unsigned>(targetAS);
77 return builder.CreateAddrSpaceCast(
78 ptr, llvm::PointerType::get(builder.getContext(), AS));
82static llvm::nvvm::CTAGroupKind
85 case NVVM::CTAGroupKind::CTA_1:
86 return llvm::nvvm::CTAGroupKind::CG_1;
87 case NVVM::CTAGroupKind::CTA_2:
88 return llvm::nvvm::CTAGroupKind::CG_2;
90 llvm_unreachable(
"unsupported cta_group value");
102 size_t numIm2ColOffsets,
104 if (tensorDims < 1 || tensorDims > 5)
105 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
113 "to use im2col mode, the tensor has to be at least 3-dimensional");
115 if (numIm2ColOffsets && (tensorDims != (numIm2ColOffsets + 2)))
117 loc,
"im2col offsets must be 2 less than number of coordinates");
122LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
123 TMAStoreMode mode = getMode();
127 if (getPredicate()) {
128 if (mode != TMAStoreMode::TILE)
129 return emitError(
"Inline-ptx lowering supported only for Tile mode.");
130 if (getL2CacheHint())
131 return emitError(
"Inline-ptx lowering unsupported with L2 cache-hint.");
136 case TMAStoreMode::TILE:
138 case TMAStoreMode::IM2COL:
140 case TMAStoreMode::TILE_SCATTER4:
142 return emitError(
"Scatter4 mode expects 5 coordinates");
147LogicalResult CpAsyncOp::verify() {
148 if (getModifier() != LoadCacheModifierKind::CG &&
149 getModifier() != LoadCacheModifierKind::CA)
150 return emitError(
"Only CG and CA cache modifiers are supported.");
151 if (getSize() != 4 && getSize() != 8 && getSize() != 16)
152 return emitError(
"expected byte size to be either 4, 8 or 16.");
153 if (getModifier() == LoadCacheModifierKind::CG && getSize() != 16)
154 return emitError(
"CG cache modifier is only support for 16 bytes copy.");
161 if (tensorDims < 1 || tensorDims > 5)
162 return emitError(loc,
"expects coordinates between 1 to 5 dimension");
164 auto checkTMALoadParams = [&](TMALoadMode mode,
bool isIm2col,
165 size_t expectedIm2colOff) -> LogicalResult {
166 if (isIm2col && (tensorDims < 3))
169 <<
" mode, the tensor has to be at least 3-dimensional";
171 if (numIm2colOff != expectedIm2colOff)
172 return emitError(loc) <<
" im2col offsets expected " << expectedIm2colOff
173 <<
" (provided " << numIm2colOff <<
")";
179 case TMALoadMode::TILE:
180 return checkTMALoadParams(mode,
false, 0);
181 case TMALoadMode::IM2COL:
182 return checkTMALoadParams(mode,
true, tensorDims - 2);
183 case TMALoadMode::IM2COL_W:
184 case TMALoadMode::IM2COL_W_128:
185 return checkTMALoadParams(mode,
true, 2);
186 case TMALoadMode::TILE_GATHER4:
187 return (tensorDims == 5)
188 ? checkTMALoadParams(mode,
false, 0)
189 :
emitError(loc,
"Gather4 mode expects 5 coordinates");
194LogicalResult CpAsyncBulkTensorPrefetchOp::verify() {
196 getMode(), getLoc());
199LogicalResult CpAsyncBulkTensorGlobalToSharedClusterOp::verify() {
200 TMALoadMode mode = getMode();
201 bool isCTAOnly = getIsCTAOnly();
202 if (getPredicate()) {
204 return emitError(
"Predicate is supported only for shared::cluster mode.");
205 if (mode != TMALoadMode::TILE && mode != TMALoadMode::IM2COL)
207 "Predicate is supported only for Tile and Im2col modes.");
209 NVVMMemorySpace expectedAS =
210 isCTAOnly ? NVVMMemorySpace::Shared : NVVMMemorySpace::SharedCluster;
211 unsigned AS = llvm::cast<LLVM::LLVMPointerType>(getDstMem().
getType())
213 if (AS != expectedAS)
216 ?
"Shared::cta destination requires address-space 3."
217 :
"Shared::cluster destination requires address-space 7.");
220 if (getMulticastMask())
221 return emitError(
"Multicast is not supported with shared::cta mode.");
223 return emitError(
"CTAGroup is not supported with shared::cta mode.");
228 getMode(), getLoc());
231LogicalResult CpAsyncBulkTensorReduceOp::verify() {
232 TMAStoreMode mode = getMode();
235 case TMAStoreMode::TILE:
237 case TMAStoreMode::IM2COL:
239 case TMAStoreMode::TILE_SCATTER4:
240 return emitError(
"Scatter mode unsupported for CpAsyncBulkTensorReduceOp");
245LogicalResult CpAsyncBulkGlobalToSharedClusterOp::verify() {
247 if (isSharedCTA && getMulticastMask())
248 return emitError(
"Multicast is not supported with shared::cta mode.");
254 NVVM::MemScopeKind scope,
255 Value retVal =
nullptr) {
256 if (scope != NVVM::MemScopeKind::CTA && scope != NVVM::MemScopeKind::CLUSTER)
257 return op->
emitError(
"mbarrier scope must be either CTA or Cluster");
260 bool hasRetValue =
static_cast<bool>(retVal);
261 if (isSharedCluster && hasRetValue)
263 "mbarrier in shared_cluster space cannot return any value");
268LogicalResult MBarrierArriveOp::verify() {
273LogicalResult MBarrierArriveDropOp::verify() {
278LogicalResult MBarrierArriveExpectTxOp::verify() {
282 if (getPredicate()) {
283 if (getScope() != NVVM::MemScopeKind::CTA)
284 return emitError(
"mbarrier scope must be CTA when using predicate");
287 return emitError(
"mbarrier in shared_cluster space is not supported when "
291 return emitError(
"return-value is not supported when using predicate");
293 if (getRelaxed() ==
true)
294 return emitError(
"mbarrier with relaxed semantics is not supported when "
301LogicalResult MBarrierArriveDropExpectTxOp::verify() {
316 inferredReturnTypes.push_back(IntegerType::get(context, 64));
321MBarrierArriveOp::inferReturnTypes(
MLIRContext *context,
322 std::optional<Location> location,
323 MBarrierArriveOp::Adaptor adaptor,
326 inferredReturnTypes);
329LogicalResult MBarrierArriveDropOp::inferReturnTypes(
330 MLIRContext *context, std::optional<Location> location,
331 MBarrierArriveDropOp::Adaptor adaptor,
334 inferredReturnTypes);
337LogicalResult MBarrierArriveExpectTxOp::inferReturnTypes(
338 MLIRContext *context, std::optional<Location> location,
339 MBarrierArriveExpectTxOp::Adaptor adaptor,
343 if (adaptor.getPredicate())
346 inferredReturnTypes);
349LogicalResult MBarrierArriveDropExpectTxOp::inferReturnTypes(
350 MLIRContext *context, std::optional<Location> location,
351 MBarrierArriveDropExpectTxOp::Adaptor adaptor,
354 inferredReturnTypes);
364 return inferred == actual;
373bool MBarrierArriveExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
377bool MBarrierArriveDropExpectTxOp::isCompatibleReturnTypes(
TypeRange l,
382LogicalResult MBarrierExpectTxOp::verify() {
386LogicalResult MBarrierCompleteTxOp::verify() {
390LogicalResult MBarrierTestWaitOp::verify() {
394LogicalResult MBarrierTryWaitOp::verify() {
398LogicalResult ConvertFloatToTF32Op::verify() {
399 using RndMode = NVVM::FPRoundingMode;
403 return emitError(
"Relu not supported with rna rounding mode.");
410 "Only {rn,rz,rna} rounding modes supported for ConvertFloatToTF32Op.");
415LogicalResult ConvertF32x2ToF6x2Op::verify() {
418 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy())) {
420 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
421 << mlir::Float6E3M2FNType::get(ctx)
422 <<
" types are supported for conversions from f32x2 to f6x2.";
427LogicalResult ConvertF32x2ToF8x2Op::verify() {
428 using RndMode = NVVM::FPRoundingMode;
429 using SatMode = NVVM::SaturationMode;
431 bool isRoundingModeRN = getRnd() == RndMode::RN;
432 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
433 bool isRoundingModeRP = getRnd() == RndMode::RP;
434 bool isSatFinite = getSat() == SatMode::SATFINITE;
436 bool hasRelu = getRelu();
441 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
443 if (!isRoundingModeRN) {
444 return emitOpError(
"Only RN rounding mode is supported for "
445 "conversions from f32x2 to ")
446 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
447 << mlir::Float8E5M2Type::get(ctx) <<
" types";
450 return emitOpError(
"Only SATFINITE saturation mode is supported "
453 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
454 << mlir::Float8E5M2Type::get(ctx) <<
" types";
458 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
459 if (!(isRoundingModeRZ || isRoundingModeRP)) {
460 return emitOpError(
"Only RZ and RP rounding modes are supported for "
461 "conversions from f32x2 to ")
462 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
465 return emitOpError(
"relu not supported for conversions to ")
466 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
472 << mlir::Float8E4M3FNType::get(ctx) <<
", "
473 << mlir::Float8E5M2Type::get(ctx) <<
", and "
474 << mlir::Float8E8M0FNUType::get(ctx)
476 "supported for conversions from f32x2 to f8x2";
480LogicalResult ConvertF16x2ToF8x2Op::verify() {
483 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy())) {
485 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
486 << mlir::Float8E5M2Type::get(ctx)
487 <<
" types are supported for conversions from f16x2 to f8x2.";
492LogicalResult ConvertBF16x2ToF8x2Op::verify() {
493 using RndMode = NVVM::FPRoundingMode;
494 using SatMode = NVVM::SaturationMode;
496 bool isRoundingModeRN = getRnd() == RndMode::RN;
497 bool isRoundingModeRZ = getRnd() == RndMode::RZ;
498 bool isRoundingModeRP = getRnd() == RndMode::RP;
499 bool isSatFinite = getSat() == SatMode::SATFINITE;
500 bool hasRelu = getRelu();
505 .Case<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(
507 if (!isRoundingModeRN)
508 return emitOpError(
"Only RN rounding mode is supported for "
509 "conversions from bf16x2 to ")
510 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
511 << mlir::Float8E5M2Type::get(ctx) <<
" types";
513 return emitOpError(
"Only SATFINITE saturation mode is supported "
514 "for conversions from bf16x2 to ")
515 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
516 << mlir::Float8E5M2Type::get(ctx) <<
" types";
519 .Case<mlir::Float8E8M0FNUType>([&](
mlir::Type) -> LogicalResult {
520 if (!(isRoundingModeRZ || isRoundingModeRP))
521 return emitOpError(
"Only RZ and RP rounding modes are supported for "
522 "conversions from bf16x2 to ")
523 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
525 return emitOpError(
"relu not supported for conversions to ")
526 << mlir::Float8E8M0FNUType::get(ctx) <<
" type";
530 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
535LogicalResult ConvertF32x2ToF4x2Op::verify() {
538 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
540 << mlir::Float4E2M1FNType::get(ctx)
541 <<
" type is supported for conversions from f32x2 to f4x2.";
546LogicalResult ConvertF8x2ToBF16x2Op::verify() {
548 if (llvm::isa<Float8E8M0FNUType>(getSrcType())) {
549 if (getSat() != SaturationMode::NONE)
551 "Only NONE saturation mode is supported for conversions from ")
552 << Float8E8M0FNUType::get(ctx) <<
" type";
553 if (getScaleFactor())
554 return emitOpError(
"scaleFactor not supported for conversions from ")
555 << Float8E8M0FNUType::get(ctx) <<
" type";
557 return emitOpError(
"relu not supported for conversions from ")
558 << Float8E8M0FNUType::get(ctx) <<
" type";
564LogicalResult PermuteOp::verify() {
565 using Mode = NVVM::PermuteMode;
566 bool hasHi =
static_cast<bool>(getHi());
573 return emitError(
"mode '") << getMode() <<
"' requires 'hi' operand.";
581 << getMode() <<
"' does not accept 'hi' operand.";
596 static constexpr FPRoundingMode validRndModes[] = {
597 FPRoundingMode::RN, FPRoundingMode::RZ, FPRoundingMode::RS};
599 if (!llvm::is_contained(validRndModes, rnd)) {
601 "Only RN, RZ, and RS rounding modes are supported for "
602 "conversions from f32x2 to ")
606 if (rnd == FPRoundingMode::RS) {
607 if (!hasRandomBits) {
608 return op->
emitOpError(
"random_bits is required for RS rounding mode.");
613 "random_bits not supported for RN and RZ rounding modes.");
620LogicalResult ConvertF32x2ToF16x2Op::verify() {
622 getRandomBits() ?
true :
false, *
this);
625LogicalResult ConvertF32x2ToBF16x2Op::verify() {
627 getRandomBits() ?
true :
false, *
this);
630LogicalResult ConvertF32x4ToF8x4Op::verify() {
633 if (!llvm::isa<mlir::Float8E4M3FNType, mlir::Float8E5M2Type>(getDstTy()))
635 << mlir::Float8E4M3FNType::get(ctx) <<
" and "
636 << mlir::Float8E5M2Type::get(ctx)
637 <<
" types are supported for conversions from f32x4 to f8x4.";
642LogicalResult ConvertF32x4ToF6x4Op::verify() {
645 if (!llvm::isa<mlir::Float6E2M3FNType, mlir::Float6E3M2FNType>(getDstTy()))
647 << mlir::Float6E2M3FNType::get(ctx) <<
" and "
648 << mlir::Float6E3M2FNType::get(ctx)
649 <<
" types are supported for conversions from f32x4 to f6x4.";
654LogicalResult ConvertF32x4ToF4x4Op::verify() {
657 if (!llvm::isa<mlir::Float4E2M1FNType>(getDstTy()))
658 return emitOpError(
"Only ") << mlir::Float4E2M1FNType::get(ctx)
659 <<
" type is supported for conversions from "
665LogicalResult BulkStoreOp::verify() {
666 if (getInitVal() != 0)
667 return emitOpError(
"only 0 is supported for initVal, got ") << getInitVal();
671LogicalResult PMEventOp::verify() {
672 auto eventId = getEventId();
673 auto maskedEventId = getMaskedEventId();
674 if (!maskedEventId && !eventId) {
675 return emitOpError() <<
"either `id` or `mask` must be set";
678 if (maskedEventId && eventId) {
679 return emitOpError() <<
"`id` and `mask` cannot be set at the same time";
683 if (eventId < 0 || eventId > 15) {
684 return emitOpError() <<
"`id` must be between 0 and 15";
688 return llvm::success();
694std::optional<mlir::NVVM::MMATypes>
695MmaOp::inferOperandMMAType(
Type operandElType,
bool isAccumulator) {
697 VectorType::get(2, Float16Type::get(operandElType.
getContext()));
698 if (operandElType.
isF64())
699 return NVVM::MMATypes::f64;
700 if (operandElType.
isF16() || operandElType == half2Type)
701 return NVVM::MMATypes::f16;
702 if (operandElType.
isF32() && isAccumulator)
703 return NVVM::MMATypes::f32;
704 if (operandElType.
isF32() && !isAccumulator)
705 return NVVM::MMATypes::tf32;
706 if (llvm::isa<IntegerType>(operandElType)) {
708 return NVVM::MMATypes::s32;
712 if (
auto structType = llvm::dyn_cast<LLVM::LLVMStructType>(operandElType)) {
713 if (structType.getBody().empty())
715 return inferOperandMMAType(structType.getBody()[0], isAccumulator);
722 return (type == MMATypes::u4 || type == MMATypes::s4);
726 return (type == MMATypes::u8 || type == MMATypes::s8);
731 type == MMATypes::s32;
734MMATypes MmaOp::accumPtxType() {
735 std::optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
736 getODSOperands(2).getTypes().front(),
true);
737 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
741MMATypes MmaOp::resultPtxType() {
742 std::optional<mlir::NVVM::MMATypes> val =
743 inferOperandMMAType(getResult().
getType(),
true);
744 assert(val.has_value() &&
"result PTX type should always be inferrable");
750 struct MMAOperandFragment {
751 StringRef operandName;
752 StringRef ptxTypeAttr;
753 SmallVector<Value, 4> regs;
754 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
755 : operandName(name), ptxTypeAttr(ptxTypeName) {}
758 std::array<MMAOperandFragment, 3> frags{
759 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
760 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
761 MMAOperandFragment(
"C",
"")};
763 mlir::NVVM::MmaOp::getOperandSegmentSizeAttr()};
765 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
766 auto &frag = frags[fragIdx];
767 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
768 for (
auto operandIdx = varOperandSpec.first;
769 operandIdx < varOperandSpec.first + varOperandSpec.second;
771 frag.regs.push_back(this->getOperand(operandIdx));
772 if (operandIdx == 0) {
773 regTypes.push_back(this->getOperand(operandIdx).
getType());
776 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
777 regTypes.back(), fragIdx >= 2);
779 ignoreAttrNames.push_back(frag.ptxTypeAttr);
782 auto printMmaOperand = [&](
const MMAOperandFragment &frag) ->
void {
783 p <<
" " << frag.operandName;
789 for (
const auto &frag : frags) {
790 printMmaOperand(frag);
798 frags[1].regs[0].getType(),
799 frags[2].regs[0].getType()},
808 std::optional<MMAIntOverflow> intOverflow,
809 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
810 std::optional<std::array<MMALayout, 2>> multiplicandLayouts) {
812 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
817 result.addOperands(operandA);
818 result.addOperands(operandB);
819 result.addOperands(operandC);
821 if (multiplicandPtxTypes) {
822 result.addAttribute(
"multiplicandAPtxType",
823 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
824 result.addAttribute(
"multiplicandBPtxType",
825 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
827 if (
auto res = inferOperandMMAType(operandA[0].
getType(),
false))
828 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
829 if (
auto res = inferOperandMMAType(operandB[0].
getType(),
false))
830 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
833 if (multiplicandLayouts) {
834 result.addAttribute(
"layoutA",
835 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[0]));
836 result.addAttribute(
"layoutB",
837 MMALayoutAttr::get(ctx, (*multiplicandLayouts)[1]));
839 result.addAttribute(
"layoutA", MMALayoutAttr::get(ctx, MMALayout::row));
840 result.addAttribute(
"layoutB", MMALayoutAttr::get(ctx, MMALayout::col));
843 if (intOverflow.has_value())
844 result.addAttribute(
"intOverflowBehavior",
845 MMAIntOverflowAttr::get(ctx, *intOverflow));
846 if (b1Op.has_value())
847 result.addAttribute(
"b1Op", MMAB1OpAttr::get(ctx, *b1Op));
849 result.addTypes(resultType);
851 MmaOp::getOperandSegmentSizeAttr(),
853 static_cast<int32_t>(operandB.size()),
854 static_cast<int32_t>(operandC.size())}));
862 struct MMAOperandFragment {
863 std::optional<MMATypes> elemtype;
864 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
865 SmallVector<Type> regTypes;
869 std::array<MMAOperandFragment, 4> frags;
875 MMAOperandFragment &frag) -> LogicalResult {
905 if (operandTypes.size() != 3)
908 "expected one type for each operand segment but got " +
909 Twine(operandTypes.size()) +
" types");
910 for (
const auto &iter : llvm::enumerate(operandTypes)) {
911 auto &frag = frags[iter.index()];
912 frag.regTypes.resize(frag.regs.size(), iter.value());
916 frag.elemtype = inferOperandMMAType(frag.regTypes[0],
923 frags[3].elemtype = inferOperandMMAType(resultType,
true);
925 std::array<StringRef, 2> names{
"multiplicandAPtxType",
926 "multiplicandBPtxType"};
927 for (
unsigned idx = 0; idx < names.size(); idx++) {
928 const auto &frag = frags[idx];
929 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
930 if (!frag.elemtype.has_value() && !attr.has_value()) {
933 "attribute " + names[idx] +
934 " is not provided explicitly and cannot be inferred");
936 if (!attr.has_value())
938 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
941 result.addTypes(resultType);
942 if (!namedAttributes.
empty())
943 result.addAttributes(namedAttributes);
944 result.addAttribute(MmaOp::getOperandSegmentSizeAttr(),
946 static_cast<int32_t>(frags[0].regs.size()),
947 static_cast<int32_t>(frags[1].regs.size()),
948 static_cast<int32_t>(frags[2].regs.size()),
953LogicalResult MmaOp::verify() {
955 auto f16Ty = Float16Type::get(context);
956 auto i32Ty = IntegerType::get(context, 32);
957 auto f16x2Ty = VectorType::get(2, f16Ty);
958 auto f32Ty = Float32Type::get(context);
959 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
960 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
963 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
966 auto f16x2x2StructTy =
967 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
969 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
971 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
973 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
974 getShapeAttr().getK()};
980 AllowedShapes allowedShapes;
981 AllowedTypes expectedA;
982 AllowedTypes expectedB;
983 AllowedTypes expectedC;
988 if (mmaShape[0] == 16) {
990 Type multiplicandFragType;
991 switch (*getMultiplicandAPtxType()) {
994 multiplicandFragType = i32Ty;
995 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
996 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1000 multiplicandFragType = i32Ty;
1001 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1002 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1006 multiplicandFragType = f16x2Ty;
1007 expectedResult.push_back(f16x2x2StructTy);
1008 expectedResult.push_back(f32x4StructTy);
1010 case MMATypes::e4m3:
1011 case MMATypes::e5m2:
1015 multiplicandFragType = i32Ty;
1016 expectedResult.push_back(f16x2x2StructTy);
1017 expectedResult.push_back(f32x4StructTy);
1031 return emitError(
"invalid shape or multiplicand type: ")
1032 << getMultiplicandAPtxType().value();
1036 expectedResult.push_back(s32x4StructTy);
1037 expectedC.emplace_back(4, i32Ty);
1038 multiplicandFragType = i32Ty;
1040 expectedC.emplace_back(2, f16x2Ty);
1041 expectedC.emplace_back(4, f32Ty);
1044 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor);
1045 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1046 expectedA.emplace_back(unitA, multiplicandFragType);
1047 expectedB.emplace_back(unitB, multiplicandFragType);
1048 allowedShapes.push_back({16, 8, kFactor});
1049 allowedShapes.push_back({16, 8, kFactor * 2});
1051 if (resultPtxType() != accumPtxType())
1056 if (mmaShape[0] == 8) {
1057 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1058 expectedA.emplace_back(2, f16x2Ty);
1059 expectedB.emplace_back(2, f16x2Ty);
1060 expectedResult.push_back(f16x2x4StructTy);
1061 expectedResult.push_back(f32x8StructTy);
1062 expectedC.emplace_back(4, f16x2Ty);
1063 expectedC.emplace_back(8, f32Ty);
1064 allowedShapes.push_back({8, 8, 4});
1066 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1067 Type f64Ty = Float64Type::get(context);
1068 expectedA.emplace_back(1, f64Ty);
1069 expectedB.emplace_back(1, f64Ty);
1070 expectedC.emplace_back(2, f64Ty);
1071 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1073 allowedShapes.push_back({8, 8, 4});
1076 expectedA.push_back({i32Ty});
1077 expectedB.push_back({i32Ty});
1078 expectedC.push_back({i32Ty, i32Ty});
1079 expectedResult.push_back(s32x2StructTy);
1081 allowedShapes.push_back({8, 8, 32});
1083 allowedShapes.push_back({8, 8, 16});
1084 if (getMultiplicandAPtxType().value() == MMATypes::b1)
1085 allowedShapes.push_back({8, 8, 128});
1089 std::string errorMessage;
1090 llvm::raw_string_ostream errorStream(errorMessage);
1093 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1094 !llvm::is_contained(allowedShapes, mmaShape)) {
1095 errorStream <<
"unimplemented variant for MMA shape <";
1096 llvm::interleaveComma(mmaShape, errorStream);
1102 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1103 for (
const auto &iter : llvm::enumerate(
1105 auto spec = this->getODSOperandIndexAndLength(iter.index());
1107 operand_type_begin() + spec.first +
1109 bool match = llvm::is_contained(iter.value(), operandTySeg);
1112 errorStream <<
"Could not match types for the "
1113 << operandNames[iter.index()]
1114 <<
" operands; expected one of ";
1115 for (
const auto &x : iter.value()) {
1116 errorStream << x.size() <<
"x" << x[0] <<
" ";
1118 errorStream <<
"but got ";
1119 llvm::interleaveComma(operandTySeg, errorStream);
1125 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1126 return expectedResultType == getResult().getType();
1129 <<
"Could not match allowed types for the result; expected one of ";
1130 llvm::interleaveComma(expectedResult, errorStream);
1131 errorStream <<
" but got " << getResult().getType();
1136 if (getMultiplicandAPtxType() == MMATypes::b1 && !getB1Op()) {
1137 return emitOpError(
"op requires " + getB1OpAttrName().strref() +
1145 if (!getIntOverflowBehavior())
1147 getIntOverflowBehaviorAttrName().strref() +
1155 (mmaShape[0] == 8 && mmaShape[1] == 8 && mmaShape[2] == 4 &&
1156 getMultiplicandAPtxType() == MMATypes::f16);
1158 if (!isM8N8K4_F16) {
1160 if (getLayoutA() != MMALayout::row || getLayoutB() != MMALayout::col) {
1161 return emitOpError(
"requires layoutA = #nvvm.mma_layout<row> and "
1162 "layoutB = #nvvm.mma_layout<col> for shape <")
1163 << mmaShape[0] <<
", " << mmaShape[1] <<
", " << mmaShape[2]
1164 <<
"> with element types " << *getMultiplicandAPtxType() <<
" and "
1165 << *getMultiplicandBPtxType()
1166 <<
". Only m8n8k4 with f16 supports other layouts.";
1173MMATypes MmaSpOp::accumPtxType() {
1174 std::optional<mlir::NVVM::MMATypes> val = MmaOp::inferOperandMMAType(
1175 getODSOperands(2).getTypes().front(),
true);
1176 assert(val.has_value() &&
"accumulator PTX type should always be inferrable");
1180MMATypes MmaSpOp::resultPtxType() {
1181 std::optional<mlir::NVVM::MMATypes> val =
1182 MmaOp::inferOperandMMAType(getResult().
getType(),
true);
1183 assert(val.has_value() &&
"result PTX type should always be inferrable");
1189 llvm::IRBuilderBase &builder) {
1190 auto thisOp = cast<NVVM::MmaSpOp>(op);
1198 auto intId = MmaSpOp::getIntrinsicID(
1199 thisOp.getShape().getM(), thisOp.getShape().getN(),
1200 thisOp.getShape().getK(), thisOp.getIntOverflowBehavior(),
1201 thisOp.getOrderedMetadata(), thisOp.getKind(),
1202 *thisOp.getMultiplicandAPtxType(), *thisOp.getMultiplicandBPtxType(),
1203 thisOp.accumPtxType(), thisOp.resultPtxType());
1205 return {intId, args};
1210 struct MMAOperandFragment {
1211 StringRef operandName;
1212 StringRef ptxTypeAttr;
1213 SmallVector<Value, 4> regs;
1214 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1215 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1218 std::array<MMAOperandFragment, 5> frags{
1219 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1220 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1221 MMAOperandFragment(
"C",
""), MMAOperandFragment(
"sparseMetadata",
""),
1222 MMAOperandFragment(
"selector",
"")};
1224 mlir::NVVM::MmaSpOp::getOperandSegmentSizeAttr()};
1227 for (
unsigned fragIdx = 0; fragIdx < 3; fragIdx++) {
1228 auto &frag = frags[fragIdx];
1229 auto varOperandSpec = getODSOperandIndexAndLength(fragIdx);
1230 for (
auto operandIdx = varOperandSpec.first;
1231 operandIdx < varOperandSpec.first + varOperandSpec.second;
1233 frag.regs.push_back(this->getOperand(operandIdx));
1234 if (operandIdx == varOperandSpec.first) {
1235 regTypes.push_back(this->getOperand(operandIdx).
getType());
1238 std::optional<MMATypes> inferredType = MmaOp::inferOperandMMAType(
1239 regTypes.back(), fragIdx >= 2);
1241 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1245 frags[3].regs.push_back(getSparseMetadata());
1246 frags[4].regs.push_back(getSparsitySelector());
1248 auto printMmaSpOperand = [&](
const MMAOperandFragment &frag) ->
void {
1249 p <<
" " << frag.operandName;
1255 for (
const auto &frag : frags)
1256 printMmaSpOperand(frag);
1261 for (
int i = 0; i < 3; ++i) {
1266 p <<
") -> " << getResult().getType();
1273 std::optional<MMAIntOverflow> intOverflow,
1274 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
1276 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1281 result.addOperands(operandA);
1282 result.addOperands(operandB);
1283 result.addOperands(operandC);
1284 result.addOperands(sparseMetadata);
1285 result.addOperands(sparsitySelector);
1287 if (multiplicandPtxTypes) {
1288 result.addAttribute(
"multiplicandAPtxType",
1289 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1290 result.addAttribute(
"multiplicandBPtxType",
1291 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1293 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
1294 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1295 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
1296 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1299 if (intOverflow.has_value())
1300 result.addAttribute(
"intOverflowBehavior",
1301 MMAIntOverflowAttr::get(ctx, *intOverflow));
1303 result.addTypes(resultType);
1305 MmaSpOp::getOperandSegmentSizeAttr(),
1307 static_cast<int32_t>(operandB.size()),
1308 static_cast<int32_t>(operandC.size()), 1,
1313 struct MMAOperandFragment {
1314 std::optional<MMATypes> elemtype;
1315 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1316 SmallVector<Type> regTypes;
1320 std::array<MMAOperandFragment, 6> frags;
1325 auto parseMmaSpOperand = [&](StringRef operandName,
1326 MMAOperandFragment &frag) -> LogicalResult {
1337 if (parseMmaSpOperand(
"A", frags[0]).
failed())
1339 if (parseMmaSpOperand(
"B", frags[1]).
failed())
1341 if (parseMmaSpOperand(
"C", frags[2]).
failed())
1343 if (parseMmaSpOperand(
"sparseMetadata", frags[3]).
failed())
1345 if (parseMmaSpOperand(
"selector", frags[4]).
failed())
1361 if (operandTypes.size() != 3)
1364 "expected one type for each operand segment but got " +
1365 Twine(operandTypes.size()) +
" types");
1366 for (
const auto &iter : llvm::enumerate(operandTypes)) {
1367 auto &frag = frags[iter.index()];
1368 frag.regTypes.resize(frag.regs.size(), iter.value());
1373 MmaOp::inferOperandMMAType(frag.regTypes[0],
1381 MmaOp::inferOperandMMAType(resultType,
true);
1396 std::array<StringRef, 2> names{
"multiplicandAPtxType",
1397 "multiplicandBPtxType"};
1398 for (
unsigned idx = 0; idx < names.size(); idx++) {
1399 const auto &frag = frags[idx];
1400 std::optional<NamedAttribute> attr = namedAttributes.
getNamed(names[idx]);
1401 if (!frag.elemtype.has_value() && !attr.has_value()) {
1404 "attribute " + names[idx] +
1405 " is not provided explicitly and cannot be inferred");
1407 if (!attr.has_value())
1409 names[idx], MMATypesAttr::get(parser.
getContext(), *frag.elemtype));
1412 result.addTypes(resultType);
1413 if (!namedAttributes.
empty())
1414 result.addAttributes(namedAttributes);
1415 result.addAttribute(MmaSpOp::getOperandSegmentSizeAttr(),
1417 static_cast<int32_t>(frags[0].regs.size()),
1418 static_cast<int32_t>(frags[1].regs.size()),
1419 static_cast<int32_t>(frags[2].regs.size()),
1426LogicalResult MmaSpOp::verify() {
1428 auto f16Ty = Float16Type::get(context);
1429 auto i32Ty = IntegerType::get(context, 32);
1430 auto f16x2Ty = VectorType::get(2, f16Ty);
1431 auto f32Ty = Float32Type::get(context);
1432 auto f16x2x4StructTy = LLVM::LLVMStructType::getLiteral(
1433 context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
1435 auto s32x4StructTy =
1436 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty, i32Ty, i32Ty});
1437 auto f32x8StructTy =
1439 auto f16x2x2StructTy =
1440 LLVM::LLVMStructType::getLiteral(context, {f16x2Ty, f16x2Ty});
1441 auto f32x4StructTy =
1442 LLVM::LLVMStructType::getLiteral(context, {f32Ty, f32Ty, f32Ty, f32Ty});
1443 auto s32x2StructTy =
1444 LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty});
1446 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(),
1447 getShapeAttr().getK()};
1453 AllowedShapes allowedShapes;
1454 AllowedTypes expectedA;
1455 AllowedTypes expectedB;
1456 AllowedTypes expectedC;
1461 if (mmaShape[0] == 16) {
1463 Type multiplicandFragType;
1464 switch (*getMultiplicandAPtxType()) {
1465 case MMATypes::tf32:
1467 multiplicandFragType = i32Ty;
1468 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1469 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1471 allowedShapes.push_back({16, 8, 8});
1472 allowedShapes.push_back({16, 8, 16});
1474 case MMATypes::bf16:
1476 multiplicandFragType = i32Ty;
1477 expectedResult.push_back(LLVM::LLVMStructType::getLiteral(
1478 context, {f32Ty, f32Ty, f32Ty, f32Ty}));
1480 allowedShapes.push_back({16, 8, 16});
1481 allowedShapes.push_back({16, 8, 32});
1485 multiplicandFragType = f16x2Ty;
1486 expectedResult.push_back(f16x2x2StructTy);
1487 expectedResult.push_back(f32x4StructTy);
1489 allowedShapes.push_back({16, 8, 16});
1490 allowedShapes.push_back({16, 8, 32});
1496 allowedShapes.push_back({16, 8, 64});
1497 allowedShapes.push_back({16, 8, 128});
1503 allowedShapes.push_back({16, 8, 32});
1504 allowedShapes.push_back({16, 8, 64});
1506 case MMATypes::e4m3:
1507 case MMATypes::e5m2:
1508 case MMATypes::e3m2:
1509 case MMATypes::e2m3:
1510 case MMATypes::e2m1:
1512 multiplicandFragType = i32Ty;
1513 expectedResult.push_back(f16x2x2StructTy);
1514 expectedResult.push_back(f32x4StructTy);
1516 allowedShapes.push_back({16, 8, 64});
1519 return emitError(
"invalid shape or multiplicand type: ")
1520 << getMultiplicandAPtxType().value();
1524 expectedResult.push_back(s32x4StructTy);
1525 expectedC.emplace_back(4, i32Ty);
1526 multiplicandFragType = i32Ty;
1527 }
else if (*getMultiplicandAPtxType() >= MMATypes::e4m3 &&
1528 *getMultiplicandAPtxType() <= MMATypes::e2m1) {
1530 expectedC.emplace_back(2, f16x2Ty);
1531 expectedC.emplace_back(4, f32Ty);
1533 expectedC.emplace_back(2, f16x2Ty);
1534 expectedC.emplace_back(4, f32Ty);
1539 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor) / 2;
1540 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor);
1541 expectedA.emplace_back(unitA, multiplicandFragType);
1542 expectedB.emplace_back(unitB, multiplicandFragType);
1544 if (resultPtxType() != accumPtxType())
1549 if (mmaShape[0] == 8) {
1550 if (*getMultiplicandAPtxType() == MMATypes::f16) {
1551 expectedA.emplace_back(2, f16x2Ty);
1552 expectedB.emplace_back(2, f16x2Ty);
1553 expectedResult.push_back(f16x2x4StructTy);
1554 expectedResult.push_back(f32x8StructTy);
1555 expectedC.emplace_back(4, f16x2Ty);
1556 expectedC.emplace_back(8, f32Ty);
1557 allowedShapes.push_back({8, 8, 4});
1559 if (*getMultiplicandAPtxType() == MMATypes::f64) {
1560 Type f64Ty = Float64Type::get(context);
1561 expectedA.emplace_back(1, f64Ty);
1562 expectedB.emplace_back(1, f64Ty);
1563 expectedC.emplace_back(2, f64Ty);
1564 expectedResult.emplace_back(LLVM::LLVMStructType::getLiteral(
1566 allowedShapes.push_back({8, 8, 4});
1569 expectedA.push_back({i32Ty});
1570 expectedB.push_back({i32Ty});
1571 expectedC.push_back({i32Ty, i32Ty});
1572 expectedResult.push_back(s32x2StructTy);
1574 allowedShapes.push_back({8, 8, 32});
1576 allowedShapes.push_back({8, 8, 16});
1580 std::string errorMessage;
1581 llvm::raw_string_ostream errorStream(errorMessage);
1584 if (expectedA.empty() || expectedB.empty() || expectedC.empty() ||
1585 !llvm::is_contained(allowedShapes, mmaShape)) {
1586 errorStream <<
"unimplemented variant for MMA shape <";
1587 llvm::interleaveComma(mmaShape, errorStream);
1593 std::array<StringRef, 3> operandNames{
"A",
"B",
"C"};
1594 for (
const auto &iter : llvm::enumerate(
1596 auto spec = this->getODSOperandIndexAndLength(iter.index());
1598 operand_type_begin() + spec.first +
1600 bool match = llvm::is_contained(iter.value(), operandTySeg);
1603 errorStream <<
"Could not match types for the "
1604 << operandNames[iter.index()]
1605 <<
" operands; expected one of ";
1606 for (
const auto &x : iter.value()) {
1607 errorStream << x.size() <<
"x" << x[0] <<
" ";
1609 errorStream <<
"but got ";
1610 llvm::interleaveComma(operandTySeg, errorStream);
1616 if (!llvm::any_of(expectedResult, [&](
Type expectedResultType) {
1617 return expectedResultType == getResult().getType();
1620 <<
"Could not match allowed types for the result; expected one of ";
1621 llvm::interleaveComma(expectedResult, errorStream);
1622 errorStream <<
" but got " << getResult().getType();
1630 if (!getIntOverflowBehavior())
1632 getIntOverflowBehaviorAttrName().strref() +
1637 if (!getSparseMetadata().
getType().isInteger(32)) {
1638 return emitOpError() <<
"sparse metadata must be i32 type";
1642 if (!getSparsitySelector().
getType().isInteger(32)) {
1643 return emitOpError() <<
"sparsity selector must be i32 type";
1655struct MMAOperandFragment {
1656 StringRef operandName;
1657 StringRef ptxTypeAttr;
1658 SmallVector<Value, 4> regs;
1659 explicit MMAOperandFragment(StringRef name, StringRef ptxTypeName)
1660 : operandName(name), ptxTypeAttr(ptxTypeName) {}
1667 p <<
" " << name <<
"[";
1686template <
typename Op>
1691 for (
unsigned fragIdx = 0; fragIdx < frags.size(); fragIdx++) {
1692 auto &frag = frags[fragIdx];
1693 auto varOperandSpec = op.getODSOperandIndexAndLength(fragIdx);
1694 for (
auto operandIdx = varOperandSpec.first;
1695 operandIdx < varOperandSpec.first + varOperandSpec.second;
1697 frag.regs.push_back(op.getOperand(operandIdx));
1698 if (fragIdx == 0 && operandIdx == varOperandSpec.first) {
1699 regTypes.push_back(op.getOperand(operandIdx).getType());
1703 regTypes.push_back(frag.regs[0].getType());
1705 std::optional<MMATypes> inferredType =
1706 MmaOp::inferOperandMMAType(regTypes.back(),
1709 ignoreAttrNames.push_back(frag.ptxTypeAttr);
1720 auto typeParser = [&]() {
1724 operandTypes.push_back(ty);
1730 if (operandTypes.size() != 3)
1732 "expected exactly 3 types");
1741 if (!attrs.
get(
"multiplicandAPtxType")) {
1742 if (
auto inferredType =
1743 MmaOp::inferOperandMMAType(operandTypes[0],
false)) {
1744 attrs.
set(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *inferredType));
1747 if (!attrs.
get(
"multiplicandBPtxType")) {
1748 if (
auto inferredType =
1749 MmaOp::inferOperandMMAType(operandTypes[1],
false)) {
1750 attrs.
set(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *inferredType));
1756template <
typename OpType>
1759 ScaleVecSize scaleVecSize,
1760 BlockScaleFormat blockScaleFormat,
1761 MMABlockScaleKind kind) {
1763 auto &properties =
result.getOrAddProperties<
typename OpType::Properties>();
1764 properties.setShape(
1766 properties.setScaleVecSize(ScaleVecSizeAttr::get(ctx, scaleVecSize));
1767 properties.setBlockScaleFormat(
1768 BlockScaleFormatAttr::get(ctx, blockScaleFormat));
1769 properties.setKind(MMABlockScaleKindAttr::get(ctx, kind));
1776 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes) {
1777 if (multiplicandPtxTypes) {
1778 result.addAttribute(
"multiplicandAPtxType",
1779 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[0]));
1780 result.addAttribute(
"multiplicandBPtxType",
1781 MMATypesAttr::get(ctx, (*multiplicandPtxTypes)[1]));
1783 if (
auto res = MmaOp::inferOperandMMAType(operandA[0].
getType(),
false))
1784 result.addAttribute(
"multiplicandAPtxType", MMATypesAttr::get(ctx, *res));
1785 if (
auto res = MmaOp::inferOperandMMAType(operandB[0].
getType(),
false))
1786 result.addAttribute(
"multiplicandBPtxType", MMATypesAttr::get(ctx, *res));
1791template <
typename OpTy>
1793 return *MmaOp::inferOperandMMAType(
1794 cast<LLVM::LLVMStructType>(op.getRes().getType()).getBody()[0],
1804 std::array<MMAOperandFragment, 3> frags{
1805 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
1806 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
1807 MMAOperandFragment(
"C",
"")};
1809 mlir::NVVM::MmaBlockScaleOp::getOperandSegmentSizeAttr()};
1814 for (
const auto &frag : frags)
1819 {getScaleAData(), getByteIdA(), getThreadIdA()});
1821 {getScaleBData(), getByteIdB(), getThreadIdB()});
1828 frags[1].regs[0].getType(),
1829 frags[2].regs[0].getType()},
1835ParseResult MmaBlockScaleOp::parse(
OpAsmParser &parser,
1837 struct LocalOperandFragment {
1838 std::optional<MMATypes> elemtype;
1839 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
1843 std::array<LocalOperandFragment, 3> frags;
1872 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
1873 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
1876 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
1886 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
1896 result.addAttributes(namedAttributes);
1900 result.addTypes(resultTypes);
1901 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
1903 static_cast<int32_t>(frags[0].regs.size()),
1904 static_cast<int32_t>(frags[1].regs.size()),
1905 static_cast<int32_t>(frags[2].regs.size()),
1916void MmaBlockScaleOp::build(
1921 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
1922 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
1923 MMABlockScaleKind kind) {
1924 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
1927 blockScaleFormat, kind);
1929 result.addOperands(operandA);
1930 result.addOperands(operandB);
1931 result.addOperands(operandC);
1933 {scaleAData, byteIdA, threadIdA, scaleBData, byteIdB, threadIdB});
1936 multiplicandPtxTypes);
1938 result.addTypes(resultType);
1939 result.addAttribute(MmaBlockScaleOp::getOperandSegmentSizeAttr(),
1941 static_cast<int32_t>(operandA.size()),
1942 static_cast<int32_t>(operandB.size()),
1943 static_cast<int32_t>(operandC.size()),
1955 auto curOp = cast<NVVM::MmaBlockScaleOp>(op);
1959 for (
Value operand : curOp.getOperandA())
1961 for (
Value operand : curOp.getOperandB())
1963 for (
Value operand : curOp.getOperandC())
1967 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
1968 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
1969 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
1970 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
1971 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
1972 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
1974 unsigned intId = MmaBlockScaleOp::getIntrinsicID(
1975 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
1976 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
1978 curOp.getBlockScaleFormat(), curOp.getKind());
1980 return {intId, args};
1983LogicalResult MmaBlockScaleOp::verify() {
1989 if (m == 16 && n == 8 && k == 64) {
1990 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
1991 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
1993 "unsupported MMATypes attribute for mma.m16n8k64.(mxf4nvf4|mxf4)");
1994 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
1995 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
1997 "unsupported ScaleVecSize attribute for mma.m16n8k64.mxf4");
1998 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2000 "unsupported BlockScaleFormat attribute for mma.m16n8k64.mxf4");
2001 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2002 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2003 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2004 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2005 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2006 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2008 "attributes for mma.m16n8k64.mxf4nvf4");
2012 }
else if (m == 16 && n == 8 && k == 32) {
2013 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2014 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2015 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2017 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2018 "attributes for mma.m16n8k32");
2031 std::array<MMAOperandFragment, 3> frags{
2032 MMAOperandFragment(
"A", getMultiplicandAPtxTypeAttrName()),
2033 MMAOperandFragment(
"B", getMultiplicandBPtxTypeAttrName()),
2034 MMAOperandFragment(
"C",
"")};
2036 mlir::NVVM::MmaSpBlockScaleOp::getOperandSegmentSizeAttr()};
2041 for (
const auto &frag : frags)
2050 {getScaleAData(), getByteIdA(), getThreadIdA()});
2052 {getScaleBData(), getByteIdB(), getThreadIdB()});
2059 frags[1].regs[0].getType(),
2060 frags[2].regs[0].getType()},
2066ParseResult MmaSpBlockScaleOp::parse(
OpAsmParser &parser,
2068 struct LocalOperandFragment {
2069 std::optional<MMATypes> elemtype;
2070 SmallVector<OpAsmParser::UnresolvedOperand, 4> regs;
2074 std::array<LocalOperandFragment, 3> frags;
2110 for (
const auto &[idx, frag] : llvm::enumerate(frags)) {
2111 frag.elemtype = MmaOp::inferOperandMMAType(operandTypes[idx],
2114 .resolveOperands(frag.regs, operandTypes[idx], parser.
getNameLoc(),
2123 .resolveOperands(metadataOperands, i32Type, parser.
getNameLoc(),
2136 .resolveOperands(scaleAOperands, scaleTypes, parser.
getNameLoc(),
2146 result.addAttributes(namedAttributes);
2151 if (!
result.attributes.get(
"orderedMetadata"))
2154 result.addTypes(resultTypes);
2155 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2157 static_cast<int32_t>(frags[0].regs.size()),
2158 static_cast<int32_t>(frags[1].regs.size()),
2159 static_cast<int32_t>(frags[2].regs.size()),
2172void MmaSpBlockScaleOp::build(
2178 std::optional<std::array<MMATypes, 2>> multiplicandPtxTypes,
2179 ScaleVecSize scaleVecSize, BlockScaleFormat blockScaleFormat,
2180 MMABlockScaleKind kind) {
2181 assert(
shape.size() == 3 &&
"expected shape to have size 3 (m, n, k)");
2184 builder,
result,
shape, scaleVecSize, blockScaleFormat, kind);
2187 result.addOperands(operandA);
2188 result.addOperands(operandB);
2189 result.addOperands(operandC);
2190 result.addOperands({sparseMetadata, sparsitySelector, scaleAData, byteIdA,
2191 threadIdA, scaleBData, byteIdB, threadIdB});
2194 multiplicandPtxTypes);
2196 result.addTypes(resultType);
2197 result.addAttribute(MmaSpBlockScaleOp::getOperandSegmentSizeAttr(),
2199 static_cast<int32_t>(operandA.size()),
2200 static_cast<int32_t>(operandB.size()),
2201 static_cast<int32_t>(operandC.size()),
2215 auto curOp = cast<NVVM::MmaSpBlockScaleOp>(op);
2219 for (
Value operand : curOp.getOperandA())
2221 for (
Value operand : curOp.getOperandB())
2223 for (
Value operand : curOp.getOperandC())
2227 args.push_back(mt.
lookupValue(curOp.getSparseMetadata()));
2228 args.push_back(mt.
lookupValue(curOp.getSparsitySelector()));
2231 args.push_back(mt.
lookupValue(curOp.getScaleAData()));
2232 args.push_back(mt.
lookupValue(curOp.getByteIdA()));
2233 args.push_back(mt.
lookupValue(curOp.getThreadIdA()));
2234 args.push_back(mt.
lookupValue(curOp.getScaleBData()));
2235 args.push_back(mt.
lookupValue(curOp.getByteIdB()));
2236 args.push_back(mt.
lookupValue(curOp.getThreadIdB()));
2238 unsigned intId = MmaSpBlockScaleOp::getIntrinsicID(
2239 curOp.getShape().getM(), curOp.getShape().getN(), curOp.getShape().getK(),
2240 *curOp.getMultiplicandAPtxType(), *curOp.getMultiplicandBPtxType(),
2242 curOp.getBlockScaleFormat(), curOp.getKind());
2244 return {intId, args};
2247LogicalResult MmaSpBlockScaleOp::verify() {
2249 if (!getOrderedMetadata()) {
2250 return emitOpError(
"'orderedMetadata' attribute is mandatory");
2258 if (m == 16 && n == 8 && k == 128) {
2259 if (getMultiplicandAPtxType() != NVVM::MMATypes::e2m1 ||
2260 getMultiplicandBPtxType() != NVVM::MMATypes::e2m1)
2262 "unsupported MMATypes attribute for mma.m16n8k128.(mxf4nvf4|mxf4)");
2263 if (getKind() == NVVM::MMABlockScaleKind::MXF4) {
2264 if (getScaleVecSize() != NVVM::ScaleVecSize::X2)
2266 "unsupported ScaleVecSize attribute for mma.m16n8k128.mxf4");
2267 if (getBlockScaleFormat() != NVVM::BlockScaleFormat::UE8M0)
2269 "unsupported BlockScaleFormat attribute for mma.m16n8k128.mxf4");
2270 }
else if (getKind() == NVVM::MMABlockScaleKind::MXF4NVF4) {
2271 if (!((getScaleVecSize() == NVVM::ScaleVecSize::X2 &&
2272 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0) ||
2273 (getScaleVecSize() == NVVM::ScaleVecSize::X4 &&
2274 (getBlockScaleFormat() == NVVM::BlockScaleFormat::UE4M3 ||
2275 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))))
2277 "attributes for mma.m16n8k128.mxf4nvf4");
2281 }
else if (m == 16 && n == 8 && k == 64) {
2282 if (!(getKind() == NVVM::MMABlockScaleKind::MXF8F6F4 &&
2283 getScaleVecSize() == NVVM::ScaleVecSize::X1 &&
2284 getBlockScaleFormat() == NVVM::BlockScaleFormat::UE8M0))
2286 emitOpError(
"unsupported Kind, ScaleVecSize and BlockScaleFormat "
2287 "attributes for mma.m16n8k64");
2294LogicalResult ShflOp::verify() {
2295 auto returnStructType = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
2297 auto verifyTypeError = [&](Twine desc,
Type expectedType,
2298 Type actualType) -> LogicalResult {
2299 return emitOpError(
"expected " + desc +
" to be of type ")
2300 << expectedType <<
" but got " << actualType <<
" instead";
2303 if (returnStructType) {
2304 if (!getReturnValueAndIsValid())
2305 return emitOpError(
"\"return_value_and_is_valid\" attribute must be "
2306 "specified when the return type is a struct type");
2308 if (returnStructType.getBody().size() != 2)
2309 return emitOpError(
"expected return type to be a two-element struct");
2312 auto resultType = returnStruct[0];
2313 if (resultType != getVal().
getType())
2314 return verifyTypeError(
"first element in the returned struct",
2315 getVal().
getType(), resultType);
2317 auto predicateType = returnStruct[1];
2318 if (!predicateType.isInteger(1))
2319 return verifyTypeError(
"second element in the returned struct",
2323 if (getReturnValueAndIsValid())
2324 return emitOpError(
"expected return type to be a two-element struct");
2327 return verifyTypeError(
"return type", getVal().
getType(),
getType());
2333ShflOp::inferReturnTypes(
MLIRContext *context, std::optional<Location> location,
2334 ShflOp::Adaptor adaptor,
2336 Type valType = adaptor.getVal().getType();
2337 if (adaptor.getReturnValueAndIsValid())
2338 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2339 context, {valType, IntegerType::get(context, 1)}));
2341 inferredReturnTypes.push_back(valType);
2346 NVVM::MMAFrag frag,
int nRow,
2349 unsigned numberElements = 0;
2352 Type f16x2 = VectorType::get(2, builder.getF16Type());
2353 if (type == NVVM::MMATypes::f16) {
2354 elementType = f16x2;
2355 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2359 }
else if (type == NVVM::MMATypes::f32) {
2360 elementType = builder.getF32Type();
2362 }
else if (type == NVVM::MMATypes::f64) {
2363 elementType = builder.getF64Type();
2364 if (frag == NVVM::MMAFrag::a || frag == NVVM::MMAFrag::b)
2368 }
else if (type == NVVM::MMATypes::tf32) {
2369 elementType = builder.getI32Type();
2371 }
else if (type == NVVM::MMATypes::s8 || type == NVVM::MMATypes::u8) {
2372 elementType = builder.getI32Type();
2373 int parallelSize = 0;
2374 if (frag == NVVM::MMAFrag::a)
2375 parallelSize = nRow;
2376 if (frag == NVVM::MMAFrag::b)
2377 parallelSize = nCol;
2380 if (parallelSize == 16)
2383 else if (parallelSize == 8)
2385 else if (parallelSize == 32)
2387 }
else if (type == NVVM::MMATypes::s32) {
2388 elementType = builder.getI32Type();
2391 assert(numberElements != 0 && elementType !=
nullptr);
2392 return std::make_pair(elementType, numberElements);
2395static std::pair<mlir::Type, unsigned>
2399 if (frag == NVVM::MMAFrag::a) {
2402 }
else if (frag == NVVM::MMAFrag::b) {
2409 assert(nRow && nCol);
2413LogicalResult NVVM::WMMALoadOp::verify() {
2414 unsigned addressSpace =
2415 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2416 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2417 addressSpace != NVVMMemorySpace::Shared)
2418 return emitOpError(
"expected source pointer in memory "
2421 if (NVVM::WMMALoadOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2422 getEltype(), getFrag()) == 0)
2423 return emitOpError() <<
"invalid attribute combination";
2428 if (typeInfo.first == f64Ty && typeInfo.second == 1) {
2430 return emitOpError(
"expected destination type to be f64");
2434 Type dstType = LLVM::LLVMStructType::getLiteral(
2437 return emitOpError(
"expected destination type is a structure of ")
2438 << typeInfo.second <<
" elements of type " << typeInfo.first;
2442LogicalResult NVVM::WMMAStoreOp::verify() {
2443 unsigned addressSpace =
2444 llvm::cast<LLVM::LLVMPointerType>(getPtr().
getType()).getAddressSpace();
2445 if (addressSpace != 0 && addressSpace != NVVMMemorySpace::Global &&
2446 addressSpace != NVVMMemorySpace::Shared)
2447 return emitOpError(
"expected operands to be a source pointer in memory "
2450 if (NVVM::WMMAStoreOp::getIntrinsicID(
getM(),
getN(), getK(), getLayout(),
2452 return emitOpError() <<
"invalid attribute combination";
2455 if (getArgs().size() != typeInfo.second)
2456 return emitOpError() <<
"expected " << typeInfo.second <<
" data operands";
2457 if (llvm::any_of(getArgs(), [&typeInfo](
Value operands) {
2458 return operands.
getType() != typeInfo.first;
2460 return emitOpError() <<
"expected data operands of type " << typeInfo.first;
2464LogicalResult NVVM::WMMAMmaOp::verify() {
2465 if (NVVM::WMMAMmaOp::getIntrinsicID(
getM(),
getN(), getK(), getLayoutA(),
2466 getLayoutB(), getEltypeA(),
2468 return emitOpError() <<
"invalid attribute combination";
2476 arguments.append(typeInfoA.second, typeInfoA.first);
2477 arguments.append(typeInfoB.second, typeInfoB.first);
2478 arguments.append(typeInfoC.second, typeInfoC.first);
2479 unsigned numArgs = arguments.size();
2480 if (getArgs().size() != numArgs)
2481 return emitOpError() <<
"expected " << numArgs <<
" arguments";
2482 for (
unsigned i = 0; i < numArgs; i++) {
2483 if (getArgs()[i].
getType() != arguments[i])
2484 return emitOpError() <<
"expected argument " << i <<
" to be of type "
2487 Type dstType = LLVM::LLVMStructType::getLiteral(
2490 return emitOpError(
"expected destination type is a structure of ")
2491 << typeInfoC.second <<
" elements of type " << typeInfoC.first;
2495LogicalResult NVVM::LdMatrixOp::verify() {
2497 if (m == 8 && n == 8) {
2498 if (num != 1 && num != 2 && num != 4) {
2499 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x8 "
2502 if (getEltType() != LdStMatrixEltType::B16) {
2503 return emitOpError(
"expected element type to be b16 for 8x8 matrix");
2505 }
else if (m == 8 && n == 16) {
2506 if (num != 1 && num != 2 && num != 4) {
2507 return emitOpError(
"expected num attribute to be 1, 2 or 4 for 8x16 "
2510 if (getLayout() != MMALayout::row) {
2511 return emitOpError(
"expected layout to be row for 8x16 matrix");
2513 if (getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2514 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2515 return emitOpError(
"expected element type to be b8x16.b4x16_p64 or "
2516 "b8x16.b6x16_p32 for 8x16 matrix");
2518 }
else if (m == 16 && n == 16) {
2519 if (num != 1 && num != 2) {
2520 return emitOpError(
"expected num attribute to be 1 or 2 for 16x16 "
2523 if (getLayout() != MMALayout::col) {
2524 return emitOpError(
"expected layout to be col for 16x16 matrix");
2526 if (getEltType() != LdStMatrixEltType::B8 &&
2527 getEltType() != LdStMatrixEltType::B8X16_B4X16_P64 &&
2528 getEltType() != LdStMatrixEltType::B8X16_B6X16_P32) {
2529 return emitOpError(
"expected element type to be b8, b8x16.b4x16_p64 or "
2530 "b8x16.b6x16_p32 for 16x16 matrix");
2533 return emitOpError(
"expected shape to be 8x8, 8x16 or 16x16");
2537 uint32_t numElements = (m == 16 && n == 16 ? num * 2 : num);
2538 if (numElements == 1 &&
getType() != i32)
2539 return emitOpError(
"expected destination type is i32");
2540 if (numElements == 2 || numElements == 4) {
2541 Type dstType = LLVM::LLVMStructType::getLiteral(
2544 return emitOpError(
"expected destination type is a structure of ")
2545 << numElements <<
" elements of type i32";
2551LogicalResult LdMatrixOp::inferReturnTypes(
2552 MLIRContext *context, std::optional<Location> location,
2554 uint32_t num = adaptor.getNum();
2555 uint32_t m = adaptor.getShape().getM();
2556 uint32_t n = adaptor.getShape().getN();
2557 uint32_t numElements = (m == 16 && n == 16) ? num * 2 : num;
2559 Type i32 = IntegerType::get(context, 32);
2560 if (numElements == 1)
2561 inferredReturnTypes.push_back(i32);
2563 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2568LogicalResult NVVM::StMatrixOp::verify() {
2569 int numMatrix = getSources().size();
2570 if (numMatrix != 1 && numMatrix != 2 && numMatrix != 4)
2571 return emitOpError(
"expected num attribute to be 1, 2 or 4");
2574 if (m == 8 && n == 8) {
2575 if (getEltType() != NVVM::LdStMatrixEltType::B16) {
2576 return emitOpError(
"expected element type to be B16 for 8x8 matrix");
2578 }
else if (m == 16 && n == 8) {
2579 if (getEltType() != NVVM::LdStMatrixEltType::B8) {
2580 return emitOpError(
"expected element type to be B8 for 16x8 matrix");
2582 if (getLayout() != NVVM::MMALayout::col) {
2583 return emitOpError(
"expected layout to be col for 16x8 matrix");
2586 return emitOpError(
"expected shape to be 8x8 or 16x8");
2592LogicalResult NVVM::MovMatrixOp::verify() {
2594 if (m != 8 || n != 8)
2596 if (getLayout() != NVVM::MMALayout::col)
2598 if (getEltType() != NVVM::LdStMatrixEltType::B16)
2599 return emitOpError(
"expected element type to be b16");
2604 if (typeA == NVVM::WGMMATypes::tf32)
2606 if (typeA == NVVM::WGMMATypes::f16 || typeA == NVVM::WGMMATypes::bf16)
2608 if (typeA == NVVM::WGMMATypes::s8 || typeA == NVVM::WGMMATypes::u8)
2610 if (typeA == NVVM::WGMMATypes::e4m3 || typeA == NVVM::WGMMATypes::e5m2)
2612 if (typeA == NVVM::WGMMATypes::b1)
2618 NVVM::WGMMATypes typeA,
2619 NVVM::WGMMATypes typeB) {
2621 case NVVM::WGMMATypes::f16:
2622 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
2623 typeB == NVVM::WGMMATypes::f16)
2626 case NVVM::WGMMATypes::tf32:
2627 if (typeD == NVVM::WGMMATypes::f32 && typeB == NVVM::WGMMATypes::tf32)
2630 case NVVM::WGMMATypes::u8:
2631 case NVVM::WGMMATypes::s8:
2632 if (typeD == NVVM::WGMMATypes::s32 &&
2633 (typeB == NVVM::WGMMATypes::u8 || typeB == NVVM::WGMMATypes::s8))
2636 case NVVM::WGMMATypes::b1:
2637 if (typeD == NVVM::WGMMATypes::s32 && typeB == NVVM::WGMMATypes::b1)
2640 case NVVM::WGMMATypes::bf16:
2641 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
2642 typeB == NVVM::WGMMATypes::bf16)
2645 case NVVM::WGMMATypes::e4m3:
2646 case NVVM::WGMMATypes::e5m2:
2647 if ((typeD == NVVM::WGMMATypes::f32 || typeD == NVVM::WGMMATypes::f16) &&
2648 (typeB == NVVM::WGMMATypes::e5m2 || typeB == NVVM::WGMMATypes::e4m3))
2651 case WGMMATypes::f32:
2652 case WGMMATypes::s32:
2653 llvm_unreachable(
"unsupported input types");
2661 72, 80, 88, 96, 104, 112, 120, 128,
2662 136, 144, 152, 160, 168, 176, 184, 192,
2663 200, 208, 216, 224, 232, 240, 248, 256};
2665 80, 96, 112, 128, 144, 160,
2666 176, 192, 208, 224, 240, 256};
2668 case WGMMATypes::f16:
2669 case WGMMATypes::tf32:
2670 case WGMMATypes::bf16:
2671 case WGMMATypes::e4m3:
2672 case WGMMATypes::e5m2:
2673 if (llvm::is_contained(allowedN, sizeN))
2676 case WGMMATypes::u8:
2677 case WGMMATypes::s8:
2678 case WGMMATypes::b1:
2679 if (llvm::is_contained(allowedNshort, sizeN))
2682 case WGMMATypes::f32:
2683 case WGMMATypes::s32:
2684 llvm_unreachable(
"unsupported input types");
2690LogicalResult NVVM::WgmmaMmaAsyncOp::verify() {
2691 Value outValue = getResults();
2692 auto stype = dyn_cast<LLVM::LLVMStructType>(outValue.
getType());
2694 return emitOpError() <<
"expected results to be struct";
2695 int outputSize = stype.getBody().size();
2696 WGMMATypes typeD = getTypeD();
2697 WGMMATypes typeA = getTypeA();
2698 WGMMATypes typeB = getTypeB();
2700 for (
Type t : stype.getBody()) {
2701 if (t != stype.getBody().front())
2703 <<
"all elements in struct must be same type but there is " << t;
2706 if (typeD != WGMMATypes::f32 && typeD != WGMMATypes::f16 &&
2707 typeD != WGMMATypes::s32) {
2708 return emitOpError() <<
"does not support the given output type " << typeD;
2710 if (typeD == WGMMATypes::s32 &&
2711 (getScaleA() == WGMMAScaleIn::neg || getScaleB() == WGMMAScaleIn::neg)) {
2712 return emitOpError() <<
"has s32 output, scaleA and scaleB cannot be neg";
2716 return emitOpError() << typeD <<
" += " << typeA <<
" * " << typeB
2717 <<
", it is not supported.";
2727 return emitOpError() <<
"shape 'k' must be " << allowedK.value()
2728 <<
" for input type " << typeA;
2732 return emitOpError() <<
"has input type " << typeA <<
" n is set to "
2733 <<
getShape().getN() <<
", it is not supported.";
2740 if ((typeA != WGMMATypes::f16 && typeA != WGMMATypes::bf16) &&
2741 (getLayoutA() == mlir::NVVM::MMALayout::col ||
2742 getLayoutB() == mlir::NVVM::MMALayout::row)) {
2744 <<
"given layouts layout_a = " << getLayoutA()
2745 <<
" and layout_b = " << getLayoutB() <<
" for input types " << typeA
2747 <<
" requires transpose. However, this is only supported for: "
2748 << MMATypes::f16 <<
" and " << MMATypes::bf16;
2752 int expectedOutput = 0;
2753 if (typeD == WGMMATypes::f32 || typeD == WGMMATypes::s32)
2754 expectedOutput =
getShape().getN() / 2;
2755 if (typeD == WGMMATypes::f16)
2756 expectedOutput =
getShape().getN() / 4;
2757 if (outputSize != expectedOutput) {
2758 return emitOpError() <<
"results " << expectedOutput
2759 <<
", however output struct has " << outputSize
2763 if (typeD != WGMMATypes::s32 &&
2764 getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
2765 NVVM::MMAIntOverflow::satfinite) {
2767 <<
" `satfinite` can be only used with s32 accumulator, however "
2768 "the current accumulator is "
2775std::string NVVM::WgmmaMmaAsyncOp::getPtx() {
2778 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
2780 StringRef outputTypeName = stringifyWGMMATypes(getTypeD());
2782 int expectedOutputRegisters = 0;
2783 if (getTypeD() == WGMMATypes::f16)
2784 expectedOutputRegisters =
getShape().getN() / 4;
2786 expectedOutputRegisters =
getShape().getN() / 2;
2789 llvm::raw_string_ostream ss(ptx);
2794 << ((expectedOutputRegisters * 2) + 2)
2796 "wgmma.mma_async.sync.aligned.m"
2797 << m <<
"n" << n <<
"k" << k <<
"." << outputTypeName <<
"." << getTypeA()
2798 <<
"." << getTypeB();
2799 if (getSatfinite().value_or(NVVM::MMAIntOverflow::wrapped) ==
2800 NVVM::MMAIntOverflow::satfinite)
2804 for (; regCnt < expectedOutputRegisters; ++regCnt) {
2805 ss <<
"$" << regCnt;
2806 if (regCnt != expectedOutputRegisters - 1)
2812 regCnt = (regCnt * 2);
2813 ss <<
" $" << (regCnt) <<
"," <<
" $" << (regCnt + 1) <<
"," <<
" p";
2814 if (getTypeD() != WGMMATypes::s32) {
2815 ss <<
", $" << (regCnt + 3) <<
", $" << (regCnt + 4);
2819 ss <<
", $" << (regCnt + 5) <<
", $" << (regCnt + 6);
2826bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
2830 bool isF16 = getTypeA() == WGMMATypes::f16 || getTypeA() == WGMMATypes::bf16;
2837 asmValues.push_back({makeConstantI32(rewriter,
static_cast<int>(getScaleD())),
2839 if (getTypeD() != WGMMATypes::s32) {
2840 asmValues.push_back(
2841 {makeConstantI32(rewriter,
2842 getScaleA() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
2844 asmValues.push_back(
2845 {makeConstantI32(rewriter,
2846 getScaleB() == NVVM::WGMMAScaleIn::neg ? -1 : 1),
2850 asmValues.push_back(
2851 {makeConstantI32(rewriter,
static_cast<int>(getLayoutA())),
2853 asmValues.push_back(
2854 {makeConstantI32(rewriter, 1 -
static_cast<int>(getLayoutB())),
2860LogicalResult NVVM::FenceProxyOp::verify() {
2861 if (getKind() == NVVM::ProxyKind::async_shared && !getSpace().has_value()) {
2862 return emitOpError() <<
"async_shared fence requires space attribute";
2864 if (getKind() != NVVM::ProxyKind::async_shared && getSpace().has_value()) {
2865 return emitOpError() <<
"only async_shared fence can have space attribute";
2870LogicalResult NVVM::FenceProxyAcquireOp::verify() {
2871 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
2872 return emitOpError(
"uni-directional proxies only support generic for "
2873 "from_proxy attribute");
2875 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
2876 return emitOpError(
"uni-directional proxies only support tensormap "
2877 "for to_proxy attribute");
2881LogicalResult NVVM::FenceProxyReleaseOp::verify() {
2882 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
2883 return emitOpError(
"uni-directional proxies only support generic for "
2884 "from_proxy attribute");
2886 if (getToProxy() != NVVM::ProxyKind::TENSORMAP)
2887 return emitOpError(
"uni-directional proxies only support tensormap "
2888 "for to_proxy attribute");
2892LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
2893 if (getFromProxy() != NVVM::ProxyKind::GENERIC)
2894 return emitOpError(
"only generic is support for from_proxy attribute");
2896 if (getToProxy() != NVVM::ProxyKind::async)
2897 return emitOpError(
"only async is supported for to_proxy attribute");
2901LogicalResult NVVM::SetMaxRegisterOp::verify() {
2902 if (getRegCount() % 8)
2903 return emitOpError(
"new register size must be multiple of 8");
2904 if (getRegCount() < 24 || getRegCount() > 256)
2905 return emitOpError(
"new register size must be in between 24 to 256");
2909LogicalResult NVVM::Tcgen05CpOp::verify() {
2910 auto mc = getMulticast();
2912 using SH = Tcgen05CpShape;
2913 using MC = Tcgen05CpMulticast;
2915 case SH::SHAPE_128x256b:
2916 case SH::SHAPE_128x128b:
2917 case SH::SHAPE_4x256b:
2919 return emitError(
"Invalid multicast type for tcgen05.cp Op");
2921 case SH::SHAPE_64x128b:
2922 if (mc != MC::WARPX2_01_23 && mc != MC::WARPX2_02_13)
2923 return emitError(
"Shape 64x128b requires multicast warpx2_01_23 or "
2924 "warpx2_02_13 for tcgen05.cp Op");
2926 case SH::SHAPE_32x128b:
2927 if (mc != MC::WARPX4)
2929 "Shape 32x128b requires multicast warpx4 for tcgen05.cp Op");
2935LogicalResult NVVM::MatchSyncOp::verify() {
2936 if (getKind() == NVVM::MatchSyncKind::all) {
2937 auto type = llvm::dyn_cast<LLVM::LLVMStructType>(
getType());
2938 if (!type || type.getBody().size() != 2 ||
2939 !type.getBody()[0].isInteger(32) || !type.getBody()[1].isInteger(1)) {
2940 return emitOpError(
"match.sync 'all' returns a two element struct with "
2941 "first element as i32 and second element as i1");
2944 if (!
getType().isInteger(32)) {
2945 return emitOpError(
"match.sync 'any' returns an i32");
2951LogicalResult MatchSyncOp::inferReturnTypes(
2952 MLIRContext *context, std::optional<Location> location,
2954 if (adaptor.getKind() == NVVM::MatchSyncKind::all)
2955 inferredReturnTypes.push_back(LLVM::LLVMStructType::getLiteral(
2957 {IntegerType::get(context, 32), IntegerType::get(context, 1)}));
2959 inferredReturnTypes.push_back(IntegerType::get(context, 32));
2963LogicalResult NVVM::VoteSyncOp::verify() {
2964 if (getKind() == NVVM::VoteSyncKind::ballot) {
2965 if (!
getType().isInteger(32)) {
2966 return emitOpError(
"vote.sync 'ballot' returns an i32");
2969 if (!
getType().isInteger(1)) {
2970 return emitOpError(
"vote.sync 'any', 'all' and 'uni' returns an i1");
2976LogicalResult VoteSyncOp::inferReturnTypes(
2977 MLIRContext *context, std::optional<Location> location,
2979 unsigned width = adaptor.getKind() == NVVM::VoteSyncKind::ballot ? 32 : 1;
2980 inferredReturnTypes.push_back(IntegerType::get(context, width));
2984LogicalResult NVVM::PrefetchOp::verify() {
2985 using MemSpace = NVVM::NVVMMemorySpace;
2986 using CacheLevel = NVVM::PrefetchCacheLevel;
2988 unsigned addressSpace =
2989 llvm::cast<LLVM::LLVMPointerType>(getAddr().
getType()).getAddressSpace();
2990 std::optional<NVVM::CacheEvictionPriority> evictPriority = getEvictPriority();
2991 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = getCacheLevel();
2993 if (getTensormap() && cacheLevel)
2994 return emitOpError(
"cannot specify both tensormap and cache level");
2996 if (getTensormap()) {
2997 if (addressSpace != MemSpace::Generic &&
2998 addressSpace != MemSpace::Constant) {
3000 "prefetch tensormap requires a generic or constant pointer");
3003 if (evictPriority) {
3005 "prefetch tensormap does not support eviction priority");
3008 if (getInParamSpace() && addressSpace != MemSpace::Generic) {
3010 "in_param_space can only be specified for a generic pointer");
3013 }
else if (cacheLevel) {
3014 if (addressSpace != MemSpace::Generic && addressSpace != MemSpace::Global &&
3015 addressSpace != MemSpace::Local) {
3016 return emitOpError(
"prefetch to cache level requires a generic, global, "
3017 "or local pointer");
3021 if (*cacheLevel != CacheLevel::L1) {
3023 "unsupported cache level, the only supported uniform "
3024 "cache level is L1");
3027 if (addressSpace != MemSpace::Generic) {
3029 "prefetch to uniform cache requires a generic pointer");
3033 if (evictPriority) {
3034 if (*cacheLevel != CacheLevel::L2)
3036 "cache eviction priority supported only for cache level L2");
3038 if (addressSpace != MemSpace::Global)
3039 return emitOpError(
"cache eviction priority requires a global pointer");
3041 if (*evictPriority != NVVM::CacheEvictionPriority::EvictNormal &&
3042 *evictPriority != NVVM::CacheEvictionPriority::EvictLast)
3044 "unsupported cache eviction priority, only evict_last and "
3045 "evict_normal are supported");
3049 return emitOpError(
"predicate supported only on prefetch tensormap");
3053 "requires specification of either cache level or tensormap");
3059LogicalResult NVVM::ClusterLaunchControlQueryCancelOp::verify() {
3060 switch (getQueryType()) {
3061 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
3063 return emitOpError(
"is_canceled query type returns an i1");
3065 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
3066 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
3067 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
3068 if (!
getType().isInteger(32)) {
3069 return emitOpError(
"get_first_cta_id_x, get_first_cta_id_y, "
3070 "get_first_cta_id_z query types return an i32");
3077LogicalResult ClusterLaunchControlQueryCancelOp::inferReturnTypes(
3078 MLIRContext *context, std::optional<Location> location,
3079 ClusterLaunchControlQueryCancelOp::Adaptor adaptor,
3082 adaptor.getQueryType() == NVVM::ClusterLaunchControlQueryType::IS_CANCELED
3085 inferredReturnTypes.push_back(IntegerType::get(context, width));
3089LogicalResult NVVM::ReduxOp::verify() {
3092 if (!reduxType.
isF32()) {
3094 return emitOpError(
"abs attribute is supported only for f32 type");
3096 return emitOpError(
"nan attribute is supported only for f32 type");
3099 NVVM::ReductionKind kind = getKind();
3101 case NVVM::ReductionKind::ADD:
3102 case NVVM::ReductionKind::AND:
3103 case NVVM::ReductionKind::OR:
3104 case NVVM::ReductionKind::XOR:
3105 case NVVM::ReductionKind::MAX:
3106 case NVVM::ReductionKind::MIN:
3107 case NVVM::ReductionKind::UMAX:
3108 case NVVM::ReductionKind::UMIN:
3111 << kind <<
"' reduction kind unsupported with " << reduxType
3112 <<
" type. Only supported type is 'i32'.";
3114 case NVVM::ReductionKind::FMIN:
3115 case NVVM::ReductionKind::FMAX:
3116 if (!reduxType.isF32())
3118 << kind <<
"' reduction kind unsupported with " << reduxType
3119 <<
" type. Only supported type is 'f32'.";
3126LogicalResult NVVM::TensormapReplaceOp::verify() {
3127 auto ord = getOrd();
3128 Value newVal = getNewValue();
3129 auto newValAttr = getNewValueAttr();
3130 auto fieldName = stringifyEnum(getField());
3132 if (ord && !llvm::is_contained({NVVM::TensormapField::BOX_DIM,
3133 NVVM::TensormapField::GLOBAL_DIM,
3134 NVVM::TensormapField::GLOBAL_STRIDE,
3135 NVVM::TensormapField::ELEMENT_STRIDE},
3137 return emitOpError(
"ordinal is not supported for ")
3138 << fieldName <<
" field";
3140 auto invalidNewVal = [&](llvm::Twine type) -> std::string {
3141 return llvm::Twine(
"new_value must be specified and must be an " + type +
3142 " for " + llvm::Twine(fieldName) +
" field")
3146 auto invalidNewValAttr = [&]() -> std::string {
3147 return (llvm::Twine(
3148 "new_value_attr must be specified and must be a valid ") +
3149 llvm::Twine(fieldName) +
" attribute for " + fieldName +
" field")
3153 switch (getField()) {
3154 case NVVM::TensormapField::GLOBAL_ADDRESS:
3158 case NVVM::TensormapField::RANK:
3162 case NVVM::TensormapField::GLOBAL_STRIDE:
3164 return emitOpError(
"ordinal is required for global_stride field");
3168 case NVVM::TensormapField::BOX_DIM:
3169 case NVVM::TensormapField::GLOBAL_DIM:
3170 case NVVM::TensormapField::ELEMENT_STRIDE:
3173 << stringifyEnum(getField()) <<
" field";
3177 case NVVM::TensormapField::ELEMTYPE:
3178 if (!(newValAttr && llvm::isa<TensormapElemtypeAttr>(*newValAttr)))
3181 case NVVM::TensormapField::INTERLEAVE_LAYOUT:
3182 if (!(newValAttr && llvm::isa<TensormapInterleaveLayoutAttr>(*newValAttr)))
3185 case NVVM::TensormapField::SWIZZLE_MODE:
3186 if (!(newValAttr && llvm::isa<TensormapSwizzleModeAttr>(*newValAttr)))
3189 case NVVM::TensormapField::SWIZZLE_ATOMICITY:
3190 if (!(newValAttr && llvm::isa<TensormapSwizzleAtomicityAttr>(*newValAttr)))
3193 case NVVM::TensormapField::FILL_MODE:
3194 if (!(newValAttr && llvm::isa<TensormapFillModeAttr>(*newValAttr)))
3202template <
typename OpType>
3204 mlir::NVVM::FPRoundingMode rndMode = op.getRnd();
3205 mlir::NVVM::SaturationMode satMode = op.getSat();
3206 bool isFTZ = op.getFtz();
3209 mlir::Type opBaseType = isa<VectorType>(opType)
3210 ? cast<VectorType>(opType).getElementType()
3213 if (opBaseType.
isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3214 return op.emitOpError(
"FTZ and saturation are not supported for "
3215 "additions/subtractions involving f64 type");
3217 if (opBaseType.
isF16() && !(rndMode == NVVM::FPRoundingMode::RN ||
3218 rndMode == NVVM::FPRoundingMode::NONE))
3219 return op.emitOpError(
"only RN rounding mode is supported for f16 and "
3220 "vector<2xf16> additions/subtractions");
3222 if (opBaseType.
isBF16()) {
3223 if (rndMode != NVVM::FPRoundingMode::RN &&
3224 rndMode != NVVM::FPRoundingMode::NONE)
3225 return op.emitOpError(
"only RN rounding mode is supported for bf16 and "
3226 "vector<2xbf16> additions/subtractions");
3227 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3228 return op.emitOpError(
"FTZ and saturation are not supported for bf16 and "
3229 "vector<2xbf16> additions/subtractions");
3236 if (opBaseType.
isF16() && isFTZ && satMode == NVVM::SaturationMode::NONE)
3237 return op.emitOpError(
"FTZ with no saturation is not supported for f16 and "
3238 "vector<2xf16> additions/subtractions");
3247LogicalResult NVVM::FmaOp::verify() {
3248 auto opType = getRes().getType();
3249 mlir::NVVM::FPRoundingMode rndMode = getRnd();
3250 mlir::NVVM::SaturationMode satMode = getSat();
3251 bool isFTZ = getFtz();
3252 bool isRelu = getRelu();
3253 bool hasOOB = getOob();
3255 auto getBaseFType = [](
Type type) ->
Type {
3256 if (isa<VectorType>(type))
3257 return cast<VectorType>(type).getElementType();
3261 auto opBaseType = getBaseFType(opType);
3263 if (rndMode == NVVM::FPRoundingMode::NONE)
3264 return emitOpError(
"rounding mode must be specified");
3266 if (isRelu && satMode == NVVM::SaturationMode::SAT)
3267 return emitOpError(
"relu and saturation are not supported together");
3269 if (hasOOB && (satMode == NVVM::SaturationMode::SAT || isFTZ))
3270 return emitOpError(
"oob is not supported with saturation or FTZ");
3272 if (!(opBaseType.isF16() || opBaseType.isBF16()) && (isRelu || hasOOB))
3273 return emitOpError(
"relu and oob are only supported for f16 and bf16");
3275 if (opBaseType.isF64() && (satMode != NVVM::SaturationMode::NONE || isFTZ))
3276 return emitOpError(
"FTZ and saturation are not supported for f64 type");
3278 if (opBaseType.isF16() && rndMode != NVVM::FPRoundingMode::RN)
3280 "only RN rounding mode is supported for f16 and vector<2xf16>");
3282 if (opBaseType.isBF16()) {
3283 if (rndMode != NVVM::FPRoundingMode::RN)
3285 "only RN rounding mode is supported for bf16 and vector<2xbf16>");
3286 if (satMode != NVVM::SaturationMode::NONE || isFTZ)
3288 "FTZ and saturation are not supported for bf16 and vector<2xbf16>");
3294LogicalResult NVVM::SqrtOp::verify() {
3295 if (getRnd() == NVVM::FPRoundingMode::NONE)
3296 return emitOpError(
"rounding mode cannot be None");
3298 if (getRes().
getType().isF64() && getFtz())
3299 return emitOpError(
"FTZ is not supported for f64");
3304LogicalResult NVVM::DivFOp::verify() {
3305 bool isApprox = getApprox();
3306 bool isFull = getFull();
3307 bool isF64 = getRes().getType().isF64();
3308 bool isFtz = getFtz();
3309 NVVM::FPRoundingMode rndMode = getRnd();
3311 if (isApprox && isFull)
3312 return emitOpError(
"'approx' and 'full' are mutually exclusive");
3314 if (isApprox || isFull) {
3316 return emitOpError(
"'approx' and 'full' forms are f32-only");
3317 if (rndMode != NVVM::FPRoundingMode::NONE)
3319 "'approx' and 'full' forms do not accept a rounding mode");
3324 if (rndMode == NVVM::FPRoundingMode::NONE)
3325 return emitOpError(
"rounding mode cannot be None for the rounded divide");
3327 return emitOpError(
"FTZ is not supported for f64");
3338 unsigned sizeInBits,
3340 field = builder.CreateZExtOrBitCast(field, builder.getInt32Ty());
3342 unsigned mask = (sizeInBits < 32 ? ((1u << sizeInBits) - 1) : 0xffffffffu);
3343 if (mask != 0xffffffffu)
3344 field = builder.CreateAnd(field, builder.getInt32(mask));
3346 field = builder.CreateZExtOrBitCast(field, builder.getInt64Ty());
3347 field = builder.CreateShl(field, start);
3349 return builder.CreateOr(
result, field);
3352void Tcgen05MmaSmemDescOp::createSmemDescriptor(
Operation &op,
3354 llvm::IRBuilderBase &builder) {
3355 auto thisOp = cast<NVVM::Tcgen05MmaSmemDescOp>(op);
3356 llvm::Value *smemDesc = builder.getInt64(0);
3361 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimOffset()), 14, 16);
3363 builder, smemDesc, mt.
lookupValue(thisOp.getStrideDimOffset()), 14, 32);
3369 builder, smemDesc, mt.
lookupValue(thisOp.getLeadingDimMode()), 1, 52);
3373 mt.
mapValue(thisOp.getRes()) = smemDesc;
3380std::string NVVM::MBarrierInitOp::getPtx() {
3382 return isShared ? std::string(
"mbarrier.init.shared.b64 [%0], %1;")
3383 : std::string(
"mbarrier.init.b64 [%0], %1;");
3386std::string NVVM::MBarrierArriveExpectTxOp::getPtx() {
3389 ? std::string(
"mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;")
3390 : std::string(
"mbarrier.arrive.expect_tx.b64 _, [%0], %1;");
3393std::string NVVM::MBarrierTryWaitParityOp::getPtx() {
3395 llvm::StringRef space = isShared ?
".shared" :
"";
3397 return llvm::formatv(
"{\n\t"
3398 ".reg .pred P1; \n\t"
3400 "mbarrier.try_wait.parity{0}.b64 P1, [%0], %1, %2; \n\t"
3401 "@P1 bra.uni DONE; \n\t"
3402 "bra.uni LAB_WAIT; \n\t"
3419 LLVM::FNegOp::create(rewriter, loc, op.getRhs().getType(), op.getRhs());
3422 op.getRnd(), op.getSat(), op.getFtz());
3441 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count
3442 : llvm::Intrinsic::nvvm_barrier_cta_sync_count;
3444 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all
3445 : llvm::Intrinsic::nvvm_barrier_cta_sync_all;
3450static llvm::Intrinsic::ID
3453 case NVVM::BarrierReduction::AND:
3454 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_and_aligned_all
3455 : llvm::Intrinsic::nvvm_barrier_cta_red_and_all;
3456 case NVVM::BarrierReduction::OR:
3457 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_or_aligned_all
3458 : llvm::Intrinsic::nvvm_barrier_cta_red_or_all;
3459 case NVVM::BarrierReduction::POPC:
3460 return aligned ? llvm::Intrinsic::nvvm_barrier_cta_red_popc_aligned_all
3461 : llvm::Intrinsic::nvvm_barrier_cta_red_popc_all;
3463 llvm_unreachable(
"unknown BarrierReduction kind");
3468 auto thisOp = cast<NVVM::BarrierOp>(op);
3469 llvm::Value *barrierId = thisOp.getBarrierId()
3471 : builder.getInt32(0);
3472 bool hasCount =
static_cast<bool>(thisOp.getNumberOfThreads());
3473 llvm::Intrinsic::ID
id =
3477 args.push_back(mt.
lookupValue(thisOp.getNumberOfThreads()));
3478 return {id, std::move(args)};
3483 auto thisOp = cast<NVVM::BarrierArriveOp>(op);
3484 llvm::Value *barrierId = thisOp.getBarrierId()
3486 : builder.getInt32(0);
3487 llvm::Value *numThreads = mt.
lookupValue(thisOp.getNumberOfThreads());
3488 llvm::Intrinsic::ID
id =
3490 ? llvm::Intrinsic::nvvm_barrier_cta_arrive_aligned_count
3491 : llvm::Intrinsic::nvvm_barrier_cta_arrive_count;
3492 return {id, {barrierId, numThreads}};
3497 auto thisOp = cast<NVVM::BarrierReductionOp>(op);
3499 thisOp.getAligned(), thisOp.getReductionOp());
3500 llvm::Value *barrierId = thisOp.getBarrierId()
3502 : builder.getInt32(0);
3505 builder.CreateICmpNE(mt.
lookupValue(thisOp.getReductionPredicate()),
3506 builder.getInt32(0))};
3507 return {id, std::move(args)};
3512 llvm::IRBuilderBase &builder) {
3513 auto thisOp = cast<NVVM::CosOp>(op);
3514 llvm::Intrinsic::ID
id = thisOp.getFtz()
3515 ? llvm::Intrinsic::nvvm_cos_approx_ftz_f
3516 : llvm::Intrinsic::nvvm_cos_approx_f;
3522 llvm::IRBuilderBase &builder) {
3523 auto thisOp = cast<NVVM::SinOp>(op);
3524 llvm::Intrinsic::ID
id = thisOp.getFtz()
3525 ? llvm::Intrinsic::nvvm_sin_approx_ftz_f
3526 : llvm::Intrinsic::nvvm_sin_approx_f;
3532 llvm::IRBuilderBase &builder) {
3533 auto thisOp = cast<NVVM::Log2Op>(op);
3534 llvm::Intrinsic::ID
id = thisOp.getFtz()
3535 ? llvm::Intrinsic::nvvm_lg2_approx_ftz_f
3536 : llvm::Intrinsic::nvvm_lg2_approx_f;
3542 llvm::IRBuilderBase &builder) {
3543 auto thisOp = cast<NVVM::Ex2Op>(op);
3544 llvm::Intrinsic::ID
id = thisOp.getFtz()
3545 ? llvm::Intrinsic::nvvm_ex2_approx_ftz
3546 : llvm::Intrinsic::nvvm_ex2_approx;
3552 llvm::IRBuilderBase &builder) {
3553 auto thisOp = cast<NVVM::RsqrtOp>(op);
3554 Type t = thisOp.getRes().getType();
3555 bool isFtz = thisOp.getFtz();
3557 llvm::Intrinsic::ID
id = [&] {
3559 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_f
3560 : llvm::Intrinsic::nvvm_rsqrt_approx_f;
3563 return isFtz ? llvm::Intrinsic::nvvm_rsqrt_approx_ftz_d
3564 : llvm::Intrinsic::nvvm_rsqrt_approx_d;
3572 llvm::IRBuilderBase &builder) {
3573 auto thisOp = cast<NVVM::SqrtOp>(op);
3574 Type t = thisOp.getRes().getType();
3575 NVVM::FPRoundingMode rndMode = thisOp.getRnd();
3576 bool isFtz = thisOp.getFtz();
3580 unsigned rndIndex =
static_cast<unsigned>(rndMode) - 1;
3582 static constexpr llvm::Intrinsic::ID f32IDs[] = {
3583 llvm::Intrinsic::nvvm_sqrt_rn_f,
3584 llvm::Intrinsic::nvvm_sqrt_rm_f,
3585 llvm::Intrinsic::nvvm_sqrt_rp_f,
3586 llvm::Intrinsic::nvvm_sqrt_rz_f,
3588 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
3589 llvm::Intrinsic::nvvm_sqrt_rn_ftz_f,
3590 llvm::Intrinsic::nvvm_sqrt_rm_ftz_f,
3591 llvm::Intrinsic::nvvm_sqrt_rp_ftz_f,
3592 llvm::Intrinsic::nvvm_sqrt_rz_ftz_f,
3594 static constexpr llvm::Intrinsic::ID f64IDs[] = {
3595 llvm::Intrinsic::nvvm_sqrt_rn_d,
3596 llvm::Intrinsic::nvvm_sqrt_rm_d,
3597 llvm::Intrinsic::nvvm_sqrt_rp_d,
3598 llvm::Intrinsic::nvvm_sqrt_rz_d,
3601 llvm::Intrinsic::ID
id =
3602 t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
3610 llvm::IRBuilderBase &builder) {
3611 auto thisOp = cast<NVVM::SqrtApproxOp>(op);
3612 llvm::Intrinsic::ID
id = thisOp.getFtz()
3613 ? llvm::Intrinsic::nvvm_sqrt_approx_ftz_f
3614 : llvm::Intrinsic::nvvm_sqrt_approx_f;
3620 llvm::IRBuilderBase &builder) {
3621 auto thisOp = cast<NVVM::DivFOp>(op);
3622 bool isFtz = thisOp.getFtz();
3624 llvm::Intrinsic::ID id;
3626 if (thisOp.getApprox()) {
3627 id = isFtz ? llvm::Intrinsic::nvvm_div_approx_ftz_f
3628 : llvm::Intrinsic::nvvm_div_approx_f;
3629 }
else if (thisOp.getFull()) {
3632 id = isFtz ? llvm::Intrinsic::nvvm_div_full_ftz
3633 : llvm::Intrinsic::nvvm_div_full;
3636 unsigned rndIndex =
static_cast<unsigned>(thisOp.getRnd()) - 1;
3638 static constexpr llvm::Intrinsic::ID f32IDs[] = {
3639 llvm::Intrinsic::nvvm_div_rn_f,
3640 llvm::Intrinsic::nvvm_div_rm_f,
3641 llvm::Intrinsic::nvvm_div_rp_f,
3642 llvm::Intrinsic::nvvm_div_rz_f,
3644 static constexpr llvm::Intrinsic::ID f32FTZIDs[] = {
3645 llvm::Intrinsic::nvvm_div_rn_ftz_f,
3646 llvm::Intrinsic::nvvm_div_rm_ftz_f,
3647 llvm::Intrinsic::nvvm_div_rp_ftz_f,
3648 llvm::Intrinsic::nvvm_div_rz_ftz_f,
3650 static constexpr llvm::Intrinsic::ID f64IDs[] = {
3651 llvm::Intrinsic::nvvm_div_rn_d,
3652 llvm::Intrinsic::nvvm_div_rm_d,
3653 llvm::Intrinsic::nvvm_div_rp_d,
3654 llvm::Intrinsic::nvvm_div_rz_d,
3656 Type t = thisOp.getRes().getType();
3657 id = t.
isF32() ? (isFtz ? f32FTZIDs[rndIndex] : f32IDs[rndIndex])
3667 llvm::IRBuilderBase &builder) {
3668 auto thisOp = cast<NVVM::PMEventOp>(op);
3672 llvm::Value *maskVal;
3673 if (
auto eventAttr = thisOp.getEventIdAttr()) {
3674 uint16_t mask =
static_cast<uint16_t
>(1u << eventAttr.getInt());
3675 maskVal = llvm::ConstantInt::get(i16Ty, mask);
3678 llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
3681 return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
3686 auto thisOp = cast<NVVM::MBarrierInitOp>(op);
3688 llvm::Intrinsic::ID
id = isShared ? llvm::Intrinsic::nvvm_mbarrier_init_shared
3689 : llvm::Intrinsic::nvvm_mbarrier_init;
3694 args.push_back(mt.
lookupValue(thisOp.getCount()));
3696 return {id, std::move(args)};
3701 auto thisOp = cast<NVVM::MBarrierInvalOp>(op);
3703 llvm::Intrinsic::ID
id = isShared
3704 ? llvm::Intrinsic::nvvm_mbarrier_inval_shared
3705 : llvm::Intrinsic::nvvm_mbarrier_inval;
3712 auto thisOp = cast<NVVM::MBarrierExpectTxOp>(op);
3715 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3718 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3720 static constexpr llvm::Intrinsic::ID IDs[] = {
3721 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cta,
3722 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cta_space_cluster,
3723 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cta,
3724 llvm::Intrinsic::nvvm_mbarrier_expect_tx_scope_cluster_space_cluster};
3729 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
3731 return {IDs[
index], std::move(args)};
3736 auto thisOp = cast<NVVM::MBarrierCompleteTxOp>(op);
3739 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3742 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3744 static constexpr llvm::Intrinsic::ID IDs[] = {
3745 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cta,
3746 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cta_space_cluster,
3747 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cta,
3748 llvm::Intrinsic::nvvm_mbarrier_complete_tx_scope_cluster_space_cluster};
3753 args.push_back(mt.
lookupValue(thisOp.getTxcount()));
3755 return {IDs[
index], std::move(args)};
3760 auto thisOp = cast<NVVM::MBarrierArriveOp>(op);
3763 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3766 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3768 static constexpr llvm::Intrinsic::ID IDs[] = {
3769 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta,
3770 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cluster,
3771 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cta,
3772 llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cluster_space_cluster};
3773 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
3774 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cta,
3775 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cta_space_cluster,
3776 llvm::Intrinsic::nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cta,
3778 nvvm_mbarrier_arrive_relaxed_scope_cluster_space_cluster};
3779 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
3783 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
3790 bool hasCount =
static_cast<bool>(thisOp.getCount());
3792 (
id == llvm::Intrinsic::nvvm_mbarrier_arrive_scope_cta_space_cta))
3793 return {llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {mbar}};
3797 llvm::Value *count =
3799 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
3800 return {id, {mbar, count}};
3805 auto thisOp = cast<NVVM::MBarrierArriveDropOp>(op);
3808 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3811 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3813 static constexpr llvm::Intrinsic::ID IDs[] = {
3814 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cta,
3815 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cta_space_cluster,
3816 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cta,
3817 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_scope_cluster_space_cluster};
3818 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
3819 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cta,
3821 nvvm_mbarrier_arrive_drop_relaxed_scope_cta_space_cluster,
3823 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cta,
3825 nvvm_mbarrier_arrive_drop_relaxed_scope_cluster_space_cluster};
3826 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
3830 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
3836 bool hasCount =
static_cast<bool>(thisOp.getCount());
3837 llvm::Value *count =
3839 : llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), 1);
3841 return {id, {mbar, count}};
3844bool MBarrierArriveExpectTxOp::getAsmValues(
3851 for (
auto val : getOperands())
3859 auto thisOp = cast<NVVM::MBarrierArriveExpectTxOp>(op);
3862 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3865 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3868 static constexpr llvm::Intrinsic::ID IDs[] = {
3869 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cta,
3870 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cta_space_cluster,
3871 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cta,
3872 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_scope_cluster_space_cluster};
3873 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
3874 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cta,
3875 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cta_space_cluster,
3876 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cta,
3877 llvm::Intrinsic::nvvm_mbarrier_arrive_expect_tx_relaxed_scope_cluster_space_cluster};
3879 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
3882 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
3883 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
3888 return {id, {mbar, txcount}};
3893 auto thisOp = cast<NVVM::MBarrierArriveDropExpectTxOp>(op);
3896 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3899 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isClusterSpace ? 1 : 0);
3902 static constexpr llvm::Intrinsic::ID IDs[] = {
3903 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cta,
3904 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cta_space_cluster,
3905 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cta,
3906 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_scope_cluster_space_cluster};
3907 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
3908 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cta,
3909 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cta_space_cluster,
3910 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cta,
3911 llvm::Intrinsic::nvvm_mbarrier_arrive_drop_expect_tx_relaxed_scope_cluster_space_cluster};
3913 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
3916 llvm::Value *txcount = mt.
lookupValue(thisOp.getTxcount());
3917 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
3922 return {id, {mbar, txcount}};
3927 auto thisOp = cast<NVVM::MBarrierArriveNocompleteOp>(op);
3929 llvm::Intrinsic::ID
id =
3930 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared
3931 : llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete;
3935 args.push_back(mt.
lookupValue(thisOp.getCount()));
3937 return {id, std::move(args)};
3942 auto thisOp = cast<NVVM::MBarrierArriveDropNocompleteOp>(op);
3944 llvm::Intrinsic::ID
id =
3945 isShared ? llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete_shared
3946 : llvm::Intrinsic::nvvm_mbarrier_arrive_drop_noComplete;
3950 args.push_back(mt.
lookupValue(thisOp.getCount()));
3952 return {id, std::move(args)};
3957 auto thisOp = cast<NVVM::MBarrierTestWaitOp>(op);
3958 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
3959 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3962 size_t index = ((isClusterScope ? 1 : 0) << 1) | (isPhaseParity ? 1 : 0);
3965 static constexpr llvm::Intrinsic::ID IDs[] = {
3966 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cta_space_cta,
3967 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cta_space_cta,
3968 llvm::Intrinsic::nvvm_mbarrier_test_wait_scope_cluster_space_cta,
3969 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_scope_cluster_space_cta};
3970 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
3971 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cta_space_cta,
3972 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cta_space_cta,
3973 llvm::Intrinsic::nvvm_mbarrier_test_wait_relaxed_scope_cluster_space_cta,
3974 llvm::Intrinsic::nvvm_mbarrier_test_wait_parity_relaxed_scope_cluster_space_cta};
3976 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
3979 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
3980 llvm::Value *input = mt.
lookupValue(thisOp.getStateOrPhase());
3985 return {id, {mbar, input}};
3990 auto thisOp = cast<NVVM::MBarrierTryWaitOp>(op);
3991 bool isPhaseParity = thisOp.getStateOrPhase().getType().isInteger(32);
3992 bool isClusterScope = thisOp.getScope() == NVVM::MemScopeKind::CLUSTER;
3993 bool hasTicks =
static_cast<bool>(thisOp.getTicks());
3997 size_t index = ((hasTicks ? 1 : 0) << 2) | ((isClusterScope ? 1 : 0) << 1) |
3998 (isPhaseParity ? 1 : 0);
4001 static constexpr llvm::Intrinsic::ID IDs[] = {
4002 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cta_space_cta,
4003 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cta_space_cta,
4004 llvm::Intrinsic::nvvm_mbarrier_try_wait_scope_cluster_space_cta,
4005 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_scope_cluster_space_cta,
4006 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cta_space_cta,
4007 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cta_space_cta,
4008 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_scope_cluster_space_cta,
4009 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_scope_cluster_space_cta};
4010 static constexpr llvm::Intrinsic::ID relaxedIDs[] = {
4011 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cta_space_cta,
4012 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cta_space_cta,
4013 llvm::Intrinsic::nvvm_mbarrier_try_wait_relaxed_scope_cluster_space_cta,
4014 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_relaxed_scope_cluster_space_cta,
4015 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cta_space_cta,
4016 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cta_space_cta,
4017 llvm::Intrinsic::nvvm_mbarrier_try_wait_tl_relaxed_scope_cluster_space_cta,
4018 llvm::Intrinsic::nvvm_mbarrier_try_wait_parity_tl_relaxed_scope_cluster_space_cta};
4020 auto id = thisOp.getRelaxed() ? relaxedIDs[
index] : IDs[
index];
4023 llvm::Value *mbar = mt.
lookupValue(thisOp.getAddr());
4030 args.push_back(mbar);
4031 args.push_back(mt.
lookupValue(thisOp.getStateOrPhase()));
4033 args.push_back(mt.
lookupValue(thisOp.getTicks()));
4035 return {id, std::move(args)};
4040 auto thisOp = cast<NVVM::CpAsyncMBarrierArriveOp>(op);
4043 llvm::Intrinsic::ID id;
4044 if (thisOp.getNoinc()) {
4045 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc_shared
4046 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc;
4048 id = isShared ? llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive_shared
4049 : llvm::Intrinsic::nvvm_cp_async_mbarrier_arrive;
4057 llvm::IRBuilderBase &builder) {
4058 auto thisOp = cast<NVVM::MovMatrixOp>(op);
4059 return {llvm::Intrinsic::nvvm_movmatrix_sync_aligned_m8n8_trans_b16,
4063#define CP_ASYNC_ID_IMPL(mod, size, suffix) \
4064 llvm::Intrinsic::nvvm_cp_async_##mod##_shared_global_##size##suffix
4066#define GET_CP_ASYNC_ID(mod, size, has_cpsize) \
4067 has_cpsize ? CP_ASYNC_ID_IMPL(mod, size, _s) : CP_ASYNC_ID_IMPL(mod, size, )
4072 llvm::Intrinsic::ID id;
4074 auto cpAsyncOp = cast<NVVM::CpAsyncOp>(op);
4075 bool hasCpSize =
static_cast<bool>(cpAsyncOp.getCpSize());
4076 switch (cpAsyncOp.getSize()) {
4084 id = (cpAsyncOp.getModifier() == NVVM::LoadCacheModifierKind::CG)
4089 llvm_unreachable(
"Invalid copy size in CpAsyncOp.");
4093 args.push_back(mt.
lookupValue(cpAsyncOp.getDst()));
4094 args.push_back(mt.
lookupValue(cpAsyncOp.getSrc()));
4096 args.push_back(mt.
lookupValue(cpAsyncOp.getCpSize()));
4103 auto thisOp = cast<NVVM::CpAsyncBulkPrefetchOp>(op);
4105 llvm::Intrinsic::ID
id = llvm::Intrinsic::nvvm_cp_async_bulk_prefetch_L2;
4108 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4112 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4113 llvm::Value *i64Unused =
4114 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4115 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4116 args.push_back(builder.getInt1(hasCacheHint));
4118 return {id, std::move(args)};
4123 auto thisOp = cast<NVVM::CpAsyncBulkGlobalToSharedClusterOp>(op);
4127 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4129 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4133 mlir::Value multicastMask = thisOp.getMulticastMask();
4134 const bool hasMulticastMask =
static_cast<bool>(multicastMask);
4137 llvm::Value *i16Unused = llvm::ConstantInt::get(builder.getInt16Ty(), 0);
4138 args.push_back(hasMulticastMask ? mt.
lookupValue(multicastMask)
4144 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4145 llvm::Value *i64Unused = llvm::ConstantInt::get(builder.getInt64Ty(), 0);
4146 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4150 args.push_back(builder.getInt1(hasMulticastMask));
4151 args.push_back(builder.getInt1(hasCacheHint));
4153 llvm::Intrinsic::ID
id =
4155 ? llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cta
4156 : llvm::Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
4158 return {id, std::move(args)};
4163 auto thisOp = cast<NVVM::CpAsyncBulkSharedCTAToGlobalOp>(op);
4165 llvm::Intrinsic::ID
id =
4166 llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global;
4169 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4170 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4174 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4175 llvm::Value *i64Unused =
4176 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4177 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4178 args.push_back(builder.getInt1(hasCacheHint));
4181 if (
mlir::Value byteMask = thisOp.getByteMask()) {
4183 id = llvm::Intrinsic::nvvm_cp_async_bulk_shared_cta_to_global_bytemask;
4186 return {id, std::move(args)};
4189bool CpAsyncBulkTensorGlobalToSharedClusterOp::getAsmValues(
4196 for (
auto val : getOperands())
4203CpAsyncBulkTensorGlobalToSharedClusterOp::getIntrinsicIDAndArgs(
4205 auto thisOp = cast<NVVM::CpAsyncBulkTensorGlobalToSharedClusterOp>(op);
4206 const bool isCTAOnly = thisOp.getIsCTAOnly();
4210 args.push_back(mt.
lookupValue(thisOp.getDstMem()));
4212 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4222 const bool hasMC =
static_cast<bool>(mcMask);
4223 llvm::Value *i16Zero =
4224 llvm::ConstantInt::get(llvm::Type::getInt16Ty(mt.
getLLVMContext()), 0);
4228 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4229 llvm::Value *i64Zero =
4230 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4236 thisOp.getGroup() ? (
static_cast<int32_t
>(*thisOp.getGroup()) + 1) : 0;
4238 llvm::ConstantInt::get(llvm::Type::getInt32Ty(mt.
getLLVMContext()), val);
4242 args.push_back(hasMC ? mt.
lookupValue(mcMask) : i16Zero);
4243 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4244 args.push_back(builder.getInt1(hasMC));
4245 args.push_back(builder.getInt1(hasCacheHint));
4249 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Zero);
4250 args.push_back(builder.getInt1(hasCacheHint));
4253 constexpr size_t numDims = 5;
4254 constexpr size_t numModes = 5;
4255 using rowTy = std::array<llvm::Intrinsic::ID, numDims + 1>;
4256 using TableTy = std::array<rowTy, numModes>;
4257 static constexpr TableTy IDTable{
4258 {{
notIntrinsic, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_1d,
4259 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_2d,
4260 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_3d,
4261 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_4d,
4262 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_5d},
4264 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_3d,
4265 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_4d,
4266 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_5d},
4268 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_3d,
4269 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_4d,
4270 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_5d},
4272 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_3d,
4273 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_4d,
4274 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_im2col_w_128_5d},
4276 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_tile_gather4_2d}}};
4278 static constexpr TableTy IDTableCTA{
4280 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_1d,
4281 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_2d,
4282 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_3d,
4283 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_4d,
4284 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_5d},
4286 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_3d,
4287 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_4d,
4288 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_5d},
4290 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_3d,
4291 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_4d,
4292 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_5d},
4294 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_3d,
4295 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_4d,
4296 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_im2col_w_128_5d},
4298 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_g2s_cta_tile_gather4_2d}}};
4301 (getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1) &&
4302 (getMaxEnumValForTMALoadMode() == std::size(IDTableCTA) - 1),
4303 "TMALoadModes must match number of rows in IDTable and IDTableCTA");
4304 size_t mode =
static_cast<size_t>(thisOp.getMode());
4305 size_t dim = thisOp.getCoordinates().size();
4306 auto id = isCTAOnly ? IDTableCTA[mode][dim] : IDTable[mode][dim];
4308 "Invalid intrinsic for CpAsyncBulkTensorGlobalToSharedClusterOp.");
4310 return {id, std::move(args)};
4315 auto thisOp = cast<NVVM::CpAsyncBulkTensorPrefetchOp>(op);
4319 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4321 for (
auto v : thisOp.getCoordinates())
4323 for (
auto v : thisOp.getIm2colOffsets())
4327 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4328 llvm::Value *i64Unused =
4329 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4330 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4331 args.push_back(builder.getInt1(hasCacheHint));
4333 const unsigned NI = llvm::Intrinsic::not_intrinsic;
4334 static constexpr llvm::Intrinsic::ID IDTable[][6] = {
4335 {NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_1d,
4336 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_2d,
4337 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_3d,
4338 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_4d,
4339 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_5d},
4341 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_3d,
4342 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_4d,
4343 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_5d},
4345 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_3d,
4346 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_4d,
4347 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_5d},
4349 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_3d,
4350 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_4d,
4351 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_im2col_w_128_5d},
4352 {NI, NI, NI, NI, NI,
4353 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_prefetch_tile_gather4_2d}};
4355 static_assert(getMaxEnumValForTMALoadMode() == std::size(IDTable) - 1,
4356 "TMALoadModes must match number of rows in IDTable");
4357 size_t mode =
static_cast<size_t>(thisOp.getMode());
4358 size_t dim = thisOp.getCoordinates().size();
4359 llvm::Intrinsic::ID
id = IDTable[mode][dim];
4360 if (
id == llvm::Intrinsic::not_intrinsic)
4361 llvm_unreachable(
"Invalid intrinsic for CpAsyncBulkTensorPrefetchOp.");
4363 return {id, std::move(args)};
4367CpAsyncBulkTensorSharedCTAToGlobalOp::getIntrinsicIDAndArgs(
4369 auto thisOp = cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOp>(op);
4373 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4374 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4376 for (
auto v : thisOp.getCoordinates())
4380 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4381 llvm::Value *i64Unused =
4382 llvm::ConstantInt::get(llvm::Type::getInt64Ty(mt.
getLLVMContext()), 0);
4383 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64Unused);
4384 args.push_back(builder.getInt1(hasCacheHint));
4386 const unsigned NI = llvm::Intrinsic::not_intrinsic;
4387 static constexpr llvm::Intrinsic::ID IDTable[][6] = {
4388 {NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_1d,
4389 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_2d,
4390 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_3d,
4391 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_4d,
4392 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_5d},
4393 {NI, NI, NI, llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_im2col_3d,
4394 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_im2col_4d,
4395 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_im2col_5d},
4396 {NI, NI, NI, NI, NI,
4397 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_2d}};
4399 static_assert(getMaxEnumValForTMAStoreMode() == std::size(IDTable) - 1,
4400 "TMAStoreModes must match number of rows in IDTable");
4401 size_t mode =
static_cast<size_t>(thisOp.getMode());
4402 size_t dim = thisOp.getCoordinates().size();
4403 llvm::Intrinsic::ID
id = IDTable[mode][dim];
4404 if (
id == llvm::Intrinsic::not_intrinsic)
4406 "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOp.");
4408 return {id, std::move(args)};
4413 auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOp>(op);
4421 args.push_back(mt.
lookupValue(thisOp.getSrcMem()));
4422 args.push_back(mt.
lookupValue(thisOp.getTmaDescriptor()));
4424 for (
Value v : thisOp.getCoordinates())
4428 const bool hasCacheHint =
static_cast<bool>(cacheHint);
4429 llvm::Value *i64ZeroValue =
4430 llvm::ConstantInt::get(llvm::Type::getInt64Ty(ctx), 0);
4431 args.push_back(hasCacheHint ? mt.
lookupValue(cacheHint) : i64ZeroValue);
4432 args.push_back(builder.getInt1(hasCacheHint));
4434 const llvm::Intrinsic::ID
notIntrinsic = llvm::Intrinsic::not_intrinsic;
4436 constexpr unsigned numRedKinds = 8;
4437 constexpr unsigned numLayouts = 2;
4438 constexpr unsigned maxDim = 5;
4439 using row = std::array<llvm::Intrinsic::ID, maxDim + 1>;
4440 using layoutTable = std::array<row, numLayouts>;
4441 using fullTable = std::array<layoutTable, numRedKinds>;
4442 static constexpr fullTable IDTable{
4445 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_1d,
4446 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_2d,
4447 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_3d,
4448 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_4d,
4449 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_tile_5d}},
4451 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_3d,
4452 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_4d,
4453 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_add_im2col_5d}}}},
4456 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_1d,
4457 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_2d,
4458 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_3d,
4459 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_4d,
4460 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_tile_5d}},
4462 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_3d,
4463 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_4d,
4464 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_min_im2col_5d}}}},
4467 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_1d,
4468 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_2d,
4469 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_3d,
4470 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_4d,
4471 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_tile_5d}},
4473 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_3d,
4474 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_4d,
4475 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_max_im2col_5d}}}},
4478 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_1d,
4479 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_2d,
4480 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_3d,
4481 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_4d,
4482 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_tile_5d}},
4484 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_3d,
4485 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_4d,
4486 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_inc_im2col_5d}}}},
4489 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_1d,
4490 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_2d,
4491 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_3d,
4492 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_4d,
4493 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_tile_5d}},
4495 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_3d,
4496 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_4d,
4497 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_dec_im2col_5d}}}},
4500 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_1d,
4501 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_2d,
4502 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_3d,
4503 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_4d,
4504 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_tile_5d}},
4506 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_3d,
4507 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_4d,
4508 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_and_im2col_5d}}}},
4511 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_1d,
4512 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_2d,
4513 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_3d,
4514 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_4d,
4515 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_tile_5d}},
4517 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_3d,
4518 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_4d,
4519 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_or_im2col_5d}}}},
4522 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_1d,
4523 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_2d,
4524 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_3d,
4525 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_4d,
4526 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_tile_5d}},
4528 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_im2col_3d,
4529 llvm::Intrinsic::nvvm_cp_async_bulk_tensor_reduce_xor_im2col_4d,
4531 nvvm_cp_async_bulk_tensor_reduce_xor_im2col_5d}}}}}};
4533 static_assert(getMaxEnumValForTMAReduxKind() == std::size(IDTable) - 1,
4534 "TMAReduxKinds must match number of rows in IDTable");
4536 size_t redKind =
static_cast<size_t>(thisOp.getRedKind());
4537 size_t mode =
static_cast<size_t>(thisOp.getMode());
4538 size_t dim = thisOp.getCoordinates().size();
4540 assert(redKind < IDTable.size() &&
4541 "Invalid redKind for CpAsyncBulkTensorReduceOp");
4542 assert(mode < IDTable[redKind].size() &&
4543 "Invalid mode for CpAsyncBulkTensorReduceOp");
4544 assert(dim < IDTable[redKind][mode].size() &&
4545 "Invalid dim for CpAsyncBulkTensorReduceOp");
4547 llvm::Intrinsic::ID intrinsicID = IDTable[redKind][mode][dim];
4550 "Invalid intrinsic for CpAsyncBulkTensorReduceOp.");
4552 return {intrinsicID, std::move(args)};
4557#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
4558 hasRelu ? llvm::Intrinsic::nvvm_f2tf32_##rnd##relu##sf \
4559 : llvm::Intrinsic::nvvm_f2tf32_##rnd##sf
4561#define GET_CVT_F2TF32_ID(rnd, relu, sf) \
4562 hasSatFinite ? CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
4563 : CVT_F2TF32_ID_IMPL(rnd, relu, )
4566ConvertFloatToTF32Op::getIntrinsicID(NVVM::FPRoundingMode rnd,
4567 NVVM::SaturationMode sat,
bool hasRelu) {
4568 using RndMode = NVVM::FPRoundingMode;
4569 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
4578 llvm_unreachable(
"Invalid RoundingMode for CvtFloatToTF32Op");
4583ConvertF32x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF4x2Op op,
4585 llvm::IRBuilderBase &builder) {
4590 bool hasRelu = op.getRelu();
4592 llvm::Intrinsic::ID intId =
4593 hasRelu ? llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_relu_satfinite
4594 : llvm::Intrinsic::nvvm_ff_to_e2m1x2_rn_satfinite;
4596 return {intId, std::move(args)};
4599#define GET_F32x2_TO_F6x2_ID(type, has_relu) \
4600 has_relu ? llvm::Intrinsic::nvvm_ff_to_##type##_rn_relu_satfinite \
4601 : llvm::Intrinsic::nvvm_ff_to_##type##_rn_satfinite
4603llvm::Intrinsic::ID ConvertF32x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
4606 .Case([&](mlir::Float6E2M3FNType) {
4609 .Case([&](mlir::Float6E3M2FNType) {
4613 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF6x2Op");
4614 return llvm::Intrinsic::not_intrinsic;
4619ConvertF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF16x2ToF4x2Op &op,
4621 llvm::IRBuilderBase &builder) {
4623 bool hasRelu = op.getRelu();
4625 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
4627 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
4628 intId = hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_relu_satfinite
4629 : llvm::Intrinsic::nvvm_f16x2_to_e2m1x2_rn_satfinite;
4634 return {intId, std::move(args)};
4638ConvertBF16x2ToF4x2Op::getIntrinsicIDAndArgs(NVVM::ConvertBF16x2ToF4x2Op &op,
4640 llvm::IRBuilderBase &builder) {
4642 bool hasRelu = op.getRelu();
4644 llvm::Intrinsic::ID intId = llvm::Intrinsic::not_intrinsic;
4646 if (llvm::isa<mlir::Float4E2M1FNType>(dstTy))
4647 intId = hasRelu ? llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite
4648 : llvm::Intrinsic::nvvm_bf16x2_to_e2m1x2_rn_satfinite;
4653 return {intId, std::move(args)};
4656llvm::Intrinsic::ID ConvertF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
4659 .Case<mlir::Float6E2M3FNType>([&](mlir::Float6E2M3FNType) {
4660 return hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_relu_satfinite
4661 : llvm::Intrinsic::nvvm_f16x2_to_e2m3x2_rn_satfinite;
4663 .Case<mlir::Float6E3M2FNType>([&](mlir::Float6E3M2FNType) {
4664 return hasRelu ? llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_relu_satfinite
4665 : llvm::Intrinsic::nvvm_f16x2_to_e3m2x2_rn_satfinite;
4668 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF6x2Op");
4669 return llvm::Intrinsic::not_intrinsic;
4673llvm::Intrinsic::ID ConvertBF16x2ToF6x2Op::getIntrinsicID(
mlir::Type dstTy,
4676 .Case<mlir::Float6E2M3FNType>([&](mlir::Float6E2M3FNType) {
4678 ? llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite
4679 : llvm::Intrinsic::nvvm_bf16x2_to_e2m3x2_rn_satfinite;
4681 .Case<mlir::Float6E3M2FNType>([&](mlir::Float6E3M2FNType) {
4683 ? llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite
4684 : llvm::Intrinsic::nvvm_bf16x2_to_e3m2x2_rn_satfinite;
4687 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF6x2Op");
4688 return llvm::Intrinsic::not_intrinsic;
4692#define GET_F32x2_TO_F8X2_US_ID(rnd, has_satf) \
4693 has_satf ? llvm::Intrinsic::nvvm_ff_to_ue8m0x2_##rnd##_satfinite \
4694 : llvm::Intrinsic::nvvm_ff_to_ue8m0x2_##rnd
4696#define GET_F32x2_TO_F8X2_S_ID(type, has_relu) \
4697 has_relu ? llvm::Intrinsic::nvvm_ff_to_##type##_rn_relu \
4698 : llvm::Intrinsic::nvvm_ff_to_##type##_rn
4701ConvertF32x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy, NVVM::FPRoundingMode rnd,
4702 NVVM::SaturationMode sat,
bool hasRelu) {
4703 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
4704 bool hasRoundingModeRZ = (rnd == NVVM::FPRoundingMode::RZ);
4705 bool hasRoundingModeRP = (rnd == NVVM::FPRoundingMode::RP);
4708 .Case([&](mlir::Float8E4M3FNType) {
4711 .Case([&](mlir::Float8E5M2Type) {
4714 .Case([&](mlir::Float8E8M0FNUType) {
4715 if (hasRoundingModeRZ)
4717 else if (hasRoundingModeRP)
4720 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
4723 llvm_unreachable(
"Invalid conversion in ConvertF32x2ToF8x2Op");
4724 return llvm::Intrinsic::not_intrinsic;
4728#define GET_F16x2_TO_F8X2_ID(type, has_relu) \
4729 has_relu ? llvm::Intrinsic::nvvm_f16x2_to_##type##_rn_relu \
4730 : llvm::Intrinsic::nvvm_f16x2_to_##type##_rn
4732llvm::Intrinsic::ID ConvertF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
4735 .Case([&](mlir::Float8E4M3FNType) {
4738 .Case([&](mlir::Float8E5M2Type) {
4742 llvm_unreachable(
"Invalid conversion in ConvertF16x2ToF8x2Op");
4743 return llvm::Intrinsic::not_intrinsic;
4748ConvertBF16x2ToF8x2Op::getIntrinsicID(
mlir::Type dstTy,
4749 NVVM::FPRoundingMode rnd,
4750 NVVM::SaturationMode sat,
bool hasRelu) {
4751 bool hasSatFinite = (sat == NVVM::SaturationMode::SATFINITE);
4753 static constexpr llvm::Intrinsic::ID ue8m0x2IDs[] = {
4754 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz,
4755 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp,
4756 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rz_satfinite,
4757 llvm::Intrinsic::nvvm_bf16x2_to_ue8m0x2_rp_satfinite,
4761 .Case<mlir::Float8E4M3FNType>([&](mlir::Float8E4M3FNType) {
4763 ? llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite
4764 : llvm::Intrinsic::nvvm_bf16x2_to_e4m3x2_rn_satfinite;
4766 .Case<mlir::Float8E5M2Type>([&](mlir::Float8E5M2Type) {
4768 ? llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite
4769 : llvm::Intrinsic::nvvm_bf16x2_to_e5m2x2_rn_satfinite;
4771 .Case<mlir::Float8E8M0FNUType>([&](mlir::Float8E8M0FNUType) {
4772 bool hasRoundingModeRP = (rnd == NVVM::FPRoundingMode::RP);
4773 unsigned index = (hasSatFinite << 1) | hasRoundingModeRP;
4774 return ue8m0x2IDs[
index];
4777 llvm_unreachable(
"Invalid conversion in ConvertBF16x2ToF8x2Op");
4778 return llvm::Intrinsic::not_intrinsic;
4784 auto curOp = cast<NVVM::ConvertF8x2ToF16x2Op>(op);
4786 bool hasRelu = curOp.getRelu();
4788 llvm::Intrinsic::ID intId =
4790 .Case([&](Float8E4M3FNType type) {
4791 return hasRelu ? llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn_relu
4792 : llvm::Intrinsic::nvvm_e4m3x2_to_f16x2_rn;
4794 .Case([&](Float8E5M2Type type) {
4795 return hasRelu ? llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn_relu
4796 : llvm::Intrinsic::nvvm_e5m2x2_to_f16x2_rn;
4799 llvm_unreachable(
"Invalid type for ConvertF8x2ToF16x2Op");
4800 return llvm::Intrinsic::not_intrinsic;
4803 llvm::Value *packedI16 =
4804 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
4805 llvm::Type::getInt16Ty(builder.getContext()));
4807 return {intId, {packedI16}};
4812 auto curOp = cast<NVVM::ConvertF8x2ToBF16x2Op>(op);
4813 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
4814 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
4815 bool hasRelu = curOp.getRelu();
4817 static constexpr llvm::Intrinsic::ID E4M3Ids[] = {
4818 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_scale_n2_ue8m0,
4819 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
4820 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
4821 llvm::Intrinsic::nvvm_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
4824 static constexpr llvm::Intrinsic::ID E5M2Ids[] = {
4825 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_scale_n2_ue8m0,
4826 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
4827 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
4828 llvm::Intrinsic::nvvm_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
4831 llvm::Intrinsic::ID intId =
4833 .Case([&](Float8E8M0FNUType type) {
4834 return llvm::Intrinsic::nvvm_ue8m0x2_to_bf16x2;
4836 .Case([&](Float8E4M3FNType type) {
4837 return E4M3Ids[hasSatfinite << 1 | hasRelu];
4839 .Case([&](Float8E5M2Type type) {
4840 return E5M2Ids[hasSatfinite << 1 | hasRelu];
4843 llvm_unreachable(
"Invalid type for ConvertF8x2ToBF16x2Op");
4844 return llvm::Intrinsic::not_intrinsic;
4846 llvm::Value *packedI16 =
4847 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
4848 llvm::Type::getInt16Ty(builder.getContext()));
4851 args.push_back(packedI16);
4852 if (!isa<Float8E8M0FNUType>(curOp.getSrcType()))
4855 : builder.getInt16(0x7f7f));
4858 return {intId, std::move(args)};
4863 auto curOp = cast<NVVM::ConvertF6x2ToF16x2Op>(op);
4865 bool hasRelu = curOp.getRelu();
4867 llvm::Intrinsic::ID intId =
4869 .Case([&](Float6E2M3FNType type) {
4870 return hasRelu ? llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn_relu
4871 : llvm::Intrinsic::nvvm_e2m3x2_to_f16x2_rn;
4873 .Case([&](Float6E3M2FNType type) {
4874 return hasRelu ? llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn_relu
4875 : llvm::Intrinsic::nvvm_e3m2x2_to_f16x2_rn;
4878 llvm_unreachable(
"Invalid type for ConvertF6x2ToF16x2Op");
4879 return llvm::Intrinsic::not_intrinsic;
4882 llvm::Value *packedI16 =
4883 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
4884 llvm::Type::getInt16Ty(builder.getContext()));
4886 return {intId, {packedI16}};
4891 auto curOp = cast<NVVM::ConvertF6x2ToBF16x2Op>(op);
4892 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
4893 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
4894 bool hasRelu = curOp.getRelu();
4896 static constexpr llvm::Intrinsic::ID E2M3Ids[] = {
4897 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_scale_n2_ue8m0,
4898 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
4899 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
4900 llvm::Intrinsic::nvvm_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
4903 static constexpr llvm::Intrinsic::ID E3M2Ids[] = {
4904 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_scale_n2_ue8m0,
4905 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
4906 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
4907 llvm::Intrinsic::nvvm_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
4910 unsigned idx = (hasSatfinite << 1) | hasRelu;
4911 llvm::Intrinsic::ID intId =
4913 .Case([&](Float6E2M3FNType type) {
return E2M3Ids[idx]; })
4914 .Case([&](Float6E3M2FNType type) {
return E3M2Ids[idx]; })
4916 llvm_unreachable(
"Invalid type for ConvertF6x2ToBF16x2Op");
4917 return llvm::Intrinsic::not_intrinsic;
4920 llvm::Value *packedI16 =
4921 builder.CreateBitCast(mt.
lookupValue(curOp.getSrc()),
4922 llvm::Type::getInt16Ty(builder.getContext()));
4925 args.push_back(packedI16);
4932 return {intId, std::move(args)};
4937 auto curOp = cast<NVVM::ConvertF4x2ToF16x2Op>(op);
4939 bool hasRelu = curOp.getRelu();
4941 llvm::Intrinsic::ID intId =
4943 .Case([&](Float4E2M1FNType type) {
4944 return hasRelu ? llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn_relu
4945 : llvm::Intrinsic::nvvm_e2m1x2_to_f16x2_rn;
4948 llvm_unreachable(
"Invalid type for ConvertF4x2ToF16x2Op");
4949 return llvm::Intrinsic::not_intrinsic;
4952 llvm::Value *extendedI16 =
4953 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
4954 llvm::Type::getInt16Ty(builder.getContext()));
4956 return {intId, {extendedI16}};
4961 auto curOp = cast<NVVM::ConvertF4x2ToBF16x2Op>(op);
4962 bool hasScale =
static_cast<bool>(curOp.getScaleFactor());
4963 bool hasSatfinite = curOp.getSat() == NVVM::SaturationMode::SATFINITE;
4964 bool hasRelu = curOp.getRelu();
4966 static constexpr llvm::Intrinsic::ID E2M1Ids[] = {
4967 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_scale_n2_ue8m0,
4968 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
4969 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
4970 llvm::Intrinsic::nvvm_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
4973 unsigned idx = (hasSatfinite << 1) | hasRelu;
4974 llvm::Intrinsic::ID intId =
4976 .Case([&](Float4E2M1FNType type) {
return E2M1Ids[idx]; })
4978 llvm_unreachable(
"Invalid type for ConvertF4x2ToBF16x2Op");
4979 return llvm::Intrinsic::not_intrinsic;
4982 llvm::Value *extendedI16 =
4983 builder.CreateZExt(mt.
lookupValue(curOp.getSrc()),
4984 llvm::Type::getInt16Ty(builder.getContext()));
4987 args.push_back(extendedI16);
4994 return {intId, std::move(args)};
4999 auto thisOp = cast<NVVM::ConvertF32x2ToS2F6x2Op>(op);
5000 bool hasRelu = thisOp.getRelu();
5001 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5003 llvm::Intrinsic::ID
id =
5005 ? llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5006 : llvm::Intrinsic::nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5012 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5013 : builder.getInt16(0x7f7f));
5014 return {id, std::move(args)};
5019 auto thisOp = cast<NVVM::ConvertBF16x2ToS2F6x2Op>(op);
5020 bool hasRelu = thisOp.getRelu();
5021 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5023 llvm::Intrinsic::ID
id =
5026 nvvm_bf16x2_to_s2f6x2_rn_relu_satfinite_scale_n2_ue8m0
5027 : llvm::Intrinsic::nvvm_bf16x2_to_s2f6x2_rn_satfinite_scale_n2_ue8m0;
5032 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5033 : builder.getInt16(0x7f7f));
5034 return {id, std::move(args)};
5039 auto thisOp = cast<NVVM::ConvertS2F6x2ToBF16x2Op>(op);
5040 bool hasRelu = thisOp.getRelu();
5041 bool hasScale =
static_cast<bool>(thisOp.getScaleFactor());
5042 bool hasSat = thisOp.getSat() == NVVM::SaturationMode::SATFINITE;
5044 static constexpr llvm::Intrinsic::ID ids[] = {
5045 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_scale_n2_ue8m0,
5046 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_scale_n2_ue8m0,
5047 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0,
5048 llvm::Intrinsic::nvvm_s2f6x2_to_bf16x2_rn_relu_satfinite_scale_n2_ue8m0,
5051 unsigned idx = (hasSat << 1) | hasRelu;
5055 llvm::Value *packedI16 =
5056 builder.CreateBitCast(mt.
lookupValue(thisOp.getSrc()),
5057 llvm::Type::getInt16Ty(builder.getContext()));
5058 args.push_back(packedI16);
5059 args.push_back(hasScale ? mt.
lookupValue(thisOp.getScaleFactor())
5060 : builder.getInt16(0x7f7f));
5062 return {ids[idx], std::move(args)};
5066Tcgen05AllocOp::getIntrinsicIDAndArgs(
Operation &op,
5069 auto curOp = cast<NVVM::Tcgen05AllocOp>(op);
5070 unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
5072 bool isShared = as == NVVMMemorySpace::Shared;
5073 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5075 llvm::Intrinsic::ID id;
5077 id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_shared_cg2
5078 : llvm::Intrinsic::nvvm_tcgen05_alloc_shared_cg1;
5080 id = is2CTAMode ? llvm::Intrinsic::nvvm_tcgen05_alloc_cg2
5081 : llvm::Intrinsic::nvvm_tcgen05_alloc_cg1;
5091llvm::Intrinsic::ID Tcgen05DeallocOp::getIntrinsicIDAndArgs(
5094 auto curOp = cast<NVVM::Tcgen05DeallocOp>(op);
5095 auto id = (curOp.getGroup() == CTAGroupKind::CTA_1)
5096 ? llvm::Intrinsic::nvvm_tcgen05_dealloc_cg1
5097 : llvm::Intrinsic::nvvm_tcgen05_dealloc_cg2;
5106#define TCGEN05_COMMIT_IMPL(cg, is_shared, mc) \
5107 is_shared ? llvm::Intrinsic::nvvm_tcgen05_commit##mc##_shared##_##cg \
5108 : llvm::Intrinsic::nvvm_tcgen05_commit##mc##_##cg
5110#define GET_TCGEN05_COMMIT_ID(cta_group, is_shared, has_mc) \
5111 has_mc ? TCGEN05_COMMIT_IMPL(cta_group, is_shared, _mc) \
5112 : TCGEN05_COMMIT_IMPL(cta_group, is_shared, )
5115Tcgen05CommitOp::getIntrinsicIDAndArgs(
Operation &op,
5118 auto curOp = cast<NVVM::Tcgen05CommitOp>(op);
5119 unsigned as = llvm::cast<LLVM::LLVMPointerType>(curOp.getAddr().getType())
5121 bool isShared = as == NVVMMemorySpace::Shared;
5122 bool hasMulticast =
static_cast<bool>(curOp.getMulticastMask());
5123 bool is2CTAMode = curOp.getGroup() == CTAGroupKind::CTA_2;
5125 llvm::Intrinsic::ID
id =
5132 args.push_back(mt.
lookupValue(curOp.getMulticastMask()));
5137#define TCGEN05_CP_IMPL(shape_mc, src_fmt, cg) \
5138 llvm::Intrinsic::nvvm_tcgen05_cp##shape_mc##src_fmt##cg
5140#define TCGEN05_CP_2CTA(shape_mc, src_fmt, is_2cta) \
5141 is_2cta ? TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg2) \
5142 : TCGEN05_CP_IMPL(shape_mc, src_fmt, _cg1)
5144#define GET_TCGEN05_CP_ID(shape_mc, src_fmt, is_2cta) \
5146 if ((src_fmt) == Tcgen05CpSrcFormat::B6x16_P32) \
5147 return TCGEN05_CP_2CTA(shape_mc, _b6x16_p32, is_2cta); \
5148 if ((src_fmt) == Tcgen05CpSrcFormat::B4x16_P64) \
5149 return TCGEN05_CP_2CTA(shape_mc, _b4x16_p64, is_2cta); \
5150 return TCGEN05_CP_2CTA(shape_mc, , is_2cta); \
5154ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
5156 llvm::IRBuilderBase &builder) {
5157 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5158 llvm::Intrinsic::nvvm_ff2f16x2_rn,
5159 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu,
5160 llvm::Intrinsic::nvvm_ff2f16x2_rn_satfinite,
5161 llvm::Intrinsic::nvvm_ff2f16x2_rn_relu_satfinite,
5163 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5164 llvm::Intrinsic::nvvm_ff2f16x2_rz,
5165 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu,
5166 llvm::Intrinsic::nvvm_ff2f16x2_rz_satfinite,
5167 llvm::Intrinsic::nvvm_ff2f16x2_rz_relu_satfinite,
5169 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5170 llvm::Intrinsic::nvvm_ff2f16x2_rs,
5171 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu,
5172 llvm::Intrinsic::nvvm_ff2f16x2_rs_satfinite,
5173 llvm::Intrinsic::nvvm_ff2f16x2_rs_relu_satfinite,
5176 unsigned hasRelu = op.getRelu() ? 1 : 0;
5177 unsigned hasSatFinite =
5178 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5181 unsigned idx = (hasSatFinite << 1) | hasRelu;
5186 if (op.getRandomBits())
5187 args.push_back(mt.
lookupValue(op.getRandomBits()));
5189 switch (op.getRnd()) {
5190 case FPRoundingMode::RN:
5191 return {rndRNIds[idx], std::move(args)};
5192 case FPRoundingMode::RZ:
5193 return {rndRZIds[idx], std::move(args)};
5194 case FPRoundingMode::RS:
5195 return {rndRSIds[idx], std::move(args)};
5197 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToF16x2Op");
5202ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
5204 llvm::IRBuilderBase &builder) {
5205 static constexpr llvm::Intrinsic::ID rndRNIds[] = {
5206 llvm::Intrinsic::nvvm_ff2bf16x2_rn,
5207 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu,
5208 llvm::Intrinsic::nvvm_ff2bf16x2_rn_satfinite,
5209 llvm::Intrinsic::nvvm_ff2bf16x2_rn_relu_satfinite,
5211 static constexpr llvm::Intrinsic::ID rndRZIds[] = {
5212 llvm::Intrinsic::nvvm_ff2bf16x2_rz,
5213 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu,
5214 llvm::Intrinsic::nvvm_ff2bf16x2_rz_satfinite,
5215 llvm::Intrinsic::nvvm_ff2bf16x2_rz_relu_satfinite,
5217 static constexpr llvm::Intrinsic::ID rndRSIds[] = {
5218 llvm::Intrinsic::nvvm_ff2bf16x2_rs,
5219 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu,
5220 llvm::Intrinsic::nvvm_ff2bf16x2_rs_satfinite,
5221 llvm::Intrinsic::nvvm_ff2bf16x2_rs_relu_satfinite,
5224 unsigned hasRelu = op.getRelu() ? 1 : 0;
5225 unsigned hasSatFinite =
5226 (op.getSat() == NVVM::SaturationMode::SATFINITE) ? 1 : 0;
5229 unsigned idx = (hasSatFinite << 1) | hasRelu;
5234 if (op.getRandomBits())
5235 args.push_back(mt.
lookupValue(op.getRandomBits()));
5237 switch (op.getRnd()) {
5238 case FPRoundingMode::RN:
5239 return {rndRNIds[idx], std::move(args)};
5240 case FPRoundingMode::RZ:
5241 return {rndRZIds[idx], std::move(args)};
5242 case FPRoundingMode::RS:
5243 return {rndRSIds[idx], std::move(args)};
5245 llvm_unreachable(
"Invalid rounding mode for ConvertF32x2ToBF16x2Op");
5249llvm::Intrinsic::ID ConvertF32x4ToF8x4Op::getIntrinsicID() {
5251 bool hasRelu = getRelu();
5254 .Case([&](mlir::Float8E4M3FNType) {
5255 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_relu_satfinite
5256 : llvm::Intrinsic::nvvm_f32x4_to_e4m3x4_rs_satfinite;
5258 .Case([&](mlir::Float8E5M2Type) {
5259 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_relu_satfinite
5260 : llvm::Intrinsic::nvvm_f32x4_to_e5m2x4_rs_satfinite;
5263 llvm_unreachable(
"Invalid F8 type in ConvertF32x4ToF8x4Op");
5264 return llvm::Intrinsic::not_intrinsic;
5268llvm::Intrinsic::ID ConvertF32x4ToF6x4Op::getIntrinsicID() {
5270 bool hasRelu = getRelu();
5273 .Case([&](mlir::Float6E2M3FNType) {
5274 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_relu_satfinite
5275 : llvm::Intrinsic::nvvm_f32x4_to_e2m3x4_rs_satfinite;
5277 .Case([&](mlir::Float6E3M2FNType) {
5278 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_relu_satfinite
5279 : llvm::Intrinsic::nvvm_f32x4_to_e3m2x4_rs_satfinite;
5282 llvm_unreachable(
"Invalid F6 type in ConvertF32x4ToF6x4Op");
5283 return llvm::Intrinsic::not_intrinsic;
5287llvm::Intrinsic::ID ConvertF32x4ToF4x4Op::getIntrinsicID() {
5289 bool hasRelu = getRelu();
5292 .Case([&](mlir::Float4E2M1FNType) {
5293 return hasRelu ? llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_relu_satfinite
5294 : llvm::Intrinsic::nvvm_f32x4_to_e2m1x4_rs_satfinite;
5297 llvm_unreachable(
"Invalid F4 type in ConvertF32x4ToF4x4Op");
5298 return llvm::Intrinsic::not_intrinsic;
5302llvm::Intrinsic::ID Tcgen05CpOp::getIntrinsicID(
Operation &op) {
5303 auto curOp = cast<NVVM::Tcgen05CpOp>(op);
5304 bool is2CTA = curOp.getGroup() == CTAGroupKind::CTA_2;
5305 auto srcFmt = curOp.getSrcFormat();
5306 auto mc = curOp.getMulticast();
5308 switch (curOp.getShape()) {
5309 case Tcgen05CpShape::SHAPE_128x256b:
5311 case Tcgen05CpShape::SHAPE_128x128b:
5313 case Tcgen05CpShape::SHAPE_4x256b:
5315 case Tcgen05CpShape::SHAPE_32x128b:
5317 case Tcgen05CpShape::SHAPE_64x128b:
5318 return (mc == Tcgen05CpMulticast::WARPX2_01_23)
5322 llvm_unreachable(
"Invalid shape in tcgen05 cp Op");
5329 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X128B)
5331 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X256B)
5336LogicalResult Tcgen05LdOp::verify() {
5338 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5341 if (
getShape() != NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && getOffset())
5342 result =
emitError(
"offset argument is only supported for shape 16x32bx2");
5344 auto resTy = getRes().getType();
5345 unsigned resLen = isa<VectorType>(resTy)
5346 ? llvm::cast<VectorType>(resTy).getNumElements()
5349 result =
emitError(llvm::formatv(
"invalid result type length {0} for shape "
5350 "{1} in tcgen05.ld Op",
5351 resLen, stringifyEnum(
getShape())));
5356LogicalResult Tcgen05StOp::verify() {
5358 if (
getShape() == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2 && !getOffset())
5361 auto valTy = getVal().getType();
5362 unsigned valLen = isa<VectorType>(valTy)
5363 ? llvm::cast<VectorType>(valTy).getNumElements()
5366 result =
emitError(llvm::formatv(
"invalid input length {0} for shape "
5367 "{1} in tcgen05.st Op",
5368 valLen, stringifyEnum(
getShape())));
5378 if (
auto rangeAttr = op->
getAttrOfType<LLVM::ConstantRangeAttr>(
"range")) {
5379 setResultRanges(
result, {rangeAttr.getLower(), rangeAttr.getUpper(),
5380 rangeAttr.getLower(), rangeAttr.getUpper()});
5390 std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
5394 const llvm::APInt &lower = rangeAttr->getLower();
5395 const llvm::APInt &upper = rangeAttr->getUpper();
5398 if (lower == upper && !lower.isMaxValue() && !lower.isMinValue()) {
5399 unsigned bitWidth = lower.getBitWidth();
5400 llvm::APInt minVal = llvm::APInt::getMinValue(bitWidth);
5401 llvm::APInt maxVal = llvm::APInt::getMaxValue(bitWidth);
5403 "invalid range attribute: Lower == Upper, but they aren't min (")
5404 << llvm::toString(minVal, 10,
false) <<
") or max ("
5405 << llvm::toString(maxVal, 10,
false)
5406 <<
") value! This is an invalid constant range.";
5413 llvm::IRBuilderBase &builder) {
5414 return builder.CreateBitCast(arg,
5415 llvm::Type::getInt32Ty(builder.getContext()));
5420 auto curOp = cast<NVVM::DotAccumulate4WayOp>(op);
5427 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
5428 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
5429 unsigned type = (isASigned << 1) | isBSigned;
5430 const llvm::Intrinsic::ID ids[] = {
5431 llvm::Intrinsic::nvvm_idp4a_u_u,
5432 llvm::Intrinsic::nvvm_idp4a_u_s,
5433 llvm::Intrinsic::nvvm_idp4a_s_u,
5434 llvm::Intrinsic::nvvm_idp4a_s_s,
5436 return {ids[type], args};
5441 auto curOp = cast<NVVM::DotAccumulate2WayOp>(op);
5446 args.push_back(builder.getInt1(curOp.getBHi()));
5449 bool isASigned = curOp.getAType() == NVVM::DotAccumulateType::SIGNED;
5450 bool isBSigned = curOp.getBType() == NVVM::DotAccumulateType::SIGNED;
5451 unsigned type = (isASigned << 1) | isBSigned;
5452 const llvm::Intrinsic::ID ids[] = {
5453 llvm::Intrinsic::nvvm_idp2a_u_u,
5454 llvm::Intrinsic::nvvm_idp2a_u_s,
5455 llvm::Intrinsic::nvvm_idp2a_s_u,
5456 llvm::Intrinsic::nvvm_idp2a_s_s,
5458 return {ids[type], args};
5462 llvm::IRBuilderBase &builder) {
5463 return builder.CreateAddrSpaceCast(
5464 addr, builder.getPtrTy(llvm::NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM));
5468PrefetchOp::getIntrinsicIDAndArgs(NVVM::PrefetchOp &op,
5470 llvm::IRBuilderBase &builder) {
5471 using MemSpace = NVVM::NVVMMemorySpace;
5472 using CacheLevel = NVVM::PrefetchCacheLevel;
5474 std::optional<NVVM::PrefetchCacheLevel> cacheLevel = op.getCacheLevel();
5475 std::optional<NVVM::CacheEvictionPriority> evictPriority =
5476 op.getEvictPriority();
5477 unsigned addressSpace =
5478 llvm::cast<LLVM::LLVMPointerType>(op.getAddr().getType())
5486 if (op.getTensormap())
5487 return {llvm::Intrinsic::nvvm_prefetch_tensormap, args};
5489 assert(cacheLevel &&
"expected cache level for non-tensormap prefetch");
5491 if (op.getUniform() && *cacheLevel == CacheLevel::L1)
5492 return {llvm::Intrinsic::nvvm_prefetchu_L1, args};
5494 if (evictPriority && *cacheLevel == CacheLevel::L2) {
5495 switch (*evictPriority) {
5496 case NVVM::CacheEvictionPriority::EvictLast:
5497 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_last, args};
5498 case NVVM::CacheEvictionPriority::EvictNormal:
5499 return {llvm::Intrinsic::nvvm_prefetch_global_L2_evict_normal, args};
5501 llvm_unreachable(
"Invalid cache eviction priority");
5505 switch (
static_cast<MemSpace
>(addressSpace)) {
5506 case MemSpace::Generic:
5507 return *cacheLevel == CacheLevel::L1
5509 :
NVVM::
IDArgPair({llvm::Intrinsic::nvvm_prefetch_L2, args});
5510 case MemSpace::Global:
5511 return *cacheLevel == CacheLevel::L1
5513 {llvm::Intrinsic::nvvm_prefetch_global_L1, args})
5515 {llvm::Intrinsic::nvvm_prefetch_global_L2, args});
5516 case MemSpace::Local:
5517 return *cacheLevel == CacheLevel::L1
5519 {llvm::Intrinsic::nvvm_prefetch_local_L1, args})
5521 {llvm::Intrinsic::nvvm_prefetch_local_L2, args});
5523 llvm_unreachable(
"Invalid pointer address space");
5527bool NVVM::InlinePtxOp::getAsmValues(
5531 for (
auto arg : getReadWriteArgs())
5533 for (
auto arg : getResults())
5535 for (
auto arg : getReadOnlyArgs())
5542NVVM::IDArgPair ClusterLaunchControlTryCancelOp::getIntrinsicIDAndArgs(
5544 auto curOp = cast<NVVM::ClusterLaunchControlTryCancelOp>(op);
5546 args.push_back(mt.
lookupValue(curOp.getSmemAddress()));
5547 args.push_back(mt.
lookupValue(curOp.getMbarrier()));
5549 llvm::Intrinsic::ID intrinsicID =
5550 curOp.getMulticast()
5552 nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
5553 : llvm::Intrinsic::nvvm_clusterlaunchcontrol_try_cancel_async_shared;
5555 return {intrinsicID, args};
5558NVVM::IDArgPair ClusterLaunchControlQueryCancelOp::getIntrinsicIDAndArgs(
5560 auto curOp = cast<NVVM::ClusterLaunchControlQueryCancelOp>(op);
5562 args.push_back(mt.
lookupValue(curOp.getTryCancelResponse()));
5564 llvm::Intrinsic::ID intrinsicID;
5566 switch (curOp.getQueryType()) {
5567 case NVVM::ClusterLaunchControlQueryType::IS_CANCELED:
5569 llvm::Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled;
5571 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_X:
5572 intrinsicID = llvm::Intrinsic::
5573 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_x;
5575 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Y:
5576 intrinsicID = llvm::Intrinsic::
5577 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_y;
5579 case NVVM::ClusterLaunchControlQueryType::GET_FIRST_CTA_ID_Z:
5580 intrinsicID = llvm::Intrinsic::
5581 nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_z;
5584 return {intrinsicID, args};
5589 llvm::IRBuilderBase &builder) {
5590 auto thisOp = cast<NVVM::PermuteOp>(op);
5591 NVVM::PermuteMode mode = thisOp.getMode();
5593 static constexpr llvm::Intrinsic::ID IDs[] = {
5594 llvm::Intrinsic::nvvm_prmt, llvm::Intrinsic::nvvm_prmt_f4e,
5595 llvm::Intrinsic::nvvm_prmt_b4e, llvm::Intrinsic::nvvm_prmt_rc8,
5596 llvm::Intrinsic::nvvm_prmt_ecl, llvm::Intrinsic::nvvm_prmt_ecr,
5597 llvm::Intrinsic::nvvm_prmt_rc16};
5599 unsigned modeIndex =
static_cast<unsigned>(mode);
5607 args.push_back(mt.
lookupValue(thisOp.getSelector()));
5609 return {IDs[modeIndex], args};
5614 auto thisOp = cast<NVVM::TensormapReplaceOp>(op);
5618 if (thisOp.getOrd())
5619 args.push_back(builder.getInt32(thisOp.getOrd().value()));
5620 if (thisOp.getNewValue())
5621 args.push_back(mt.
lookupValue(thisOp.getNewValue()));
5622 if (
auto attr = thisOp.getNewValueAttr()) {
5625 .Case<TensormapElemtypeAttr, TensormapInterleaveLayoutAttr,
5626 TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
5627 TensormapFillModeAttr>([](
auto attr) {
5628 return static_cast<unsigned>(attr.getValue());
5630 .Default([](
auto attr) {
5631 llvm_unreachable(
"Invalid attribute type");
5634 args.push_back(builder.getInt32(val));
5637 static constexpr llvm::Intrinsic::ID IDs[] = {
5638 llvm::Intrinsic::nvvm_tensormap_replace_global_address,
5639 llvm::Intrinsic::nvvm_tensormap_replace_rank,
5640 llvm::Intrinsic::nvvm_tensormap_replace_box_dim,
5641 llvm::Intrinsic::nvvm_tensormap_replace_global_dim,
5642 llvm::Intrinsic::nvvm_tensormap_replace_global_stride,
5643 llvm::Intrinsic::nvvm_tensormap_replace_element_stride,
5644 llvm::Intrinsic::nvvm_tensormap_replace_elemtype,
5645 llvm::Intrinsic::nvvm_tensormap_replace_interleave_layout,
5646 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_mode,
5647 llvm::Intrinsic::nvvm_tensormap_replace_swizzle_atomicity,
5648 llvm::Intrinsic::nvvm_tensormap_replace_fill_mode,
5651 unsigned fieldIndex =
static_cast<unsigned>(thisOp.getField());
5653 return {IDs[fieldIndex], args};
5662 llvm::IRBuilderBase &builder) {
5664 auto thisOp = cast<NVVM::Tcgen05MMAOp>(op);
5667 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
5670 const bool isATensor = isa<llvm::PointerType>(
A->getType());
5673 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
5674 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
5675 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
5677 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
5678 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
5679 using IsATensorArray = std::array<CtaGroupArray, 2>;
5680 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
5681 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
5684 static constexpr HasDisableOutputLaneArray tcgen05MMAIDs = {
5690 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic},
5692 {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
notIntrinsic}}},
5696 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
5697 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
5701 llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
5702 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift,
5708 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic},
5710 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
notIntrinsic}}},
5714 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
5715 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
5719 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
5720 llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift,
5726 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1,
5729 {llvm::Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2,
5734 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1,
5736 nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift,
5741 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2,
5743 nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift,
5749 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1,
5753 nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2,
5758 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1,
5760 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift},
5764 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2,
5766 nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift,
5769 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
5770 bool hasScaleInputD = ScaleInputD !=
nullptr;
5772 llvm::Value *DisableOutputLane =
5774 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
5776 const unsigned ctaGroup =
5779 llvm::Intrinsic::ID ID =
5780 tcgen05MMAIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
5781 [ctaGroup - 1][thisOp.getAShift()];
5783 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMAOp.");
5786 args.push_back(ScaleInputD);
5788 if (hasDisableOutputLane)
5789 args.push_back(DisableOutputLane);
5791 args.push_back(builder.getInt32(
static_cast<unsigned>(thisOp.getKind())));
5793 if (!hasDisableOutputLane)
5794 args.push_back(builder.getInt32(ctaGroup));
5797 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
5804 NVVM::CTAGroupKind ctaGroup,
bool hasAShift,
5805 NVVM::Tcgen05MMACollectorOp collectorOp,
Location loc) {
5807 if (disableOutputLane) {
5808 mlir::VectorType disableOutputLaneType =
5809 cast<mlir::VectorType>(disableOutputLane.
getType());
5810 if ((ctaGroup == NVVM::CTAGroupKind::CTA_1 &&
5811 disableOutputLaneType.getNumElements() != 4) ||
5812 (ctaGroup == NVVM::CTAGroupKind::CTA_2 &&
5813 disableOutputLaneType.getNumElements() != 8))
5814 return emitError(loc) <<
"Disable Output Lane of length "
5815 << disableOutputLaneType.getNumElements()
5816 <<
" is incompatible with CtaGroupAttr";
5819 if (hasAShift && !isATensor)
5821 loc,
"A-shift can be applied only when matrix A is in tensor memory");
5823 if (hasAShift ==
true && (collectorOp == Tcgen05MMACollectorOp::FILL ||
5824 collectorOp == Tcgen05MMACollectorOp::USE))
5826 loc,
"Cannot use collector buffer operation fill or use with ashift");
5831LogicalResult Tcgen05MMAOp::verify() {
5833 getDisableOutputLane(), getCtaGroup(), getAShift(),
5834 getCollectorOp(), getLoc());
5844 auto thisOp = cast<NVVM::Tcgen05MMASparseOp>(op);
5847 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
5850 bool isATensor = isa<llvm::PointerType>(
A->getType());
5853 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
5854 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
5855 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
5856 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
5858 using EnableAShiftArray = std::array<llvm::Intrinsic::ID, 2>;
5859 using CtaGroupArray = std::array<EnableAShiftArray, 2>;
5860 using IsATensorArray = std::array<CtaGroupArray, 2>;
5861 using HasScaleInputDArray = std::array<IsATensorArray, 2>;
5862 using HasDisableOutputLaneArray = std::array<HasScaleInputDArray, 2>;
5865 static constexpr HasDisableOutputLaneArray tcgen05MMASparseIDs = {
5871 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic},
5873 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
notIntrinsic}}},
5877 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
5878 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
5882 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
5883 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift,
5889 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
5892 {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
5897 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
5898 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
5902 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
5903 llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift,
5910 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1,
5914 nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2,
5919 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1,
5921 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift,
5926 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2,
5928 nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift,
5934 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1,
5938 nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2,
5943 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1,
5945 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift},
5949 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2,
5951 nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift,
5954 llvm::Value *ScaleInputD = mt.
lookupValue(thisOp.getScaleInputD());
5955 bool hasScaleInputD = ScaleInputD !=
nullptr;
5957 llvm::Value *DisableOutputLane =
5959 bool hasDisableOutputLane = DisableOutputLane !=
nullptr;
5964 llvm::Intrinsic::ID ID =
5965 tcgen05MMASparseIDs[hasDisableOutputLane][hasScaleInputD][isATensor]
5966 [ctaGroup - 1][thisOp.getAShift()];
5968 assert(ID !=
notIntrinsic &&
"Invalid intrinsic for Tcgen05MMASparseOp.");
5971 args.push_back(ScaleInputD);
5973 if (hasDisableOutputLane)
5974 args.push_back(DisableOutputLane);
5976 args.push_back(builder.getInt32(
static_cast<unsigned>(thisOp.getKind())));
5978 if (!hasDisableOutputLane)
5979 args.push_back(builder.getInt32(ctaGroup));
5982 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
5987LogicalResult Tcgen05MMASparseOp::verify() {
5989 getDisableOutputLane(), getCtaGroup(), getAShift(),
5990 getCollectorOp(), getLoc());
6000 auto thisOp = cast<NVVM::Tcgen05MMABlockScaleOp>(op);
6003 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6006 bool isATensor = isa<llvm::PointerType>(
A->getType());
6009 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6010 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6011 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6012 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6013 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6014 args.push_back(builder.getInt32(
6017 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6019 auto kind = thisOp.getKind();
6020 auto blockScale = thisOp.getBlockScale();
6021 llvm::Intrinsic::ID ID = [&]() {
6022 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6023 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6024 return isATensor ? llvm::Intrinsic::
6025 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale
6027 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale;
6028 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6031 nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32
6033 nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32;
6035 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6036 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6038 ? llvm::Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale
6039 : llvm::Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale;
6040 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6041 return isATensor ? llvm::Intrinsic::
6042 nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32
6044 nvvm_tcgen05_mma_shared_mxf4_block_scale_block32;
6046 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6047 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6050 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32
6052 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32;
6054 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6057 nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16
6059 nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16;
6062 llvm_unreachable(
"Invalid tcgen05.mma.block_scale attributes");
6069 NVVM::Tcgen05MMACollectorOp collectorOp, NVVM::Tcgen05MMAKind kind,
6070 NVVM::Tcgen05MMABlockScale blockScale,
Location loc) {
6071 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT &&
6072 kind == NVVM::Tcgen05MMAKind::MXF4NVF4)
6073 return emitError(loc,
"mxf4nvf4 requires block scale attribute");
6075 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16 &&
6076 kind != NVVM::Tcgen05MMAKind::MXF4NVF4)
6078 llvm::formatv(
"{} kind does not support block16 attribute",
6079 stringifyEnum(kind)));
6084LogicalResult Tcgen05MMABlockScaleOp::verify() {
6086 getBlockScale(), getLoc());
6096 auto thisOp = cast<NVVM::Tcgen05MMASparseBlockScaleOp>(op);
6099 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6102 bool isATensor = isa<llvm::PointerType>(
A->getType());
6105 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6106 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6107 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6108 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6109 args.push_back(mt.
lookupValue(thisOp.getScaleA()));
6110 args.push_back(mt.
lookupValue(thisOp.getScaleB()));
6111 args.push_back(builder.getInt32(
6114 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6116 auto kind = thisOp.getKind();
6117 auto blockScale = thisOp.getBlockScale();
6118 llvm::Intrinsic::ID ID = [&]() {
6119 if (kind == NVVM::Tcgen05MMAKind::MXF8F6F4) {
6120 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6121 return isATensor ? llvm::Intrinsic::
6122 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale
6124 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale;
6125 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6128 nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32
6130 nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32;
6132 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4) {
6133 if (blockScale == NVVM::Tcgen05MMABlockScale::DEFAULT) {
6134 return isATensor ? llvm::Intrinsic::
6135 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale
6137 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale;
6138 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6141 nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32
6143 nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32;
6145 }
else if (kind == NVVM::Tcgen05MMAKind::MXF4NVF4) {
6146 if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK32) {
6149 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32
6151 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32;
6153 }
else if (blockScale == NVVM::Tcgen05MMABlockScale::BLOCK16) {
6156 nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16
6158 nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16;
6161 llvm_unreachable(
"Invalid tcgen05.mma.sp.block_scale attributes");
6167LogicalResult Tcgen05MMASparseBlockScaleOp::verify() {
6169 getBlockScale(), getLoc());
6179 auto thisOp = cast<NVVM::Tcgen05MMAWsOp>(op);
6182 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6185 bool isATensor = isa<llvm::PointerType>(
A->getType());
6188 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6189 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6190 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6192 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6196 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor_zero_col_mask
6197 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared_zero_col_mask;
6199 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_tensor
6200 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_shared;
6202 args.push_back(builder.getInt32(
static_cast<unsigned>(thisOp.getKind())));
6204 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6206 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6218 auto thisOp = cast<NVVM::Tcgen05MMAWsSparseOp>(op);
6221 args.push_back(mt.
lookupValue(thisOp.getMatrixD()));
6224 bool isATensor = isa<llvm::PointerType>(
A->getType());
6227 args.push_back(mt.
lookupValue(thisOp.getMatrixB()));
6228 args.push_back(mt.
lookupValue(thisOp.getIdesc()));
6229 args.push_back(mt.
lookupValue(thisOp.getEnableInputD()));
6230 args.push_back(mt.
lookupValue(thisOp.getSparseMetadata()));
6232 mlir::Value ZeroColMask = thisOp.getZeroColMask();
6237 ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor_zero_col_mask
6238 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared_zero_col_mask;
6240 ID = isATensor ? llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_tensor
6241 : llvm::Intrinsic::nvvm_tcgen05_mma_ws_sp_shared;
6243 args.push_back(builder.getInt32(
static_cast<unsigned>(thisOp.getKind())));
6245 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorBBuffer())));
6247 builder.getInt32(
static_cast<unsigned>(thisOp.getCollectorOp())));
6256#define TCGEN05LDRED(SHAPE, NUM, TYPE) \
6257 llvm::Intrinsic::nvvm_tcgen05_ld_red_##SHAPE##_##NUM##_##TYPE
6261 auto thisOp = cast<NVVM::Tcgen05LdRedOp>(op);
6264 mlir::VectorType VecResTy =
6265 cast<mlir::VectorType>(thisOp.getData().getType());
6266 unsigned Num = VecResTy.getNumElements();
6267 bool IsFloat = thisOp.getRedVal().getType().isF32();
6269 llvm::Intrinsic::ID Shape32x32b[][2] = {
6280 llvm::Intrinsic::ID Shape16x32bx2[][2] = {
6291 NVVM::Tcgen05LdStShape
shape = thisOp.getShape();
6292 unsigned ID = [&]() {
6295 unsigned idx = std::log2(Num);
6297 case NVVM::Tcgen05LdStShape::SHAPE_32X32B:
6298 return Shape32x32b[idx][IsFloat];
6299 case NVVM::Tcgen05LdStShape::SHAPE_16X32BX2:
6300 return Shape16x32bx2[idx][IsFloat];
6302 llvm_unreachable(
"unhandled tcgen05.ld lowering");
6308 if (
shape == NVVM::Tcgen05LdStShape::SHAPE_16X32BX2)
6309 args.push_back(mt.
lookupValue(thisOp.getOffset()));
6312 builder.getInt32(thisOp.getOp() == NVVM::ReductionKind::MIN ? 0 : 1));
6315 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getAbs())));
6316 args.push_back(builder.getInt1(
static_cast<unsigned>(thisOp.getNan())));
6321LogicalResult Tcgen05LdRedOp::verify() {
6322 VectorType data = cast<VectorType>(getData().
getType());
6323 Type redVal = getRedVal().getType();
6325 if (data.getElementType() != redVal)
6327 "type of reduction value and element type of vector data should match");
6329 if (getOp() != NVVM::ReductionKind::MIN &&
6330 getOp() != NVVM::ReductionKind::MAX)
6331 return emitError(
"only min and max reduction kinds are supported");
6333 if (redVal.
isInteger() && (getAbs() || getNan())) {
6334 return emitError(
"abs or nan is only applicable for f32 type");
6344struct NVVMInlinerInterface final : DialectInlinerInterface {
6345 using DialectInlinerInterface::DialectInlinerInterface;
6346 bool isLegalToInline(Operation *, Region *,
bool, IRMapping &)
const final {
6353void NVVMDialect::initialize() {
6356#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
6359#define GET_ATTRDEF_LIST
6360#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
6365 allowUnknownOperations();
6366 addInterfaces<NVVMInlinerInterface>();
6367 declarePromisedInterface<ConvertToLLVMPatternInterface, NVVMDialect>();
6368 declarePromisedInterface<gpu::TargetAttrInterface, NVVMTargetAttr>();
6371LogicalResult NVVMDialect::verifyOperationAttribute(
Operation *op,
6373 StringAttr attrName = attr.
getName();
6375 if (attrName == NVVMDialect::getKernelFuncAttrName()) {
6376 if (!isa<LLVM::LLVMFuncOp>(op)) {
6377 return op->
emitError() <<
"'" << NVVMDialect::getKernelFuncAttrName()
6378 <<
"' attribute attached to unexpected op";
6383 if (attrName == NVVMDialect::getMaxntidAttrName() ||
6384 attrName == NVVMDialect::getReqntidAttrName() ||
6385 attrName == NVVMDialect::getClusterDimAttrName()) {
6386 auto values = llvm::dyn_cast<DenseI32ArrayAttr>(attr.
getValue());
6387 if (!values || values.empty() || values.size() > 3) {
6390 <<
"' attribute must be integer array with maximum 3 index";
6395 if (attrName == NVVMDialect::getMinctasmAttrName() ||
6396 attrName == NVVMDialect::getMaxnregAttrName() ||
6397 attrName == NVVMDialect::getClusterMaxBlocksAttrName()) {
6398 if (!llvm::dyn_cast<IntegerAttr>(attr.
getValue())) {
6400 <<
"'" << attrName <<
"' attribute must be integer constant";
6404 if (attrName == NVVMDialect::getBlocksAreClustersAttrName()) {
6405 if (!op->
hasAttr(NVVMDialect::getReqntidAttrName()) ||
6406 !op->
hasAttr(NVVMDialect::getClusterDimAttrName())) {
6408 <<
"'" << attrName <<
"' attribute must be used along with " <<
"'"
6409 << NVVMDialect::getReqntidAttrName() <<
"' and " <<
"'"
6410 << NVVMDialect::getClusterDimAttrName() <<
"'";
6417LogicalResult NVVMDialect::verifyRegionArgAttribute(
Operation *op,
6418 unsigned regionIndex,
6421 auto funcOp = dyn_cast<FunctionOpInterface>(op);
6425 bool isKernel = op->
hasAttr(NVVMDialect::getKernelFuncAttrName());
6426 StringAttr attrName = argAttr.
getName();
6427 if (attrName == NVVM::NVVMDialect::getGridConstantAttrName()) {
6431 <<
"' attribute must be present only on kernel arguments";
6433 if (!isa<UnitAttr>(argAttr.
getValue()))
6434 return op->
emitError() <<
"'" << attrName <<
"' must be a unit attribute";
6435 if (!funcOp.getArgAttr(argIndex, LLVM::LLVMDialect::getByValAttrName())) {
6438 <<
"' attribute requires the argument to also have attribute '"
6439 << LLVM::LLVMDialect::getByValAttrName() <<
"'";
6450unsigned NVVMMemorySpaceAttr::getAddressSpace()
const {
6451 return static_cast<unsigned>(getValue());
6454bool NVVMMemorySpaceAttr::isValidLoad(
6455 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
6456 const ::mlir::DataLayout *dataLayout,
6462bool NVVMMemorySpaceAttr::isValidStore(
6463 Type type, ptr::AtomicOrdering ordering, std::optional<int64_t> alignment,
6464 const ::mlir::DataLayout *dataLayout,
6470bool NVVMMemorySpaceAttr::isValidAtomicOp(
6471 ptr::AtomicBinOp op,
Type type, ptr::AtomicOrdering ordering,
6472 std::optional<int64_t> alignment, const ::mlir::DataLayout *dataLayout,
6475 assert(
false &&
"unimplemented, see TODO in the source.");
6479bool NVVMMemorySpaceAttr::isValidAtomicXchg(
6480 Type type, ptr::AtomicOrdering successOrdering,
6481 ptr::AtomicOrdering failureOrdering, std::optional<int64_t> alignment,
6482 const ::mlir::DataLayout *dataLayout,
6485 assert(
false &&
"unimplemented, see TODO in the source.");
6489bool NVVMMemorySpaceAttr::isValidAddrSpaceCast(
6493 assert(
false &&
"unimplemented, see TODO in the source.");
6497bool NVVMMemorySpaceAttr::isValidPtrIntCast(
6502 assert(
false &&
"unimplemented, see TODO in the source.");
6511 int optLevel, StringRef triple, StringRef chip,
6512 StringRef features, DictionaryAttr flags,
6514 if (optLevel < 0 || optLevel > 3) {
6515 emitError() <<
"The optimization level must be a number between 0 and 3.";
6518 if (triple.empty()) {
6519 emitError() <<
"The target triple cannot be empty.";
6523 emitError() <<
"The target chip cannot be empty.";
6526 if (files && !llvm::all_of(files, [](::mlir::Attribute attr) {
6527 return mlir::isa_and_nonnull<StringAttr>(attr);
6529 emitError() <<
"All the elements in the `link` array must be strings.";
6535LogicalResult NVVMTargetAttr::verifyTarget(
Operation *gpuModule) {
6536 if (!getVerifyTarget())
6539 auto gpuModuleOp = llvm::dyn_cast<gpu::GPUModuleOp>(gpuModule);
6542 "NVVM target attribute must be attached to a GPU module");
6545 const unsigned targetFullSmVersion =
6549 "Minimum NVVM target SM version is sm_20");
6553 ->
walk([&](Operation *op) {
6554 if (
auto reqOp = llvm::dyn_cast<NVVM::RequiresSMInterface>(op)) {
6555 const NVVMCheckSMVersion requirement =
6556 reqOp.getRequiredMinSMVersion();
6558 op->
emitOpError() <<
"is not supported on " << getChip();
6570#define GET_OP_CLASSES
6571#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
6573#define GET_ATTRDEF_CLASSES
6574#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
p<< " : "<< getMemRefType()<< ", "<< getType();}static LogicalResult verifyVectorMemoryOp(Operation *op, MemRefType memrefType, VectorType vectorType) { if(memrefType.getElementType() !=vectorType.getElementType()) return op-> emitOpError("requires memref and vector types of the same elemental type")
Given a list of lists of parsed operands, populates uniqueOperands with unique operands.
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)
#define GET_F32x2_TO_F8X2_US_ID(rnd, has_satf)
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 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)
#define GET_F32x2_TO_F6x2_ID(type, has_relu)
static llvm::Value * getAsPackedI32(llvm::Value *arg, llvm::IRBuilderBase &builder)
#define GET_F16x2_TO_F8X2_ID(type, has_relu)
static LogicalResult verifyMBarrierArriveLikeOp(Operation *op, Value addr, NVVM::MemScopeKind scope, Value retVal=nullptr)
static llvm::Value * castPtrToAddrSpace(llvm::IRBuilderBase &builder, llvm::Value *ptr, NVVMMemorySpace targetAS)
static LogicalResult isAllowedWGMMADataType(NVVM::WGMMATypes typeD, NVVM::WGMMATypes typeA, NVVM::WGMMATypes typeB)
static llvm::Intrinsic::ID getBarrierReductionIntrinsic(bool aligned, NVVM::BarrierReduction kind)
Maps the (aligned, kind) pair to the @llvm.nvvm.barrier.cta.red.
static void inferAndSetMultiplicandTypes(MLIRContext *ctx, NamedAttrList &attrs, const SmallVectorImpl< Type > &operandTypes)
static LogicalResult parseMmaOperand(OpAsmParser &parser, StringRef operandName, SmallVectorImpl< OpAsmParser::UnresolvedOperand > ®s)
static std::pair< mlir::Type, unsigned > inferMMATypeFromMNK(NVVM::MMATypes type, NVVM::MMAFrag frag, int m, int n, int k, MLIRContext *context)
static bool isInt8PtxType(MMATypes type)
#define TCGEN05LDRED(SHAPE, NUM, TYPE)
static bool isInt4PtxType(MMATypes type)
static bool isIntegerPtxType(MMATypes type)
#define GET_F32x2_TO_F8X2_S_ID(type, has_relu)
static MMATypes inferPtxTypeFromResult(OpTy op)
static LogicalResult verifyConstantRangeAttr(Operation *op, std::optional< LLVM::ConstantRangeAttr > rangeAttr)
Verify the range attribute satisfies LLVM ConstantRange constructor requirements for NVVM SpecialRang...
static LogicalResult parseMmaTypeSignature(OpAsmParser &parser, SmallVectorImpl< Type > &operandTypes)
static FailureOr< int > getAllowedSizeK(NVVM::WGMMATypes typeA)
static bool isPtrInSharedClusterSpace(mlir::Value ptr)
#define GET_CP_ASYNC_ID(mod, size, has_cpsize)
static unsigned isValidVectorLength(NVVM::Tcgen05LdStShape shape, unsigned vecLen)
#define GET_TCGEN05_COMMIT_ID(cta_group, is_shared, has_mc)
static LogicalResult verifyConvertF32x2ToFP16x2Op(Twine dstType, FPRoundingMode rnd, bool hasRandomBits, Operation *op)
static void nvvmInferResultRanges(Operation *op, Value result, ArrayRef<::mlir::ConstantIntRanges > argRanges, SetIntRangeFn setResultRanges)
Infer the result ranges for the NVVM SpecialRangeableRegisterOp that might have ConstantRangeAttr.
static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims, bool isIm2Col, size_t numIm2ColOffsets, Location loc)
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.
virtual SMLoc getCurrentLocation()=0
Get the location of the next token and store it into the argument.
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)
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.
Attribute set(StringAttr name, Attribute value)
If the an attribute exists with the specified name, change it to the new value.
NamedAttribute represents a combination of a name and an Attribute value.
StringAttr getName() const
Return the name of the attribute.
Attribute getValue() const
Return the value of the attribute.
The OpAsmParser has methods for interacting with the asm parser: parsing things from it,...
ParseResult resolveOperands(Operands &&operands, Type type, SmallVectorImpl< Value > &result)
Resolve a list of operands to SSA values, emitting an error on failure, or appending the results to t...
virtual ParseResult parseOperandList(SmallVectorImpl< UnresolvedOperand > &result, Delimiter delimiter=Delimiter::None, bool allowResultNumber=true, int requiredOperandCount=-1)=0
Parse zero or more SSA comma-separated operand references with a specified surrounding delimiter,...
This is a pure-virtual base class that exposes the asmprinter hooks necessary to implement a custom p...
void printOperands(const ContainerType &container)
Print a comma separated list of operands.
virtual void printOptionalAttrDict(ArrayRef< NamedAttribute > attrs, ArrayRef< StringRef > elidedAttrs={})=0
If the specified operation has attributes, print out an attribute dictionary with their values.
This class helps build Operations.
This provides public APIs that all operations should have.
Operation is the basic unit of execution within MLIR.
AttrClass getAttrOfType(StringAttr name)
bool hasAttr(StringAttr name)
Return true if the operation has an attribute with the provided name, false otherwise.
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...
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()
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.
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 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.