MLIR 24.0.0git
WmmaOpsToNvvm.cpp
Go to the documentation of this file.
1//===------ WmmaOpsToNVVM.cpp - WMMA LD/ST/Compute to NVVM lowering -------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file contains definitions of patterns to lower GPU Subgroup MMA ops to
10// NVVM Dialect.
11//
12//===----------------------------------------------------------------------===//
13
20#include "mlir/IR/Types.h"
21
22using namespace mlir;
23
24namespace {
25
26/// Checks if all the operands of the op being lowered are of LLVM Types. The
27/// types are expected to be converted by the `LLVMTypeConverter` before the op
28/// is actually lowered. If the type of an operands is not already converted it
29/// hints a missing typeConversion and failure is returned in that case.
30static LogicalResult areAllLLVMTypes(Operation *op, ValueRange operands,
31 ConversionPatternRewriter &rewriter) {
32 if (!llvm::all_of(operands, [](Value value) {
33 return LLVM::isCompatibleType(value.getType());
34 })) {
35 return rewriter.notifyMatchFailure(
36 op, "cannot convert if operands aren't of LLVM type.");
37 }
38
39 return success();
40}
41
42/// Error string to emit when an unimplemented WMMA variant is encountered.
43static constexpr StringRef kInvalidCaseStr = "Unsupported WMMA variant.";
44
45static NVVM::MMAFrag convertOperand(StringRef operandName) {
46 if (operandName == "AOp")
47 return NVVM::MMAFrag::a;
48 if (operandName == "BOp")
49 return NVVM::MMAFrag::b;
50 if (operandName == "COp")
51 return NVVM::MMAFrag::c;
52 llvm_unreachable("Unknown operand name");
53}
54
55static NVVM::MMATypes getElementType(gpu::MMAMatrixType type) {
56 if (type.getElementType().isF16())
57 return NVVM::MMATypes::f16;
58 if (type.getElementType().isF32())
59 return type.getOperand() == "COp" ? NVVM::MMATypes::f32
60 : NVVM::MMATypes::tf32;
61 if (type.getElementType().isF64())
62 return NVVM::MMATypes::f64;
63 if (type.getElementType().isSignedInteger(8))
64 return NVVM::MMATypes::s8;
66 return NVVM::MMATypes::u8;
67 // Accumulator type is signless and implies signed.
68 if (type.getElementType().isInteger(32))
69 return NVVM::MMATypes::s32;
70 llvm_unreachable("Unsupported type");
71}
72
73/// This class implements the conversion of GPU MMA loadOp to wmma.load op
74/// in the NVVM dialect. The conversion not only emits the NVVM op but also
75/// emits code that is necessary to store the data in the destination memref
76/// after it has been loaded.
77struct WmmaLoadOpToNVVMLowering
78 : public ConvertOpToLLVMPattern<gpu::SubgroupMmaLoadMatrixOp> {
79 using ConvertOpToLLVMPattern<
80 gpu::SubgroupMmaLoadMatrixOp>::ConvertOpToLLVMPattern;
81
82 LogicalResult
83 matchAndRewrite(gpu::SubgroupMmaLoadMatrixOp subgroupMmaLoadMatrixOp,
84 OpAdaptor adaptor,
85 ConversionPatternRewriter &rewriter) const override {
86 Operation *op = subgroupMmaLoadMatrixOp.getOperation();
87 if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
88 return failure();
89
90 // Get the shape of the MMAMatrix type being returned. The shape will
91 // choose which intrinsic this op will be lowered to.
92 NVVM::MMALayout layout = subgroupMmaLoadMatrixOp.getTranspose()
93 ? NVVM::MMALayout::col
94 : NVVM::MMALayout::row;
95 gpu::MMAMatrixType retType =
96 cast<gpu::MMAMatrixType>(subgroupMmaLoadMatrixOp.getRes().getType());
97 ArrayRef<int64_t> retTypeShape = retType.getShape();
98 int64_t m = 0;
99 int64_t n = 0;
100 int64_t k = 0;
101 NVVM::MMATypes eltype = getElementType(retType);
102 // NVVM intrinsics require to give mxnxk dimensions, infer the missing
103 // dimension based on the valid intrinsics available.
104 if (retType.getOperand() == "AOp") {
105 m = retTypeShape[0];
106 k = retTypeShape[1];
107 n = NVVM::WMMALoadOp::inferNDimension(m, k, eltype);
108 } else if (retType.getOperand() == "BOp") {
109 k = retTypeShape[0];
110 n = retTypeShape[1];
111 m = NVVM::WMMALoadOp::inferMDimension(k, n, eltype);
112 } else if (retType.getOperand() == "COp") {
113 m = retTypeShape[0];
114 n = retTypeShape[1];
115 k = NVVM::WMMALoadOp::inferKDimension(m, n, eltype);
116 }
117 NVVM::MMAFrag frag = convertOperand(retType.getOperand());
118 // Check that there is an exisiting instruction for the combination we need.
119 if (NVVM::WMMALoadOp::getIntrinsicID(m, n, k, layout, eltype, frag) == 0)
120 return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
121
122 Type resType = convertMMAToLLVMType(retType);
123 Location loc = op->getLoc();
124
125 // Create nvvm.mma_load op according to the operand types.
126 Value dataPtr = getStridedElementPtr(
127 rewriter, loc,
128 cast<MemRefType>(subgroupMmaLoadMatrixOp.getSrcMemref().getType()),
129 adaptor.getSrcMemref(), adaptor.getIndices());
130
131 // The NVVM op takes the leading dimension as an `i32`. Reject a value that
132 // does not fit rather than silently truncating it.
133 int64_t leadDimension =
134 subgroupMmaLoadMatrixOp.getLeadDimension().getSExtValue();
135 if (!llvm::isInt<32>(leadDimension))
136 return rewriter.notifyMatchFailure(
137 op, "leading dimension does not fit into an i32");
138 Value leadingDim = LLVM::ConstantOp::create(
139 rewriter, loc, rewriter.getI32Type(),
140 rewriter.getI32IntegerAttr(static_cast<int32_t>(leadDimension)));
141 rewriter.replaceOpWithNewOp<NVVM::WMMALoadOp>(
142 op, resType, dataPtr, leadingDim, m, n, k, layout, eltype, frag);
143 return success();
144 }
145};
146
147/// This class implements the conversion of GPU MMA storeOp to wmma.store op
148/// in the NVVM dialect. The conversion not only emits the NVVM op but also
149/// emits code that is necessary to unpack the data in the source and
150/// convert the data in the format that is needed by the NVVM op.
151struct WmmaStoreOpToNVVMLowering
152 : public ConvertOpToLLVMPattern<gpu::SubgroupMmaStoreMatrixOp> {
153 using ConvertOpToLLVMPattern<
154 gpu::SubgroupMmaStoreMatrixOp>::ConvertOpToLLVMPattern;
155
156 LogicalResult
157 matchAndRewrite(gpu::SubgroupMmaStoreMatrixOp subgroupMmaStoreMatrixOp,
158 OpAdaptor adaptor,
159 ConversionPatternRewriter &rewriter) const override {
160 Operation *op = subgroupMmaStoreMatrixOp.getOperation();
161 if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
162 return failure();
163
164 Location loc = op->getLoc();
165
166 SmallVector<Value, 4> storeOpOperands;
167 // Get the shape of the MMAMatrix type being stored. The shape will
168 // choose which intrinsic this op will be lowered to.
169 gpu::MMAMatrixType srcType =
170 cast<gpu::MMAMatrixType>(subgroupMmaStoreMatrixOp.getSrc().getType());
171 ArrayRef<int64_t> srcTypeShape = srcType.getShape();
172 NVVM::MMALayout layout = subgroupMmaStoreMatrixOp.getTranspose()
173 ? NVVM::MMALayout::col
174 : NVVM::MMALayout::row;
175 NVVM::MMATypes eltype = getElementType(srcType);
176 int64_t m = srcTypeShape[0];
177 int64_t n = srcTypeShape[1];
178 int64_t k = NVVM::WMMAStoreOp::inferKDimension(m, n, eltype);
179 if (NVVM::WMMAStoreOp::getIntrinsicID(m, n, k, layout, eltype) == 0)
180 return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
181
182 auto matrixType = cast<LLVM::LLVMStructType>(adaptor.getSrc().getType());
183 for (unsigned i = 0, e = matrixType.getBody().size(); i < e; ++i) {
184 Value toUse =
185 LLVM::ExtractValueOp::create(rewriter, loc, adaptor.getSrc(), i);
186 storeOpOperands.push_back(toUse);
187 }
188
189 Value dataPtr = getStridedElementPtr(
190 rewriter, loc,
191 cast<MemRefType>(subgroupMmaStoreMatrixOp.getDstMemref().getType()),
192 adaptor.getDstMemref(), adaptor.getIndices());
193 // The NVVM op takes the leading dimension as an `i32`. Reject a value that
194 // does not fit rather than silently truncating it.
195 int64_t leadDimension =
196 subgroupMmaStoreMatrixOp.getLeadDimension().getSExtValue();
197 if (!llvm::isInt<32>(leadDimension))
198 return rewriter.notifyMatchFailure(
199 op, "leading dimension does not fit into an i32");
200 Value leadingDim = LLVM::ConstantOp::create(
201 rewriter, loc, rewriter.getI32Type(),
202 rewriter.getI32IntegerAttr(static_cast<int32_t>(leadDimension)));
203 rewriter.replaceOpWithNewOp<NVVM::WMMAStoreOp>(
204 op, dataPtr, m, n, k, layout, eltype, storeOpOperands, leadingDim);
205 return success();
206 }
207};
208
209/// This class implements the conversion of GPU MMA computeOp to wmma.mma op
210/// in the NVVM dialect.
211struct WmmaMmaOpToNVVMLowering
212 : public ConvertOpToLLVMPattern<gpu::SubgroupMmaComputeOp> {
213 using ConvertOpToLLVMPattern<
214 gpu::SubgroupMmaComputeOp>::ConvertOpToLLVMPattern;
215
216 LogicalResult
217 matchAndRewrite(gpu::SubgroupMmaComputeOp subgroupMmaComputeOp,
218 OpAdaptor adaptor,
219 ConversionPatternRewriter &rewriter) const override {
220 Operation *op = subgroupMmaComputeOp.getOperation();
221 if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
222 return failure();
223
224 Location loc = op->getLoc();
225
226 // The wmma.mma intrinsic in llvm requires the operands as individual
227 // values. So individual elements from the memrefs need to be extracted and
228 // then passed on to the intrinsic call. Emit llvm ops to extract individual
229 // values form lowered memrefs.
230 SmallVector<Value> unpackedOps;
231 auto unpackOp = [&](Value operand) {
232 // f64 a and b fragments are not structs but scalars.
233 if (!isa<LLVM::LLVMStructType>(operand.getType())) {
234 unpackedOps.push_back(operand);
235 return;
236 }
237 // every other type is lowered to an LLVM struct, extract the values.
238 auto structType = cast<LLVM::LLVMStructType>(operand.getType());
239 for (size_t i = 0, e = structType.getBody().size(); i < e; ++i) {
240 Value toUse = LLVM::ExtractValueOp::create(rewriter, loc, operand, i);
241 unpackedOps.push_back(toUse);
242 }
243 };
244
245 // Get the shapes of the MMAMatrix type being used. The shapes will
246 // choose which intrinsic this op will be lowered to.
247 gpu::MMAMatrixType aType =
248 cast<gpu::MMAMatrixType>(subgroupMmaComputeOp.getOpA().getType());
249 ArrayRef<int64_t> aTypeShape = aType.getShape();
250 gpu::MMAMatrixType cType =
251 cast<gpu::MMAMatrixType>(subgroupMmaComputeOp.getOpC().getType());
252 ArrayRef<int64_t> cTypeShape = cType.getShape();
253 int64_t m = cTypeShape[0];
254 int64_t n = cTypeShape[1];
255 int64_t k = aTypeShape[1];
256 NVVM::MMALayout aLayout = subgroupMmaComputeOp.getATranspose()
257 ? NVVM::MMALayout::col
258 : NVVM::MMALayout::row;
259 NVVM::MMALayout bLayout = subgroupMmaComputeOp.getBTranspose()
260 ? NVVM::MMALayout::col
261 : NVVM::MMALayout::row;
262 NVVM::MMATypes sourceType = getElementType(aType);
263 NVVM::MMATypes destType = getElementType(cType);
264 if (NVVM::WMMAMmaOp::getIntrinsicID(m, n, k, aLayout, bLayout, sourceType,
265 destType) == 0)
266 return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
267
268 NVVM::MMATypes bElementType = getElementType(
269 cast<gpu::MMAMatrixType>(subgroupMmaComputeOp.getOpB().getType()));
270 if (bElementType != sourceType)
271 return rewriter.notifyMatchFailure(
272 op, "WMMA compute op input matrix element types must match.");
273
274 unpackOp(adaptor.getOpA());
275 unpackOp(adaptor.getOpB());
276 unpackOp(adaptor.getOpC());
277
278 rewriter.replaceOpWithNewOp<NVVM::WMMAMmaOp>(
279 op, adaptor.getOpC().getType(), m, n, k, aLayout, bLayout, sourceType,
280 destType, unpackedOps);
281 return success();
282 }
283};
284
285/// Convert GPU MMA ConstantMatrixOp to a chain of InsertValueOp.
286struct WmmaConstantOpToNVVMLowering
287 : public ConvertOpToLLVMPattern<gpu::SubgroupMmaConstantMatrixOp> {
288 using ConvertOpToLLVMPattern<
289 gpu::SubgroupMmaConstantMatrixOp>::ConvertOpToLLVMPattern;
290
291 LogicalResult
292 matchAndRewrite(gpu::SubgroupMmaConstantMatrixOp subgroupMmaConstantOp,
293 OpAdaptor adaptor,
294 ConversionPatternRewriter &rewriter) const override {
295 if (failed(areAllLLVMTypes(subgroupMmaConstantOp.getOperation(),
296 adaptor.getOperands(), rewriter)))
297 return failure();
298 Location loc = subgroupMmaConstantOp.getLoc();
299 Value cst = adaptor.getOperands()[0];
300 Type type = convertMMAToLLVMType(
301 cast<gpu::MMAMatrixType>(subgroupMmaConstantOp.getType()));
302 // If the element is not a struct, it means it's a scalar f64.
303 auto structType = dyn_cast<LLVM::LLVMStructType>(type);
304 if (!structType) {
305 rewriter.replaceOp(subgroupMmaConstantOp, cst);
306 return success();
307 }
308 // If the element type is a vector create a vector from the operand.
309 if (auto vecType = dyn_cast<VectorType>(structType.getBody()[0])) {
310 Value vecCst = LLVM::PoisonOp::create(rewriter, loc, vecType);
311 for (int64_t vecEl = 0; vecEl < vecType.getNumElements(); vecEl++) {
312 Value idx = LLVM::ConstantOp::create(rewriter, loc,
313 rewriter.getI32Type(), vecEl);
314 vecCst = LLVM::InsertElementOp::create(rewriter, loc, vecType, vecCst,
315 cst, idx);
316 }
317 cst = vecCst;
318 }
319 Value matrixStruct = LLVM::PoisonOp::create(rewriter, loc, structType);
320 for (size_t i : llvm::seq(size_t(0), structType.getBody().size())) {
321 matrixStruct =
322 LLVM::InsertValueOp::create(rewriter, loc, matrixStruct, cst, i);
323 }
324 rewriter.replaceOp(subgroupMmaConstantOp, matrixStruct);
325 return success();
326 }
327};
328
329static Value createMinMaxF(OpBuilder &builder, Location loc, Value lhs,
330 Value rhs, bool isMin) {
331 auto floatType = cast<FloatType>(getElementTypeOrSelf(lhs.getType()));
332 Type i1Type = builder.getI1Type();
333 if (auto vecType = dyn_cast<VectorType>(lhs.getType()))
334 i1Type = VectorType::get(vecType.getShape(), i1Type);
335 Value cmp = LLVM::FCmpOp::create(
336 builder, loc, i1Type,
337 isMin ? LLVM::FCmpPredicate::olt : LLVM::FCmpPredicate::ogt, lhs, rhs);
338 Value sel = LLVM::SelectOp::create(builder, loc, cmp, lhs, rhs);
339 Value isNan = LLVM::FCmpOp::create(builder, loc, i1Type,
340 LLVM::FCmpPredicate::uno, lhs, rhs);
341 Value nan = LLVM::ConstantOp::create(
342 builder, loc, lhs.getType(),
343 builder.getFloatAttr(floatType,
344 APFloat::getQNaN(floatType.getFloatSemantics())));
345 return LLVM::SelectOp::create(builder, loc, isNan, nan, sel);
346}
347
348static Value createScalarOp(OpBuilder &builder, Location loc,
349 gpu::MMAElementwiseOp op,
350 ArrayRef<Value> operands) {
351 switch (op) {
352 case gpu::MMAElementwiseOp::ADDF:
353 return LLVM::FAddOp::create(builder, loc, operands[0].getType(), operands);
354 case gpu::MMAElementwiseOp::MULF:
355 return LLVM::FMulOp::create(builder, loc, operands[0].getType(), operands);
356 case gpu::MMAElementwiseOp::DIVF:
357 return LLVM::FDivOp::create(builder, loc, operands[0].getType(), operands);
358 case gpu::MMAElementwiseOp::MAXF:
359 return createMinMaxF(builder, loc, operands[0], operands[1],
360 /*isMin=*/false);
361 case gpu::MMAElementwiseOp::MINF:
362 return createMinMaxF(builder, loc, operands[0], operands[1],
363 /*isMin=*/true);
364 default:
365 llvm_unreachable("unknown op");
366 }
367}
368
369/// Convert GPU MMA elementwise ops to extract + op + insert.
370struct WmmaElementwiseOpToNVVMLowering
371 : public ConvertOpToLLVMPattern<gpu::SubgroupMmaElementwiseOp> {
372 using ConvertOpToLLVMPattern<
373 gpu::SubgroupMmaElementwiseOp>::ConvertOpToLLVMPattern;
374
375 LogicalResult
376 matchAndRewrite(gpu::SubgroupMmaElementwiseOp subgroupMmaElementwiseOp,
377 OpAdaptor adaptor,
378 ConversionPatternRewriter &rewriter) const override {
379 if (failed(areAllLLVMTypes(subgroupMmaElementwiseOp.getOperation(),
380 adaptor.getOperands(), rewriter)))
381 return failure();
382 Location loc = subgroupMmaElementwiseOp.getLoc();
383 size_t numOperands = adaptor.getOperands().size();
384 Type destType = convertMMAToLLVMType(
385 cast<gpu::MMAMatrixType>(subgroupMmaElementwiseOp.getType()));
386
387 // If the element is not a struct, it means it's a scalar f64.
388 LLVM::LLVMStructType structDestTy =
389 dyn_cast<LLVM::LLVMStructType>(destType);
390 if (!structDestTy) {
391 SmallVector<Value> operands;
392 for (auto operand : adaptor.getOperands()) {
393 operands.push_back(operand);
394 }
395 Value element = createScalarOp(
396 rewriter, loc, subgroupMmaElementwiseOp.getOpType(), operands);
397 rewriter.replaceOp(subgroupMmaElementwiseOp, element);
398 return success();
399 }
400 Value matrixStruct = LLVM::PoisonOp::create(rewriter, loc, structDestTy);
401 for (size_t i = 0, e = structDestTy.getBody().size(); i < e; ++i) {
402 SmallVector<Value> extractedOperands;
403 for (size_t opIdx = 0; opIdx < numOperands; opIdx++) {
404 extractedOperands.push_back(LLVM::ExtractValueOp::create(
405 rewriter, loc, adaptor.getOperands()[opIdx], i));
406 }
407 Value element =
408 createScalarOp(rewriter, loc, subgroupMmaElementwiseOp.getOpType(),
409 extractedOperands);
410 matrixStruct =
411 LLVM::InsertValueOp::create(rewriter, loc, matrixStruct, element, i);
412 }
413 rewriter.replaceOp(subgroupMmaElementwiseOp, matrixStruct);
414 return success();
415 }
416};
417
418} // namespace
419
420/// Return the LLVMStructureType corresponding to the MMAMatrixType `type`.
422 NVVM::MMAFrag frag = convertOperand(type.getOperand());
423 NVVM::MMATypes eltType = getElementType(type);
424 auto nRow = type.getShape()[0];
425 auto nCol = type.getShape()[1];
426 std::pair<Type, unsigned> typeInfo =
427 NVVM::inferMMAType(eltType, frag, nRow, nCol, type.getContext());
428 // Special handling for f64 a and b fragments
429 Type f64Ty = Float64Type::get(type.getContext());
430 if (typeInfo.first == f64Ty && typeInfo.second == 1) {
431 return f64Ty;
432 }
433 return LLVM::LLVMStructType::getLiteral(
434 type.getContext(), SmallVector<Type, 8>(typeInfo.second, typeInfo.first));
435}
436
438 const LLVMTypeConverter &converter, RewritePatternSet &patterns,
439 PatternBenefit benefit) {
440 patterns.add<WmmaLoadOpToNVVMLowering, WmmaMmaOpToNVVMLowering,
441 WmmaStoreOpToNVVMLowering, WmmaConstantOpToNVVMLowering,
442 WmmaElementwiseOpToNVVMLowering>(converter, benefit);
443}
return success()
static LogicalResult areAllLLVMTypes(Operation *op, ValueRange operands, ConversionPatternRewriter &rewriter)
lhs
static Type getElementType(Type type, ArrayRef< int32_t > indices, function_ref< InFlightDiagnostic(StringRef)> emitErrorFn)
Walks the given type hierarchy with the given indices, potentially down to component granularity,...
Definition SPIRVOps.cpp:229
FloatAttr getFloatAttr(Type type, double value)
Definition Builders.cpp:263
IntegerType getI1Type()
Definition Builders.cpp:61
Utility class for operation conversions targeting the LLVM dialect that match exactly one source oper...
Definition Pattern.h:233
Conversion from types to the LLVM IR dialect.
This class defines the main interface for locations in MLIR and acts as a non-nullable wrapper around...
Definition Location.h:76
This class helps build Operations.
Definition Builders.h:210
Operation is the basic unit of execution within MLIR.
Definition Operation.h:87
Location getLoc()
The source location the operation was defined or derived from.
Definition Operation.h:240
This class represents the benefit of a pattern match in a unitless scheme that ranges from 0 (very li...
RewritePatternSet & add(ConstructorArg &&arg, ConstructorArgs &&...args)
Add an instance of each of the pattern types 'Ts' to the pattern list with the given arguments.
Instances of the Type class are uniqued, have an immutable identifier and an optional mutable compone...
Definition Types.h:74
bool isF64() const
Definition Types.cpp:41
bool isSignedInteger() const
Return true if this is a signed integer type (with the specified width).
Definition Types.cpp:78
bool isF32() const
Definition Types.cpp:40
bool isUnsignedInteger() const
Return true if this is an unsigned integer type (with the specified width).
Definition Types.cpp:90
bool isInteger() const
Return true if this is an integer type (with the specified width).
Definition Types.cpp:58
bool isF16() const
Definition Types.cpp:38
This class provides an abstraction over the different types of ranges over Values.
Definition ValueRange.h:389
This class represents an instance of an SSA value in the MLIR system, representing a computable value...
Definition Value.h:96
Type getType() const
Return the type of this value.
Definition Value.h:105
MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations.
Definition GPUDialect.h:143
ArrayRef< int64_t > getShape() const
Get shape of the matrix.
Type getElementType() const
Get elementType of a single element.
StringRef getOperand() const
The general form of operation this type supports is given by the equation C += A*B.
Value getStridedElementPtr(OpBuilder &builder, Location loc, const LLVMTypeConverter &converter, MemRefType type, Value memRefDesc, ValueRange indices, LLVM::GEPNoWrapFlags noWrapFlags=LLVM::GEPNoWrapFlags::none)
Performs the index computation to get to the element at indices of the memory pointed to by memRefDes...
Definition Pattern.cpp:608
bool isCompatibleType(Type type)
Returns true if the given type is compatible with the LLVM dialect.
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.
detail::InFlightRemark failed(Location loc, RemarkOpts opts)
Report an optimization remark that failed.
Definition Remarks.h:717
Include the generated interface declarations.
Type getType(OpFoldResult ofr)
Returns the int type of the integer in ofr.
Definition Utils.cpp:307
Type convertMMAToLLVMType(gpu::MMAMatrixType type)
Return the LLVMStructureType corresponding to the MMAMatrixType type.
Type getElementTypeOrSelf(Type type)
Return the element type or return the type itself.
void populateGpuWMMAToNVVMConversionPatterns(const LLVMTypeConverter &converter, RewritePatternSet &patterns, PatternBenefit benefit=1)
Collect a set of patterns to convert WMMA ops from GPU dialect to NVVM.