MLIR  18.0.0git
AMDGPUToROCDL.cpp
Go to the documentation of this file.
1 //===- AMDGPUToROCDL.cpp - AMDGPU to ROCDL dialect conversion -------===//
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 
10 
17 #include "mlir/IR/BuiltinTypes.h"
18 #include "mlir/IR/TypeUtilities.h"
19 #include "mlir/Pass/Pass.h"
20 
21 #include "llvm/ADT/STLExtras.h"
22 #include <optional>
23 
24 namespace mlir {
25 #define GEN_PASS_DEF_CONVERTAMDGPUTOROCDL
26 #include "mlir/Conversion/Passes.h.inc"
27 } // namespace mlir
28 
29 using namespace mlir;
30 using namespace mlir::amdgpu;
31 
33  Location loc, int32_t value) {
34  Type llvmI32 = rewriter.getI32Type();
35  return rewriter.create<LLVM::ConstantOp>(loc, llvmI32, value);
36 }
37 
39  bool value) {
40  Type llvmI1 = rewriter.getI1Type();
41  return rewriter.createOrFold<LLVM::ConstantOp>(loc, llvmI1, value);
42 }
43 
44 namespace {
45 /// Define lowering patterns for raw buffer ops
46 template <typename GpuOp, typename Intrinsic>
47 struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
48  RawBufferOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
49  : ConvertOpToLLVMPattern<GpuOp>(converter), chipset(chipset) {}
50 
51  Chipset chipset;
52  static constexpr uint32_t maxVectorOpWidth = 128;
53 
55  matchAndRewrite(GpuOp gpuOp, typename GpuOp::Adaptor adaptor,
56  ConversionPatternRewriter &rewriter) const override {
57  Location loc = gpuOp.getLoc();
58  Value memref = adaptor.getMemref();
59  Value unconvertedMemref = gpuOp.getMemref();
60  MemRefType memrefType = cast<MemRefType>(unconvertedMemref.getType());
61 
62  if (chipset.majorVersion < 9)
63  return gpuOp.emitOpError("raw buffer ops require GCN or higher");
64 
65  Value storeData = adaptor.getODSOperands(0)[0];
66  if (storeData == memref) // no write component to this op
67  storeData = Value();
68  Type wantedDataType;
69  if (storeData)
70  wantedDataType = storeData.getType();
71  else
72  wantedDataType = gpuOp.getODSResults(0)[0].getType();
73 
74  Value atomicCmpData = Value();
75  // Operand index 1 of a load is the indices, trying to read them can crash.
76  if (storeData) {
77  Value maybeCmpData = adaptor.getODSOperands(1)[0];
78  if (maybeCmpData != memref)
79  atomicCmpData = maybeCmpData;
80  }
81 
82  Type llvmWantedDataType = this->typeConverter->convertType(wantedDataType);
83 
84  Type i32 = rewriter.getI32Type();
85  Type llvmI32 = this->typeConverter->convertType(i32);
86  Type llvmI16 = this->typeConverter->convertType(rewriter.getI16Type());
87 
88  int64_t elementByteWidth = memrefType.getElementTypeBitWidth() / 8;
89  Value byteWidthConst = createI32Constant(rewriter, loc, elementByteWidth);
90 
91  // If we want to load a vector<NxT> with total size <= 32
92  // bits, use a scalar load and bitcast it. Similarly, if bitsize(T) < 32
93  // and the total load size is >= 32, use a vector load of N / (bitsize(T) /
94  // 32) x i32 and bitcast. Also, the CAS intrinsic requires integer operands,
95  // so bitcast any floats to integers. On top of all this, cast bfloat
96  // (vectors) to i16 since the backend doesn't currently support bfloat on
97  // these operations.
98  Type llvmBufferValType = llvmWantedDataType;
99  if (wantedDataType.isBF16())
100  llvmBufferValType = rewriter.getI16Type();
101  if (auto wantedVecType = dyn_cast<VectorType>(wantedDataType))
102  if (wantedVecType.getElementType().isBF16())
103  llvmBufferValType = wantedVecType.clone(rewriter.getI16Type());
104  if (atomicCmpData) {
105  if (isa<VectorType>(wantedDataType))
106  return gpuOp.emitOpError("vector compare-and-swap does not exist");
107  if (auto floatType = dyn_cast<FloatType>(wantedDataType))
108  llvmBufferValType = this->getTypeConverter()->convertType(
109  rewriter.getIntegerType(floatType.getWidth()));
110  }
111  if (auto dataVector = dyn_cast<VectorType>(wantedDataType)) {
112  uint32_t elemBits = dataVector.getElementTypeBitWidth();
113  uint32_t totalBits = elemBits * dataVector.getNumElements();
114  if (totalBits > maxVectorOpWidth)
115  return gpuOp.emitOpError(
116  "Total width of loads or stores must be no more than " +
117  Twine(maxVectorOpWidth) + " bits, but we call for " +
118  Twine(totalBits) +
119  " bits. This should've been caught in validation");
120  if (elemBits < 32) {
121  if (totalBits > 32) {
122  if (totalBits % 32 != 0)
123  return gpuOp.emitOpError("Load or store of more than 32-bits that "
124  "doesn't fit into words. Can't happen\n");
125  llvmBufferValType = this->typeConverter->convertType(
126  VectorType::get(totalBits / 32, i32));
127  } else {
128  llvmBufferValType = this->typeConverter->convertType(
129  rewriter.getIntegerType(totalBits));
130  }
131  }
132  }
133 
135  if (storeData) {
136  if (llvmBufferValType != llvmWantedDataType) {
137  Value castForStore =
138  rewriter.create<LLVM::BitcastOp>(loc, llvmBufferValType, storeData);
139  args.push_back(castForStore);
140  } else {
141  args.push_back(storeData);
142  }
143  }
144 
145  if (atomicCmpData) {
146  if (llvmBufferValType != llvmWantedDataType) {
147  Value castForCmp = rewriter.create<LLVM::BitcastOp>(
148  loc, llvmBufferValType, atomicCmpData);
149  args.push_back(castForCmp);
150  } else {
151  args.push_back(atomicCmpData);
152  }
153  }
154 
155  // Construct buffer descriptor from memref, attributes
156  int64_t offset = 0;
157  SmallVector<int64_t, 5> strides;
158  if (failed(getStridesAndOffset(memrefType, strides, offset)))
159  return gpuOp.emitOpError("Can't lower non-stride-offset memrefs");
160 
161  MemRefDescriptor memrefDescriptor(memref);
162 
163  Value ptr = memrefDescriptor.alignedPtr(rewriter, loc);
164  // The stride value is always 0 for raw buffers. This also disables
165  // swizling.
166  Value stride = rewriter.createOrFold<LLVM::ConstantOp>(
167  loc, llvmI16, rewriter.getI16IntegerAttr(0));
168  Value numRecords;
169  if (memrefType.hasStaticShape()) {
170  numRecords = createI32Constant(
171  rewriter, loc,
172  static_cast<int32_t>(memrefType.getNumElements() * elementByteWidth));
173  } else {
174  Value maxIndex;
175  for (uint32_t i = 0, e = memrefType.getRank(); i < e; ++i) {
176  Value size = memrefDescriptor.size(rewriter, loc, i);
177  Value stride = memrefDescriptor.stride(rewriter, loc, i);
178  stride = rewriter.create<LLVM::MulOp>(loc, stride, byteWidthConst);
179  Value maxThisDim = rewriter.create<LLVM::MulOp>(loc, size, stride);
180  maxIndex = maxIndex ? rewriter.create<LLVM::MaximumOp>(loc, maxIndex,
181  maxThisDim)
182  : maxThisDim;
183  }
184  numRecords = rewriter.create<LLVM::TruncOp>(loc, llvmI32, maxIndex);
185  }
186 
187  // Flag word:
188  // bits 0-11: dst sel, ignored by these intrinsics
189  // bits 12-14: data format (ignored, must be nonzero, 7=float)
190  // bits 15-18: data format (ignored, must be nonzero, 4=32bit)
191  // bit 19: In nested heap (0 here)
192  // bit 20: Behavior on unmap (0 means "return 0 / ignore")
193  // bits 21-22: Index stride for swizzles (N/A)
194  // bit 23: Add thread ID (0)
195  // bit 24: Reserved to 1 (RDNA) or 0 (CDNA)
196  // bits 25-26: Reserved (0)
197  // bit 27: Buffer is non-volatile (CDNA only)
198  // bits 28-29: Out of bounds select (0 = structured, 1 = check index, 2 =
199  // none, 3 = either swizzles or testing against offset field) RDNA only
200  // bits 30-31: Type (must be 0)
201  uint32_t flags = (7 << 12) | (4 << 15);
202  if (chipset.majorVersion >= 10) {
203  flags |= (1 << 24);
204  uint32_t oob = adaptor.getBoundsCheck() ? 3 : 2;
205  flags |= (oob << 28);
206  }
207  Value flagsConst = createI32Constant(rewriter, loc, flags);
208  Type rsrcType = LLVM::LLVMPointerType::get(rewriter.getContext(), 8);
209  Value resource = rewriter.createOrFold<ROCDL::MakeBufferRsrcOp>(
210  loc, rsrcType, ptr, stride, numRecords, flagsConst);
211  args.push_back(resource);
212 
213  // Indexing (voffset)
214  Value voffset = createI32Constant(rewriter, loc, 0);
215  for (auto pair : llvm::enumerate(adaptor.getIndices())) {
216  size_t i = pair.index();
217  Value index = pair.value();
218  Value strideOp;
219  if (ShapedType::isDynamic(strides[i])) {
220  strideOp = rewriter.create<LLVM::MulOp>(
221  loc, memrefDescriptor.stride(rewriter, loc, i), byteWidthConst);
222  } else {
223  strideOp =
224  createI32Constant(rewriter, loc, strides[i] * elementByteWidth);
225  }
226  index = rewriter.create<LLVM::MulOp>(loc, index, strideOp);
227  voffset = rewriter.create<LLVM::AddOp>(loc, voffset, index);
228  }
229  if (adaptor.getIndexOffset()) {
230  int32_t indexOffset = *gpuOp.getIndexOffset() * elementByteWidth;
231  Value extraOffsetConst = createI32Constant(rewriter, loc, indexOffset);
232  voffset =
233  voffset ? rewriter.create<LLVM::AddOp>(loc, voffset, extraOffsetConst)
234  : extraOffsetConst;
235  }
236  args.push_back(voffset);
237 
238  Value sgprOffset = adaptor.getSgprOffset();
239  if (!sgprOffset)
240  sgprOffset = createI32Constant(rewriter, loc, 0);
241  if (ShapedType::isDynamic(offset))
242  sgprOffset = rewriter.create<LLVM::AddOp>(
243  loc, memrefDescriptor.offset(rewriter, loc), sgprOffset);
244  else if (offset > 0)
245  sgprOffset = rewriter.create<LLVM::AddOp>(
246  loc, sgprOffset, createI32Constant(rewriter, loc, offset));
247  args.push_back(sgprOffset);
248 
249  // bit 0: GLC = 0 (atomics drop value, less coherency)
250  // bits 1-2: SLC, DLC = 0 (similarly)
251  // bit 3: swizzled (0 for raw)
252  args.push_back(createI32Constant(rewriter, loc, 0));
253 
254  llvm::SmallVector<Type, 1> resultTypes(gpuOp->getNumResults(),
255  llvmBufferValType);
256  Operation *lowered = rewriter.create<Intrinsic>(loc, resultTypes, args,
258  if (lowered->getNumResults() == 1) {
259  Value replacement = lowered->getResult(0);
260  if (llvmBufferValType != llvmWantedDataType) {
261  replacement = rewriter.create<LLVM::BitcastOp>(loc, llvmWantedDataType,
262  replacement);
263  }
264  rewriter.replaceOp(gpuOp, replacement);
265  } else {
266  rewriter.eraseOp(gpuOp);
267  }
268  return success();
269  }
270 };
271 
272 struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
274 
276  matchAndRewrite(LDSBarrierOp op, LDSBarrierOp::Adaptor adaptor,
277  ConversionPatternRewriter &rewriter) const override {
278  auto asmDialectAttr = LLVM::AsmDialectAttr::get(rewriter.getContext(),
279  LLVM::AsmDialect::AD_ATT);
280  const char *asmStr = "s_waitcnt lgkmcnt(0)\ns_barrier";
281  const char *constraints = "";
282  rewriter.replaceOpWithNewOp<LLVM::InlineAsmOp>(
283  op,
284  /*resultTypes=*/TypeRange(), /*operands=*/ValueRange(),
285  /*asm_string=*/asmStr, constraints, /*has_side_effects=*/true,
286  /*is_align_stack=*/false, /*asm_dialect=*/asmDialectAttr,
287  /*operand_attrs=*/ArrayAttr());
288  return success();
289  }
290 };
291 } // namespace
292 
293 /// If `input` is a vector of bytes, concatentate those bytes in little-endian
294 /// order to form a single integer of size 8 * [vector length]. This works
295 /// around a wart in the AMDGPU intrinsics where operations that logically take
296 /// vectors of bytes instead integers. Since we do not want to expose this
297 /// implementation detail to MLIR, we correct for it here.
298 ///
299 /// In addition, convert vectors of LLVM bfloats to vectors of i16, since AMDGPU
300 /// MFMA intrinsics pre-date the bfloat type.
302  Location loc, Value input) {
303  Type inputType = input.getType();
304  if (auto vectorType = dyn_cast<VectorType>(inputType)) {
305  if (vectorType.getElementType().isBF16())
306  return rewriter.create<LLVM::BitcastOp>(
307  loc, vectorType.clone(rewriter.getI16Type()), input);
308 
309  if (!vectorType.getElementType().isInteger(8))
310  return input;
311  int64_t numBytes = vectorType.getNumElements();
312  Type destType = rewriter.getIntegerType(numBytes * 8);
313  Value result = rewriter.create<LLVM::ConstantOp>(
314  loc, destType, rewriter.getIntegerAttr(destType, 0));
315  for (int64_t i = 0; i < numBytes; ++i) {
316  Value idxConst = createI32Constant(rewriter, loc, i);
317  Value element =
318  rewriter.create<LLVM::ExtractElementOp>(loc, input, idxConst);
319  Value extended = rewriter.create<LLVM::ZExtOp>(loc, destType, element);
320  Value shiftConst = rewriter.create<LLVM::ConstantOp>(
321  loc, destType, rewriter.getIntegerAttr(destType, i * 8));
322  Value shifted = rewriter.create<LLVM::ShlOp>(loc, extended, shiftConst);
323  result = rewriter.create<LLVM::OrOp>(loc, result, shifted);
324  }
325  return result;
326  }
327  return input;
328 }
329 
330 /// Push an input operand. If it is a float type, nothing to do. If it is
331 /// an integer type, then we need to also push its signdness (1 for signed, 0
332 /// for unsigned) and we need to pack the input 16xi8 vector into a 4xi32
333 /// vector. We also need to convert bfloat inputs to i16 to account for the lack
334 /// of bfloat support in the WMMA intrinsics themselves.
336  Location loc,
337  const TypeConverter *typeConverter,
338  bool isUnsigned, Value llvmInput,
339  SmallVector<Value, 4> &operands) {
340  Type inputType = llvmInput.getType();
341  auto vectorType = inputType.dyn_cast<VectorType>();
342  Type elemType = vectorType.getElementType();
343 
344  if (elemType.isBF16())
345  llvmInput = rewriter.create<LLVM::BitcastOp>(
346  loc, vectorType.clone(rewriter.getI16Type()), llvmInput);
347  if (!elemType.isInteger(8)) {
348  operands.push_back(llvmInput);
349  return;
350  }
351 
352  int64_t numBytes = vectorType.getNumElements();
353  Type i32 = rewriter.getI32Type();
354  VectorType vectorType32bits = VectorType::get(numBytes * 8 / 32, i32);
355  auto llvmVectorType32bits = typeConverter->convertType(vectorType32bits);
356 
357  Value result = rewriter.createOrFold<LLVM::BitcastOp>(
358  loc, llvmVectorType32bits, llvmInput);
359 
360  // if element type is 8-bit signed or unsigned, ignore the isUnsigned flag
361  bool localIsUnsigned = isUnsigned;
362  if (elemType.isUnsignedInteger(8)) {
363  localIsUnsigned = true;
364  } else if (elemType.isSignedInteger(8)) {
365  localIsUnsigned = false;
366  }
367  Value sign = createI1Constant(rewriter, loc, !localIsUnsigned);
368  operands.push_back(sign);
369  operands.push_back(result);
370 }
371 
372 /// Push the output operand. For many cases this is only pushing the output in
373 /// the operand list. But when we have f16 -> f16 or bf16 -> bf16 intrinsics,
374 /// since the same numbers of VGPRs is used, we need to decide if to store the
375 /// result in the upper 16 bits of the VGPRs or in the lower part. To store the
376 /// result in the lower 16 bits, set subwordOffset to 1, otherwise result will
377 /// be stored it in the upper part
379  Location loc,
380  const TypeConverter *typeConverter,
381  Value output, int32_t subwordOffset,
382  bool clamp, SmallVector<Value, 4> &operands) {
383  Type inputType = output.getType();
384  auto vectorType = inputType.dyn_cast<VectorType>();
385  Type elemType = vectorType.getElementType();
386  if (elemType.isBF16())
387  output = rewriter.create<LLVM::BitcastOp>(
388  loc, vectorType.clone(rewriter.getI16Type()), output);
389  operands.push_back(output);
390  if (elemType.isF16() || elemType.isBF16() || elemType.isInteger(16)) {
391  operands.push_back(createI1Constant(rewriter, loc, subwordOffset));
392  } else if (elemType.isInteger(32)) {
393  operands.push_back(createI1Constant(rewriter, loc, clamp));
394  }
395 }
396 
397 /// Return the `rocdl` intrinsic corresponding to a MFMA operation `mfma`
398 /// if one exists. This includes checking to ensure the intrinsic is supported
399 /// on the architecture you are compiling for.
400 static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
401  Chipset chipset) {
402  uint32_t m = mfma.getM(), n = mfma.getN(), k = mfma.getK(),
403  b = mfma.getBlocks();
404  Type sourceElem = mfma.getSourceA().getType();
405  if (auto sourceType = dyn_cast<VectorType>(sourceElem))
406  sourceElem = sourceType.getElementType();
407  Type destElem = mfma.getDestC().getType();
408  if (auto destType = dyn_cast<VectorType>(destElem))
409  destElem = destType.getElementType();
410 
411  if (sourceElem.isF32() && destElem.isF32()) {
412  if (mfma.getReducePrecision() && chipset.minorVersion >= 0x40) {
413  if (m == 32 && n == 32 && k == 4 && b == 1)
414  return ROCDL::mfma_f32_32x32x4_xf32::getOperationName();
415  if (m == 16 && n == 16 && k == 8 && b == 1)
416  return ROCDL::mfma_f32_16x16x8_xf32::getOperationName();
417  }
418  if (m == 32 && n == 32 && k == 1 && b == 2)
419  return ROCDL::mfma_f32_32x32x1f32::getOperationName();
420  if (m == 16 && n == 16 && k == 1 && b == 4)
421  return ROCDL::mfma_f32_16x16x1f32::getOperationName();
422  if (m == 4 && n == 4 && k == 1 && b == 16)
423  return ROCDL::mfma_f32_4x4x1f32::getOperationName();
424  if (m == 32 && n == 32 && k == 2 && b == 1)
425  return ROCDL::mfma_f32_32x32x2f32::getOperationName();
426  if (m == 16 && n == 16 && k == 4 && b == 1)
427  return ROCDL::mfma_f32_16x16x4f32::getOperationName();
428  }
429 
430  if (sourceElem.isF16() && destElem.isF32()) {
431  if (m == 32 && n == 32 && k == 4 && b == 2)
432  return ROCDL::mfma_f32_32x32x4f16::getOperationName();
433  if (m == 16 && n == 16 && k == 4 && b == 4)
434  return ROCDL::mfma_f32_16x16x4f16::getOperationName();
435  if (m == 4 && n == 4 && k == 4 && b == 16)
436  return ROCDL::mfma_f32_4x4x4f16::getOperationName();
437  if (m == 32 && n == 32 && k == 8 && b == 1)
438  return ROCDL::mfma_f32_32x32x8f16::getOperationName();
439  if (m == 16 && n == 16 && k == 16 && b == 1)
440  return ROCDL::mfma_f32_16x16x16f16::getOperationName();
441  }
442 
443  if (sourceElem.isBF16() && destElem.isF32() && chipset.minorVersion >= 0x0a) {
444  if (m == 32 && n == 32 && k == 4 && b == 2)
445  return ROCDL::mfma_f32_32x32x4bf16_1k::getOperationName();
446  if (m == 16 && n == 16 && k == 4 && b == 4)
447  return ROCDL::mfma_f32_16x16x4bf16_1k::getOperationName();
448  if (m == 4 && n == 4 && k == 4 && b == 16)
449  return ROCDL::mfma_f32_4x4x4bf16_1k::getOperationName();
450  if (m == 32 && n == 32 && k == 8 && b == 1)
451  return ROCDL::mfma_f32_32x32x8bf16_1k::getOperationName();
452  if (m == 16 && n == 16 && k == 16 && b == 1)
453  return ROCDL::mfma_f32_16x16x16bf16_1k::getOperationName();
454  }
455 
456  if (sourceElem.isBF16() && destElem.isF32()) {
457  if (m == 32 && n == 32 && k == 2 && b == 2)
458  return ROCDL::mfma_f32_32x32x2bf16::getOperationName();
459  if (m == 16 && n == 16 && k == 2 && b == 4)
460  return ROCDL::mfma_f32_16x16x2bf16::getOperationName();
461  if (m == 4 && n == 4 && k == 2 && b == 16)
462  return ROCDL::mfma_f32_4x4x2bf16::getOperationName();
463  if (m == 32 && n == 32 && k == 4 && b == 1)
464  return ROCDL::mfma_f32_32x32x4bf16::getOperationName();
465  if (m == 16 && n == 16 && k == 8 && b == 1)
466  return ROCDL::mfma_f32_16x16x8bf16::getOperationName();
467  }
468 
469  if (isa<IntegerType>(sourceElem) && destElem.isInteger(32)) {
470  if (m == 32 && n == 32 && k == 4 && b == 2)
471  return ROCDL::mfma_i32_32x32x4i8::getOperationName();
472  if (m == 16 && n == 16 && k == 4 && b == 4)
473  return ROCDL::mfma_i32_16x16x4i8::getOperationName();
474  if (m == 4 && n == 4 && k == 4 && b == 16)
475  return ROCDL::mfma_i32_4x4x4i8::getOperationName();
476  if (m == 32 && n == 32 && k == 8 && b == 1)
477  return ROCDL::mfma_i32_32x32x8i8::getOperationName();
478  if (m == 16 && n == 16 && k == 16 && b == 1)
479  return ROCDL::mfma_i32_16x16x16i8::getOperationName();
480  if (m == 32 && n == 32 && k == 16 && b == 1 && chipset.minorVersion >= 0x40)
481  return ROCDL::mfma_i32_32x32x16_i8::getOperationName();
482  if (m == 16 && n == 16 && k == 32 && b == 1 && chipset.minorVersion >= 0x40)
483  return ROCDL::mfma_i32_16x16x32_i8::getOperationName();
484  }
485 
486  if (sourceElem.isF64() && destElem.isF64() && chipset.minorVersion >= 0x0a) {
487  if (m == 16 && n == 16 && k == 4 && b == 1)
488  return ROCDL::mfma_f64_16x16x4f64::getOperationName();
489  if (m == 4 && n == 4 && k == 4 && b == 4)
490  return ROCDL::mfma_f64_4x4x4f64::getOperationName();
491  }
492 
493  if (sourceElem.isFloat8E5M2FNUZ() && destElem.isF32() &&
494  chipset.minorVersion >= 0x40) {
495  // Known to be correct because there are no scalar f8 instructions and
496  // because a length mismatch will have been caught by the verifier.
497  Type sourceBElem =
498  cast<VectorType>(mfma.getSourceB().getType()).getElementType();
499  if (m == 16 && n == 16 && k == 32 && b == 1) {
500  if (sourceBElem.isFloat8E5M2FNUZ())
501  return ROCDL::mfma_f32_16x16x32_bf8_bf8::getOperationName();
502  if (sourceBElem.isFloat8E4M3FNUZ())
503  return ROCDL::mfma_f32_16x16x32_bf8_fp8::getOperationName();
504  }
505  if (m == 32 && n == 32 && k == 16 && b == 1) {
506  if (sourceBElem.isFloat8E5M2FNUZ())
507  return ROCDL::mfma_f32_32x32x16_bf8_bf8::getOperationName();
508  if (sourceBElem.isFloat8E4M3FNUZ())
509  return ROCDL::mfma_f32_32x32x16_bf8_fp8::getOperationName();
510  }
511  }
512 
513  if (sourceElem.isFloat8E4M3FNUZ() && destElem.isF32() &&
514  chipset.minorVersion >= 0x40) {
515  Type sourceBElem =
516  cast<VectorType>(mfma.getSourceB().getType()).getElementType();
517  if (m == 16 && n == 16 && k == 32 && b == 1) {
518  if (sourceBElem.isFloat8E5M2FNUZ())
519  return ROCDL::mfma_f32_16x16x32_fp8_bf8::getOperationName();
520  if (sourceBElem.isFloat8E4M3FNUZ())
521  return ROCDL::mfma_f32_16x16x32_fp8_fp8::getOperationName();
522  }
523  if (m == 32 && n == 32 && k == 16 && b == 1) {
524  if (sourceBElem.isFloat8E5M2FNUZ())
525  return ROCDL::mfma_f32_32x32x16_fp8_bf8::getOperationName();
526  if (sourceBElem.isFloat8E4M3FNUZ())
527  return ROCDL::mfma_f32_32x32x16_fp8_fp8::getOperationName();
528  }
529  }
530 
531  return std::nullopt;
532 }
533 
534 /// Return the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
535 /// if one exists. This includes checking to ensure the intrinsic is supported
536 /// on the architecture you are compiling for.
537 static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
538  Chipset chipset) {
539 
540  auto sourceVectorType = wmma.getSourceA().getType().dyn_cast<VectorType>();
541  auto destVectorType = wmma.getDestC().getType().dyn_cast<VectorType>();
542  auto elemSourceType = sourceVectorType.getElementType();
543  auto elemDestType = destVectorType.getElementType();
544 
545  if (elemSourceType.isF16() && elemDestType.isF32()) {
546  return ROCDL::wmma_f32_16x16x16_f16::getOperationName();
547  }
548  if (elemSourceType.isBF16() && elemDestType.isF32()) {
549  return ROCDL::wmma_f32_16x16x16_bf16::getOperationName();
550  } else if (elemSourceType.isF16() && elemDestType.isF16()) {
551  return ROCDL::wmma_f16_16x16x16_f16::getOperationName();
552  } else if (elemSourceType.isBF16() && elemDestType.isBF16()) {
553  return ROCDL::wmma_bf16_16x16x16_bf16::getOperationName();
554  } else if (elemSourceType.isInteger(8) && elemDestType.isInteger(32)) {
555  return ROCDL::wmma_i32_16x16x16_iu8::getOperationName();
556  }
557  return std::nullopt;
558 }
559 
560 namespace {
561 struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
562  MFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
563  : ConvertOpToLLVMPattern<MFMAOp>(converter), chipset(chipset) {}
564 
565  Chipset chipset;
566 
568  matchAndRewrite(MFMAOp op, MFMAOpAdaptor adaptor,
569  ConversionPatternRewriter &rewriter) const override {
570  Location loc = op.getLoc();
571  Type outType = typeConverter->convertType(op.getDestD().getType());
572  Type intrinsicOutType = outType;
573  if (auto outVecType = dyn_cast<VectorType>(outType))
574  if (outVecType.getElementType().isBF16())
575  intrinsicOutType = outVecType.clone(rewriter.getI16Type());
576 
577  if (chipset.majorVersion != 9 || chipset.minorVersion < 0x08)
578  return op->emitOpError("MFMA only supported on gfx908+");
579  uint32_t getBlgpField = static_cast<uint32_t>(op.getBlgp());
580  if (op.getNegateA() || op.getNegateB() || op.getNegateC()) {
581  if (chipset.minorVersion < 0x40)
582  return op.emitOpError("negation unsupported on older than gfx840");
583  getBlgpField |=
584  op.getNegateA() | (op.getNegateB() << 1) | (op.getNegateC() << 2);
585  }
586  std::optional<StringRef> maybeIntrinsic = mfmaOpToIntrinsic(op, chipset);
587  if (!maybeIntrinsic.has_value())
588  return op.emitOpError("no intrinsic matching MFMA size on given chipset");
589  OperationState loweredOp(loc, *maybeIntrinsic);
590  loweredOp.addTypes(intrinsicOutType);
591  loweredOp.addOperands(
592  {mfmaConcatIfNeeded(rewriter, loc, adaptor.getSourceA()),
593  mfmaConcatIfNeeded(rewriter, loc, adaptor.getSourceB()),
594  adaptor.getDestC(), createI32Constant(rewriter, loc, op.getCbsz()),
595  createI32Constant(rewriter, loc, op.getAbid()),
596  createI32Constant(rewriter, loc, getBlgpField)});
597  Value lowered = rewriter.create(loweredOp)->getResult(0);
598  if (outType != intrinsicOutType)
599  lowered = rewriter.create<LLVM::BitcastOp>(loc, outType, lowered);
600  rewriter.replaceOp(op, lowered);
601  return success();
602  }
603 };
604 
605 struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
606  WMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
607  : ConvertOpToLLVMPattern<WMMAOp>(converter), chipset(chipset) {}
608 
609  Chipset chipset;
610 
612  matchAndRewrite(WMMAOp op, WMMAOpAdaptor adaptor,
613  ConversionPatternRewriter &rewriter) const override {
614  Location loc = op.getLoc();
615  Type outType = typeConverter->convertType(op.getDestD().getType());
616 
617  if (chipset.majorVersion != 11)
618  return op->emitOpError("WMMA only supported on gfx11");
619 
620  std::optional<StringRef> maybeIntrinsic = wmmaOpToIntrinsic(op, chipset);
621 
622  if (!maybeIntrinsic.has_value())
623  return op.emitOpError("no intrinsic matching WMMA on the given chipset");
624 
625  OperationState loweredOp(loc, *maybeIntrinsic);
626  loweredOp.addTypes(outType);
627 
628  SmallVector<Value, 4> operands;
629  wmmaPushInputOperand(rewriter, loc, typeConverter, op.getUnsignedA(),
630  adaptor.getSourceA(), operands);
631  wmmaPushInputOperand(rewriter, loc, typeConverter, op.getUnsignedB(),
632  adaptor.getSourceB(), operands);
633  wmmaPushOutputOperand(rewriter, loc, typeConverter, adaptor.getDestC(),
634  op.getSubwordOffset(), op.getClamp(), operands);
635 
636  loweredOp.addOperands(operands);
637  Operation *lowered = rewriter.create(loweredOp);
638  rewriter.replaceOp(op, lowered->getResults());
639 
640  return success();
641  }
642 };
643 
644 namespace {
645 struct ExtPackedFp8OpLowering final
646  : public ConvertOpToLLVMPattern<ExtPackedFp8Op> {
647  ExtPackedFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
648  : ConvertOpToLLVMPattern<amdgpu::ExtPackedFp8Op>(converter),
649  chipset(chipset) {}
650  Chipset chipset;
651 
653  matchAndRewrite(ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
654  ConversionPatternRewriter &rewriter) const override;
655 };
656 
657 struct PackedTrunc2xFp8OpLowering final
658  : public ConvertOpToLLVMPattern<PackedTrunc2xFp8Op> {
659  PackedTrunc2xFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
660  : ConvertOpToLLVMPattern<amdgpu::PackedTrunc2xFp8Op>(converter),
661  chipset(chipset) {}
662  Chipset chipset;
663 
665  matchAndRewrite(PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
666  ConversionPatternRewriter &rewriter) const override;
667 };
668 
669 struct PackedStochRoundFp8OpLowering final
670  : public ConvertOpToLLVMPattern<PackedStochRoundFp8Op> {
671  PackedStochRoundFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
672  : ConvertOpToLLVMPattern<amdgpu::PackedStochRoundFp8Op>(converter),
673  chipset(chipset) {}
674  Chipset chipset;
675 
677  matchAndRewrite(PackedStochRoundFp8Op op,
678  PackedStochRoundFp8OpAdaptor adaptor,
679  ConversionPatternRewriter &rewriter) const override;
680 };
681 } // end namespace
682 
683 LogicalResult ExtPackedFp8OpLowering::matchAndRewrite(
684  ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
685  ConversionPatternRewriter &rewriter) const {
686  Location loc = op.getLoc();
687  if (chipset.majorVersion != 9 || chipset.minorVersion < 0x40)
688  return rewriter.notifyMatchFailure(
689  loc, "Fp8 conversion instructions are not available on target "
690  "architecture and their emulation is not implemented");
691  Type v4i8 =
692  getTypeConverter()->convertType(VectorType::get(4, rewriter.getI8Type()));
693  Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
694  Type f32 = getTypeConverter()->convertType(op.getResult().getType());
695 
696  Value source = adaptor.getSource();
697  auto sourceVecType = op.getSource().getType().dyn_cast<VectorType>();
698  Type sourceElemType = getElementTypeOrSelf(op.getSource());
699  // Extend to a v4i8
700  if (!sourceVecType || sourceVecType.getNumElements() < 4) {
701  Value longVec = rewriter.create<LLVM::UndefOp>(loc, v4i8);
702  if (!sourceVecType) {
703  longVec = rewriter.create<LLVM::InsertElementOp>(
704  loc, longVec, source, createI32Constant(rewriter, loc, 0));
705  } else {
706  for (int32_t i = 0, e = sourceVecType.getNumElements(); i < e; ++i) {
707  Value idx = createI32Constant(rewriter, loc, i);
708  Value elem = rewriter.create<LLVM::ExtractElementOp>(loc, source, idx);
709  longVec =
710  rewriter.create<LLVM::InsertElementOp>(loc, longVec, elem, idx);
711  }
712  }
713  source = longVec;
714  }
715  Value i32Source = rewriter.create<LLVM::BitcastOp>(loc, i32, source);
716  Value wordSel = createI32Constant(rewriter, loc, op.getIndex());
717  if (sourceElemType.isFloat8E5M2FNUZ()) {
718  rewriter.replaceOpWithNewOp<ROCDL::CvtF32Bf8Op>(op, f32, i32Source,
719  wordSel);
720  } else if (sourceElemType.isFloat8E4M3FNUZ()) {
721  rewriter.replaceOpWithNewOp<ROCDL::CvtF32Fp8Op>(op, f32, i32Source,
722  wordSel);
723  }
724  return success();
725 }
726 
727 LogicalResult PackedTrunc2xFp8OpLowering::matchAndRewrite(
728  PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
729  ConversionPatternRewriter &rewriter) const {
730  Location loc = op.getLoc();
731  if (chipset.majorVersion != 9 || chipset.minorVersion < 0x40)
732  return rewriter.notifyMatchFailure(
733  loc, "Fp8 conversion instructions are not available on target "
734  "architecture and their emulation is not implemented");
735  Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
736 
737  Type resultType = op.getResult().getType();
738  Type resultElemType = getElementTypeOrSelf(resultType);
739 
740  Value sourceA = adaptor.getSourceA();
741  Value sourceB = adaptor.getSourceB();
742  if (!sourceB)
743  sourceB = rewriter.create<LLVM::UndefOp>(loc, sourceA.getType());
744  Value existing = adaptor.getExisting();
745  if (existing)
746  existing = rewriter.create<LLVM::BitcastOp>(loc, i32, existing);
747  else
748  existing = rewriter.create<LLVM::UndefOp>(loc, i32);
749  Value wordSel = createI1Constant(rewriter, loc, op.getWordIndex());
750 
751  Value result;
752  if (resultElemType.isFloat8E5M2FNUZ())
753  result = rewriter.create<ROCDL::CvtPkBf8F32Op>(loc, i32, sourceA, sourceB,
754  existing, wordSel);
755  else if (resultElemType.isFloat8E4M3FNUZ())
756  result = rewriter.create<ROCDL::CvtPkFp8F32Op>(loc, i32, sourceA, sourceB,
757  existing, wordSel);
758 
759  result = rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
760  op, getTypeConverter()->convertType(resultType), result);
761  return success();
762 }
763 
764 LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
765  PackedStochRoundFp8Op op, PackedStochRoundFp8OpAdaptor adaptor,
766  ConversionPatternRewriter &rewriter) const {
767  Location loc = op.getLoc();
768  if (chipset.majorVersion != 9 || chipset.minorVersion < 0x40)
769  return rewriter.notifyMatchFailure(
770  loc, "Fp8 conversion instructions are not available on target "
771  "architecture and their emulation is not implemented");
772  Type i32 = getTypeConverter()->convertType(rewriter.getI32Type());
773 
774  Type resultType = op.getResult().getType();
775  Type resultElemType = getElementTypeOrSelf(resultType);
776 
777  Value source = adaptor.getSource();
778  Value stoch = adaptor.getStochiasticParam();
779  Value existing = adaptor.getExisting();
780  if (existing)
781  existing = rewriter.create<LLVM::BitcastOp>(loc, i32, existing);
782  else
783  existing = rewriter.create<LLVM::UndefOp>(loc, i32);
784  Value byteSel = createI32Constant(rewriter, loc, op.getStoreIndex());
785 
786  Value result;
787  if (resultElemType.isFloat8E5M2FNUZ())
788  result = rewriter.create<ROCDL::CvtSrBf8F32Op>(loc, i32, source, stoch,
789  existing, byteSel);
790  else if (resultElemType.isFloat8E4M3FNUZ())
791  result = rewriter.create<ROCDL::CvtSrFp8F32Op>(loc, i32, source, stoch,
792  existing, byteSel);
793 
794  result = rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
795  op, getTypeConverter()->convertType(resultType), result);
796  return success();
797 }
798 
799 struct ConvertAMDGPUToROCDLPass
800  : public impl::ConvertAMDGPUToROCDLBase<ConvertAMDGPUToROCDLPass> {
801  ConvertAMDGPUToROCDLPass() = default;
802 
803  void runOnOperation() override {
804  MLIRContext *ctx = &getContext();
805  FailureOr<Chipset> maybeChipset = Chipset::parse(chipset);
806  if (failed(maybeChipset)) {
807  emitError(UnknownLoc::get(ctx), "Invalid chipset name: " + chipset);
808  return signalPassFailure();
809  }
810 
811  RewritePatternSet patterns(ctx);
812  LLVMTypeConverter converter(ctx);
813  populateAMDGPUToROCDLConversionPatterns(converter, patterns, *maybeChipset);
815  target.addIllegalDialect<::mlir::amdgpu::AMDGPUDialect>();
816  target.addLegalDialect<::mlir::LLVM::LLVMDialect>();
817  target.addLegalDialect<::mlir::ROCDL::ROCDLDialect>();
818  if (failed(applyPartialConversion(getOperation(), target,
819  std::move(patterns))))
820  signalPassFailure();
821  }
822 };
823 } // namespace
824 
826  RewritePatternSet &patterns,
827  Chipset chipset) {
828  converter.addConversion([](BFloat16Type t) -> Type {
829  return IntegerType::get(t.getContext(), 16);
830  });
831  converter.addConversion([&converter](VectorType t) -> std::optional<Type> {
832  if (!t.getElementType().isBF16())
833  return std::nullopt;
834  return converter.convertType(t.clone(IntegerType::get(t.getContext(), 16)));
835  });
836 
837  patterns.add<LDSBarrierOpLowering>(converter);
838  patterns
839  .add<RawBufferOpLowering<RawBufferLoadOp, ROCDL::RawPtrBufferLoadOp>,
840  RawBufferOpLowering<RawBufferStoreOp, ROCDL::RawPtrBufferStoreOp>,
841  RawBufferOpLowering<RawBufferAtomicFaddOp,
842  ROCDL::RawPtrBufferAtomicFaddOp>,
843  RawBufferOpLowering<RawBufferAtomicFmaxOp,
844  ROCDL::RawPtrBufferAtomicFmaxOp>,
845  RawBufferOpLowering<RawBufferAtomicSmaxOp,
846  ROCDL::RawPtrBufferAtomicSmaxOp>,
847  RawBufferOpLowering<RawBufferAtomicUminOp,
848  ROCDL::RawPtrBufferAtomicUminOp>,
849  RawBufferOpLowering<RawBufferAtomicCmpswapOp,
850  ROCDL::RawPtrBufferAtomicCmpSwap>,
851  MFMAOpLowering, WMMAOpLowering, ExtPackedFp8OpLowering,
852  PackedTrunc2xFp8OpLowering, PackedStochRoundFp8OpLowering>(converter,
853  chipset);
854 }
855 
856 std::unique_ptr<Pass> mlir::createConvertAMDGPUToROCDLPass() {
857  return std::make_unique<ConvertAMDGPUToROCDLPass>();
858 }
static Value mfmaConcatIfNeeded(ConversionPatternRewriter &rewriter, Location loc, Value input)
If input is a vector of bytes, concatentate those bytes in little-endian order to form a single integ...
static std::optional< StringRef > wmmaOpToIntrinsic(WMMAOp wmma, Chipset chipset)
Return the rocdl intrinsic corresponding to a WMMA operation wmma if one exists.
static void wmmaPushInputOperand(ConversionPatternRewriter &rewriter, Location loc, const TypeConverter *typeConverter, bool isUnsigned, Value llvmInput, SmallVector< Value, 4 > &operands)
Push an input operand.
static Value createI1Constant(ConversionPatternRewriter &rewriter, Location loc, bool value)
static std::optional< StringRef > mfmaOpToIntrinsic(MFMAOp mfma, Chipset chipset)
Return the rocdl intrinsic corresponding to a MFMA operation mfma if one exists.
static void wmmaPushOutputOperand(ConversionPatternRewriter &rewriter, Location loc, const TypeConverter *typeConverter, Value output, int32_t subwordOffset, bool clamp, SmallVector< Value, 4 > &operands)
Push the output operand.
static Value createI32Constant(ConversionPatternRewriter &rewriter, Location loc, int32_t value)
static MLIRContext * getContext(OpFoldResult val)
static Value clamp(ImplicitLocOpBuilder &builder, Value value, Value lowerBound, Value upperBound)
IntegerType getI16Type()
Definition: Builders.cpp:81
IntegerAttr getIntegerAttr(Type type, int64_t value)
Definition: Builders.cpp:238
IntegerAttr getI16IntegerAttr(int16_t value)
Definition: Builders.cpp:230
IntegerType getI32Type()
Definition: Builders.cpp:83
IntegerType getIntegerType(unsigned width)
Definition: Builders.cpp:87
MLIRContext * getContext() const
Definition: Builders.h:55
IntegerType getI1Type()
Definition: Builders.cpp:73
IntegerType getI8Type()
Definition: Builders.cpp:79
This class implements a pattern rewriter for use with ConversionPatterns.
void replaceOp(Operation *op, ValueRange newValues) override
PatternRewriter hook for replacing an operation.
LogicalResult notifyMatchFailure(Location loc, function_ref< void(Diagnostic &)> reasonCallback) override
PatternRewriter hook for notifying match failure reasons.
void eraseOp(Operation *op) override
PatternRewriter hook for erasing a dead operation.
Utility class for operation conversions targeting the LLVM dialect that match exactly one source oper...
Definition: Pattern.h:139
This class provides support for representing a failure result, or a valid value of type T.
Definition: LogicalResult.h:78
Derived class that automatically populates legalization information for different LLVM ops.
Conversion from types to the LLVM IR dialect.
Definition: TypeConverter.h:33
LogicalResult convertType(Type t, SmallVectorImpl< Type > &results) const
Convert the given type.
This class defines the main interface for locations in MLIR and acts as a non-nullable wrapper around...
Definition: Location.h:63
MLIRContext is the top-level object for a collection of MLIR operations.
Definition: MLIRContext.h:60
Helper class to produce LLVM dialect operations extracting or inserting elements of a MemRef descript...
Definition: MemRefBuilder.h:33
void createOrFold(SmallVectorImpl< Value > &results, Location location, Args &&...args)
Create an operation of specific op type at the current insertion point, and immediately try to fold i...
Definition: Builders.h:505
Operation * create(const OperationState &state)
Creates an operation given the fields represented as an OperationState.
Definition: Builders.cpp:446
Operation is the basic unit of execution within MLIR.
Definition: Operation.h:88
OpResult getResult(unsigned idx)
Get the 'idx'th result of this operation.
Definition: Operation.h:402
Location getLoc()
The source location the operation was defined or derived from.
Definition: Operation.h:223
result_range getResults()
Definition: Operation.h:410
InFlightDiagnostic emitOpError(const Twine &message={})
Emit an error with the op name prefixed, like "'dim' op " which is convenient for verifiers.
Definition: Operation.cpp:640
unsigned getNumResults()
Return the number of results held by this operation.
Definition: Operation.h:399
RewritePatternSet & add(ConstructorArg &&arg, ConstructorArgs &&...args)
Add an instance of each of the pattern types 'Ts' to the pattern list with the given arguments.
OpTy replaceOpWithNewOp(Operation *op, Args &&...args)
Replaces the result op with a new op that is created without verification.
Definition: PatternMatch.h:539
Type conversion class.
void addConversion(FnT &&callback)
Register a conversion function.
LogicalResult convertType(Type t, SmallVectorImpl< Type > &results) const
Convert the given type.
This class provides an abstraction over the various different ranges of value types.
Definition: TypeRange.h:36
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:52
bool isInteger(unsigned width) const
Return true if this is an integer type with the specified width.
Definition: Types.cpp:59
U dyn_cast() const
Definition: Types.h:329
bool isF32() const
Definition: Types.cpp:51
bool isFloat8E4M3FNUZ() const
Definition: Types.cpp:42
bool isF16() const
Definition: Types.cpp:49
bool isBF16() const
Definition: Types.cpp:48
bool isFloat8E5M2FNUZ() const
Definition: Types.cpp:39
This class provides an abstraction over the different types of ranges over Values.
Definition: ValueRange.h:378
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:125
constexpr void enumerate(std::tuple< Tys... > &tuple, CallbackT &&callback)
Definition: Matchers.h:285
Include the generated interface declarations.
LogicalResult applyPartialConversion(ArrayRef< Operation * > ops, const ConversionTarget &target, const FrozenRewritePatternSet &patterns, DenseSet< Operation * > *unconvertedOps=nullptr)
Below we define several entry points for operation conversion.
std::unique_ptr< Pass > createConvertAMDGPUToROCDLPass()
InFlightDiagnostic emitError(Location loc)
Utility method to emit an error message using this location.
LogicalResult getStridesAndOffset(MemRefType t, SmallVectorImpl< int64_t > &strides, int64_t &offset)
Returns the strides of the MemRef if the layout map is in strided form.
LogicalResult success(bool isSuccess=true)
Utility function to generate a LogicalResult.
Definition: LogicalResult.h:56
Type getElementTypeOrSelf(Type type)
Return the element type or return the type itself.
void populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter, RewritePatternSet &patterns, amdgpu::Chipset chipset)
Note: The ROCDL target does not support the LLVM bfloat type at this time and so this function will a...
auto get(MLIRContext *context, Ts &&...params)
Helper method that injects context only if needed, this helps unify some of the attribute constructio...
bool failed(LogicalResult result)
Utility function that returns true if the provided LogicalResult corresponds to a failure value.
Definition: LogicalResult.h:72
This class represents an efficient way to signal success or failure.
Definition: LogicalResult.h:26
This represents an operation in an abstracted form, suitable for use with the builder APIs.
static FailureOr< Chipset > parse(StringRef name)
Definition: Chipset.cpp:16
unsigned majorVersion
Definition: Chipset.h:21
unsigned minorVersion
Definition: Chipset.h:22