MLIR 24.0.0git
GPUDialect.cpp
Go to the documentation of this file.
1//===- GPUDialect.cpp - MLIR Dialect for GPU Kernels implementation -------===//
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 implements the GPU kernel-related dialect and its operations.
10//
11//===----------------------------------------------------------------------===//
12
14
21#include "mlir/IR/Attributes.h"
22#include "mlir/IR/Builders.h"
24#include "mlir/IR/BuiltinOps.h"
26#include "mlir/IR/Diagnostics.h"
28#include "mlir/IR/Matchers.h"
31#include "mlir/IR/SymbolTable.h"
37#include "llvm/ADT/STLExtras.h"
38#include "llvm/ADT/TypeSwitch.h"
39#include "llvm/Support/ErrorHandling.h"
40#include "llvm/Support/FormatVariadic.h"
41#include "llvm/Support/InterleavedRange.h"
42#include <cassert>
43#include <numeric>
44#include <optional>
45
46using namespace mlir;
47using namespace mlir::gpu;
48
49#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
50
51namespace {
52// Attribute storage definitions live in MLIRGPUBase. Keep interface metadata
53// here: some attribute interfaces build GPU/Arith/Math IR, and their models
54// belong with the operation implementation rather than the representations.
55template <typename... Attrs>
56void registerAttributeMetadata(
57 Dialect &dialect,
58 function_ref<void(TypeID, AbstractAttribute &&)> addAttribute) {
59 (addAttribute(Attrs::getTypeID(), AbstractAttribute::get<Attrs>(dialect)),
60 ...);
61}
62
63struct GPUDialectBytecodeInterface : BytecodeDialectInterface {
64 using BytecodeDialectInterface::BytecodeDialectInterface;
65
66 void writeVersion(DialectBytecodeWriter &writer) const override {
67 writer.writeVarInt(1);
68 }
69
70 std::unique_ptr<DialectVersion>
71 readVersion(DialectBytecodeReader &reader) const override {
72 uint64_t version;
73 if (failed(reader.readVarInt(version)))
74 return nullptr;
75 if (version != 1) {
76 reader.emitError("unsupported gpu dialect bytecode version");
77 return nullptr;
78 }
79 return std::make_unique<DialectVersion>();
80 }
81};
82} // namespace
83
84int64_t GPUMappingMaskAttr::getMaxNumPhysicalIds() const { return 64; }
85
86/// 8 4 0
87/// Example mask : 0 0 0 1 1 0 1 0 0
88///
89/// Active physical (resp. logical) is 2 (0), 4 (1) and 5 (2).
90/// Logical id for e.g. 5 (2) constructs filter (1 << 5 - 1).
91///
92/// Example mask : 0 0 0 1 1 0 1 0 0
93/// Example filter: 0 0 0 0 1 1 1 1 1
94/// Intersection : 0 0 0 0 1 0 1 0 0
95/// PopCnt : 2
96Value GPUMappingMaskAttr::createLogicalLinearMappingId(
97 OpBuilder &b, Value physicalLinearMappingId) const {
98 Location loc = physicalLinearMappingId.getLoc();
99 Value mask =
100 arith::ConstantOp::create(b, loc, b.getI64IntegerAttr(getMask()));
101 Value one = arith::ConstantOp::create(b, loc, b.getI64IntegerAttr(1));
102 Value filter = arith::ShLIOp::create(b, loc, one, physicalLinearMappingId);
103 filter = arith::SubIOp::create(b, loc, filter, one);
104 Value filteredId = arith::AndIOp::create(b, loc, mask, filter);
105 return math::CtPopOp::create(b, loc, filteredId);
106}
107
108/// 8 4 0
109/// Example mask : 0 0 0 1 1 0 1 0 0
110///
111/// Active physical (resp. logical) is 2 (0), 4 (1) and 5 (2).
112/// Logical id for e.g. 5 (2) constructs filter (1 << 5).
113///
114/// Example mask : 0 0 0 1 1 0 1 0 0
115/// Example filter: 0 0 0 1 0 0 0 0 0
116/// Intersection : 0 0 0 1 0 0 0 0 0
117/// Cmp : 1
118Value GPUMappingMaskAttr::createIsActiveIdPredicate(
119 OpBuilder &b, Value physicalLinearMappingId) const {
120 Location loc = physicalLinearMappingId.getLoc();
121 Value mask =
122 arith::ConstantOp::create(b, loc, b.getI64IntegerAttr(getMask()));
123 Value one = arith::ConstantOp::create(b, loc, b.getI64IntegerAttr(1));
124 Value filter = arith::ShLIOp::create(b, loc, one, physicalLinearMappingId);
125 Value filtered = arith::AndIOp::create(b, loc, mask, filter);
126 Value zero = arith::ConstantOp::create(b, loc, b.getI64IntegerAttr(0));
127 return arith::CmpIOp::create(b, loc, arith::CmpIPredicate::ne, filtered,
128 zero);
129}
130
131bool GPUDialect::isKernel(Operation *op) {
132 if (auto gpuFunc = dyn_cast<GPUFuncOp>(op))
133 return gpuFunc.isKernel();
134 return static_cast<bool>(
135 op->getDiscardableAttrOfType<UnitAttr>(getKernelFuncAttrName()));
136}
137
138namespace {
139/// This class defines the interface for handling inlining with gpu
140/// operations.
141struct GPUInlinerInterface : public DialectInlinerInterface {
142 using DialectInlinerInterface::DialectInlinerInterface;
143
144 /// All gpu dialect ops can be inlined.
145 bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
146 return true;
147 }
148};
149} // namespace
150
151void GPUDialect::initialize() {
152 addInterfaces<GPUDialectBytecodeInterface>();
153 addTypes<AsyncTokenType>();
154 addTypes<MMAMatrixType>();
155 addTypes<NamedBarrierType>();
156 addTypes<SparseDnTensorHandleType>();
157 addTypes<SparseSpMatHandleType>();
158 addTypes<SparseSpGEMMOpHandleType>();
159 addOperations<
160#define GET_OP_LIST
161#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
162 >();
163 // Match addAttributes while keeping the storage and interface instantiations
164 // in their respective implementation libraries.
165 registerAttributeMetadata<
166#define GET_ATTRDEF_LIST
167#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
168 >(*this, [this](TypeID id, AbstractAttribute &&attribute) {
169 addAttribute(id, std::move(attribute));
170 });
172 addInterfaces<GPUInlinerInterface>();
173 declarePromisedInterface<bufferization::BufferDeallocationOpInterface,
174 TerminatorOp>();
175 declarePromisedInterfaces<ValueBoundsOpInterface, ClusterDimOp,
176 ClusterDimBlocksOp, ClusterIdOp, ClusterBlockIdOp,
177 BlockDimOp, BlockIdOp, GridDimOp, ThreadIdOp,
178 LaneIdOp, SubgroupIdOp, GlobalIdOp, NumSubgroupsOp,
179 SubgroupSizeOp, LaunchOp, SubgroupBroadcastOp>();
180 declarePromisedInterfaces<memref::IndexedAccessOpInterface,
181 SubgroupMmaLoadMatrixOp,
182 SubgroupMmaStoreMatrixOp>();
183}
184
185static std::string getSparseHandleKeyword(SparseHandleKind kind) {
186 switch (kind) {
188 return "sparse.dntensor_handle";
190 return "sparse.spmat_handle";
192 return "sparse.spgemmop_handle";
193 }
194 llvm_unreachable("unknown sparse handle kind");
195 return "";
196}
197
198Type GPUDialect::parseType(DialectAsmParser &parser) const {
199 // Parse the main keyword for the type.
200 StringRef keyword;
201 if (parser.parseKeyword(&keyword))
202 return Type();
203 MLIRContext *context = getContext();
204
205 // Handle 'async token' types.
206 if (keyword == "async.token")
207 return AsyncTokenType::get(context);
208
209 if (keyword == "mma_matrix") {
210 SMLoc beginLoc = parser.getNameLoc();
211
212 // Parse '<'.
213 if (parser.parseLess())
214 return nullptr;
215
216 // Parse the size and elementType.
218 Type elementType;
219 if (parser.parseDimensionList(shape, /*allowDynamic=*/false) ||
220 parser.parseType(elementType))
221 return nullptr;
222
223 // Parse ','
224 if (parser.parseComma())
225 return nullptr;
226
227 // Parse operand.
228 std::string operand;
229 if (failed(parser.parseOptionalString(&operand)))
230 return nullptr;
231
232 // Parse '>'.
233 if (parser.parseGreater())
234 return nullptr;
235
237 parser.getEncodedSourceLoc(beginLoc)),
238 shape, elementType, operand);
239 }
240
241 if (keyword == "named_barrier")
242 return NamedBarrierType::get(context);
243
245 return SparseDnTensorHandleType::get(context);
247 return SparseSpMatHandleType::get(context);
249 return SparseSpGEMMOpHandleType::get(context);
250
251 parser.emitError(parser.getNameLoc(), "unknown gpu type: " + keyword);
252 return Type();
253}
254// TODO: print refined type here. Notice that should be corresponding to the
255// parser
256void GPUDialect::printType(Type type, DialectAsmPrinter &os) const {
257 TypeSwitch<Type>(type)
258 .Case<AsyncTokenType>([&](Type) { os << "async.token"; })
259 .Case<NamedBarrierType>([&](Type) { os << "named_barrier"; })
260 .Case<SparseDnTensorHandleType>([&](Type) {
262 })
263 .Case<SparseSpMatHandleType>(
265 .Case<SparseSpGEMMOpHandleType>([&](Type) {
267 })
268 .Case([&](MMAMatrixType fragTy) {
269 os << "mma_matrix<";
270 auto shape = fragTy.getShape();
271 for (auto dim = shape.begin(), e = shape.end() - 1; dim != e; ++dim)
272 os << *dim << 'x';
273 os << shape.back() << 'x' << fragTy.getElementType();
274 os << ", \"" << fragTy.getOperand() << "\"" << '>';
275 })
276 .DefaultUnreachable("unexpected 'gpu' type kind");
277}
278
279static LogicalResult verifyKnownLaunchSizeAttr(Operation *op,
280 NamedAttribute attr) {
281 auto array = dyn_cast<DenseI32ArrayAttr>(attr.getValue());
282 if (!array)
283 return op->emitOpError(Twine(attr.getName()) +
284 " must be a dense i32 array");
285 if (array.size() != 3)
286 return op->emitOpError(Twine(attr.getName()) +
287 " must contain exactly 3 elements");
288 return success();
289}
290
291LogicalResult GPUDialect::verifyOperationAttribute(Operation *op,
292 NamedAttribute attr) {
293 if (attr.getName() == getKnownBlockSizeAttrHelper().getName())
294 return verifyKnownLaunchSizeAttr(op, attr);
295 if (attr.getName() == getKnownGridSizeAttrHelper().getName())
296 return verifyKnownLaunchSizeAttr(op, attr);
297 if (attr.getName() == getKnownClusterSizeAttrHelper().getName())
298 return verifyKnownLaunchSizeAttr(op, attr);
299 if (!llvm::isa<UnitAttr>(attr.getValue()) ||
300 attr.getName() != getContainerModuleAttrName())
301 return success();
302
303 auto module = dyn_cast<ModuleOp>(op);
304 if (!module)
305 return op->emitError("expected '")
306 << getContainerModuleAttrName() << "' attribute to be attached to '"
307 << ModuleOp::getOperationName() << '\'';
308 return success();
309}
310
311/// Parses an optional list of async operands with an optional leading keyword.
312/// (`async`)? (`[` ssa-id-list `]`)?
313///
314/// This method is used by the tablegen assembly format for async ops as well.
315static ParseResult parseAsyncDependencies(
316 OpAsmParser &parser, Type &asyncTokenType,
318 auto loc = parser.getCurrentLocation();
319 if (succeeded(parser.parseOptionalKeyword("async"))) {
320 if (parser.getNumResults() == 0)
321 return parser.emitError(loc, "needs to be named when marked 'async'");
322 asyncTokenType = parser.getBuilder().getType<AsyncTokenType>();
323 }
324 return parser.parseOperandList(asyncDependencies,
326}
327
328/// Prints optional async dependencies with its leading keyword.
329/// (`async`)? (`[` ssa-id-list `]`)?
330// Used by the tablegen assembly format for several async ops.
332 Type asyncTokenType,
333 OperandRange asyncDependencies) {
334 if (asyncTokenType)
335 printer << "async";
336 if (asyncDependencies.empty())
337 return;
338 if (asyncTokenType)
339 printer << ' ';
340 printer << llvm::interleaved_array(asyncDependencies);
341}
342
343// GPU Memory attributions functions shared by LaunchOp and GPUFuncOp.
344/// Parses a GPU function memory attribution.
345///
346/// memory-attribution ::= (`workgroup` `(` ssa-id-and-type-list `)`)?
347/// (`private` `(` ssa-id-and-type-list `)`)?
348///
349/// Note that this function parses only one of the two similar parts, with the
350/// keyword provided as argument.
351static ParseResult
352parseAttributions(OpAsmParser &parser, StringRef keyword,
354 // If we could not parse the keyword, just assume empty list and succeed.
355 if (failed(parser.parseOptionalKeyword(keyword)))
356 return success();
357
359 /*allowType=*/true);
360}
361
362static void printAttributions(OpAsmPrinter &p, StringRef keyword,
364 ArrayAttr attributes = {}) {
365 if (values.empty())
366 return;
367
368 p << ' ' << keyword << '(';
369 llvm::interleaveComma(
370 llvm::enumerate(values), p, [&p, attributes](auto pair) {
371 BlockArgument v = pair.value();
372 p << v << " : " << v.getType();
373
374 size_t attributionIndex = pair.index();
375 DictionaryAttr attrs;
376 if (attributes && attributionIndex < attributes.size())
377 attrs = llvm::cast<DictionaryAttr>(attributes[attributionIndex]);
378 if (attrs)
379 p.printOptionalAttrDict(attrs.getValue());
380 });
381 p << ')';
382}
383
384/// Verifies a GPU function memory attribution.
385static LogicalResult verifyAttributions(Operation *op,
386 ArrayRef<BlockArgument> attributions,
387 gpu::AddressSpace memorySpace) {
388 for (Value v : attributions) {
389 auto type = llvm::dyn_cast<MemRefType>(v.getType());
390 if (!type)
391 return op->emitOpError() << "expected memref type in attribution";
392
393 // We can only verify the address space if it hasn't already been lowered
394 // from the AddressSpaceAttr to a target-specific numeric value.
395 auto addressSpace =
396 llvm::dyn_cast_or_null<gpu::AddressSpaceAttr>(type.getMemorySpace());
397 if (!addressSpace)
398 continue;
399 if (addressSpace.getValue() != memorySpace)
400 return op->emitOpError()
401 << "expected memory space " << stringifyAddressSpace(memorySpace)
402 << " in attribution";
403 }
404 return success();
405}
406
407//===----------------------------------------------------------------------===//
408// AllReduceOp
409//===----------------------------------------------------------------------===//
410
411static LogicalResult verifyReduceOpAndType(gpu::AllReduceOperation opName,
412 Type resType) {
413 using Kind = gpu::AllReduceOperation;
414 if (llvm::is_contained({Kind::MINNUMF, Kind::MAXNUMF, Kind::MINIMUMF,
415 Kind::MAXIMUMF, Kind::MINIMUMNUMF, Kind::MAXIMUMNUMF},
416 opName)) {
417 if (!isa<FloatType>(resType))
418 return failure();
419 }
420
421 if (llvm::is_contained({Kind::MINSI, Kind::MINUI, Kind::MAXSI, Kind::MAXUI,
422 Kind::AND, Kind::OR, Kind::XOR},
423 opName)) {
424 if (!isa<IntegerType>(resType))
425 return failure();
426 }
427
428 return success();
429}
430
431LogicalResult gpu::AllReduceOp::verifyRegions() {
432 if (getBody().empty() != getOp().has_value())
433 return emitError("expected either an op attribute or a non-empty body");
434 if (!getBody().empty()) {
435 if (getBody().getNumArguments() != 2)
436 return emitError("expected two region arguments");
437 for (auto argument : getBody().getArguments()) {
438 if (argument.getType() != getType())
439 return emitError("incorrect region argument type");
440 }
441 unsigned yieldCount = 0;
442 for (Block &block : getBody()) {
443 if (auto yield = dyn_cast<gpu::YieldOp>(block.getTerminator())) {
444 if (yield.getNumOperands() != 1)
445 return emitError("expected one gpu.yield operand");
446 if (yield.getOperand(0).getType() != getType())
447 return emitError("incorrect gpu.yield type");
448 ++yieldCount;
449 }
450 }
451 if (yieldCount == 0)
452 return emitError("expected gpu.yield op in region");
453 } else {
454 gpu::AllReduceOperation opName = *getOp();
455 if (failed(verifyReduceOpAndType(opName, getType()))) {
456 return emitError() << '`' << gpu::stringifyAllReduceOperation(opName)
457 << "` reduction operation is not compatible with type "
458 << getType();
459 }
460 }
461
462 return success();
463}
464
466 auto launchOp = dyn_cast<gpu::LaunchOp>(op->getParentOp());
467 if (!launchOp)
468 return false;
469
470 Region &body = launchOp.getBody();
471 assert(!body.empty() && "Invalid region");
472
473 // Only convert ops in gpu::launch entry block for now.
474 return op->getBlock() == &body.front();
475}
476
477OpFoldResult gpu::AllReduceOp::fold(FoldAdaptor /*adaptor*/) {
478 if (!getUniform() && canMakeGroupOpUniform(*this)) {
479 setUniform(true);
480 return getResult();
481 }
482
483 return nullptr;
484}
485
486//===----------------------------------------------------------------------===//
487// SubgroupReduceOp
488//===----------------------------------------------------------------------===//
489
490LogicalResult gpu::SubgroupReduceOp::verify() {
491 Type elemType = getType();
492 if (auto vecTy = dyn_cast<VectorType>(elemType)) {
493 if (vecTy.isScalable())
494 return emitOpError() << "is not compatible with scalable vector types";
495
496 elemType = vecTy.getElementType();
497 }
498
499 gpu::AllReduceOperation opName = getOp();
500 if (failed(verifyReduceOpAndType(opName, elemType))) {
501 return emitError() << '`' << gpu::stringifyAllReduceOperation(opName)
502 << "` reduction operation is not compatible with type "
503 << getType();
504 }
505
506 auto clusterSize = getClusterSize();
507 if (clusterSize) {
508 uint32_t size = *clusterSize;
509 if (!llvm::isPowerOf2_32(size)) {
510 return emitOpError() << "cluster size " << size
511 << " is not a power of two";
512 }
513 }
514
515 uint32_t stride = getClusterStride();
516 if (stride != 1 && !clusterSize) {
517 return emitOpError() << "cluster stride can only be specified if cluster "
518 "size is specified";
519 }
520 if (!llvm::isPowerOf2_32(stride)) {
521 return emitOpError() << "cluster stride " << stride
522 << " is not a power of two";
523 }
524
525 return success();
526}
527
528OpFoldResult gpu::SubgroupReduceOp::fold(FoldAdaptor /*adaptor*/) {
529 if (getClusterSize() == 1)
530 return getValue();
531
532 if (!getUniform() && canMakeGroupOpUniform(*this)) {
533 setUniform(true);
534 return getResult();
535 }
536
537 return nullptr;
538}
539
540//===----------------------------------------------------------------------===//
541// AsyncOpInterface
542//===----------------------------------------------------------------------===//
543
545 op->insertOperands(0, {token});
546 if (!op->template hasTrait<OpTrait::AttrSizedOperandSegments>())
547 return;
548 auto attrName =
550 auto sizeAttr = dyn_cast_or_null<DenseI32ArrayAttr>(
551 op->getInherentAttr(attrName).value_or(Attribute{}));
552
553 // Async dependencies is the only variadic operand.
554 if (!sizeAttr)
555 return;
556
557 SmallVector<int32_t, 8> sizes(sizeAttr.asArrayRef());
558 ++sizes.front();
559 op->setInherentAttr(StringAttr::get(op->getContext(), attrName),
561}
562
563//===----------------------------------------------------------------------===//
564// LaunchOp
565//===----------------------------------------------------------------------===//
566
567void LaunchOp::build(OpBuilder &builder, OperationState &result,
568 Value gridSizeX, Value gridSizeY, Value gridSizeZ,
569 Value getBlockSizeX, Value getBlockSizeY,
570 Value getBlockSizeZ, Value dynamicSharedMemorySize,
571 Type asyncTokenType, ValueRange asyncDependencies,
572 Value asyncObject, TypeRange workgroupAttributions,
573 TypeRange privateAttributions, Value clusterSizeX,
574 Value clusterSizeY, Value clusterSizeZ,
575 FlatSymbolRefAttr module, FlatSymbolRefAttr function) {
576 OpBuilder::InsertionGuard g(builder);
577
578 if (!workgroupAttributions.empty())
579 result.addAttribute(
580 getWorkgroupAttributionsAttrName(result.name),
581 builder.getI64IntegerAttr(workgroupAttributions.size()));
582
583 // Add Op operands.
584 result.addOperands(asyncDependencies);
585 if (asyncTokenType)
586 result.types.push_back(builder.getType<AsyncTokenType>());
587
588 // Add grid and block sizes as op operands, followed by the data operands.
589 result.addOperands({gridSizeX, gridSizeY, gridSizeZ, getBlockSizeX,
590 getBlockSizeY, getBlockSizeZ});
591 if (clusterSizeX)
592 result.addOperands(clusterSizeX);
593 if (clusterSizeY)
594 result.addOperands(clusterSizeY);
595 if (clusterSizeZ)
596 result.addOperands(clusterSizeZ);
597 if (dynamicSharedMemorySize)
598 result.addOperands(dynamicSharedMemorySize);
599 if (asyncObject)
600 result.addOperands(asyncObject);
601
602 // Add optional module and function attributes.
603 if (module)
604 result.addAttribute(getModuleAttrName(result.name), module);
605 if (function)
606 result.addAttribute(getFunctionAttrName(result.name), function);
607
608 // Create a kernel body region with kNumConfigRegionAttributes + N memory
609 // attributions, where the first kNumConfigRegionAttributes arguments have
610 // `index` type and the rest have the same types as the data operands.
611 Region *kernelRegion = result.addRegion();
612 Block *body = builder.createBlock(kernelRegion);
613 // TODO: Allow passing in proper locations here.
614 for (unsigned i = 0; i < kNumConfigRegionAttributes; ++i)
615 body->addArgument(builder.getIndexType(), result.location);
616 // Add WorkGroup & Private attributions to the region arguments.
617 for (Type argTy : workgroupAttributions)
618 body->addArgument(argTy, result.location);
619 for (Type argTy : privateAttributions)
620 body->addArgument(argTy, result.location);
621 // Fill OperandSegmentSize Attribute.
622 SmallVector<int32_t, 12> segmentSizes(12, 1);
623 segmentSizes.front() = asyncDependencies.size();
624 segmentSizes[7] = clusterSizeX ? 1 : 0;
625 segmentSizes[8] = clusterSizeY ? 1 : 0;
626 segmentSizes[9] = clusterSizeZ ? 1 : 0;
627 segmentSizes[10] = dynamicSharedMemorySize ? 1 : 0;
628 segmentSizes[11] = asyncObject ? 1 : 0;
629 result.addAttribute(getOperandSegmentSizeAttr(),
630 builder.getDenseI32ArrayAttr(segmentSizes));
631}
632
633KernelDim3 LaunchOp::getBlockIds() {
634 assert(!getBody().empty() && "LaunchOp body must not be empty.");
635 auto args = getBody().getArguments();
636 return KernelDim3{args[0], args[1], args[2]};
637}
638
639KernelDim3 LaunchOp::getThreadIds() {
640 assert(!getBody().empty() && "LaunchOp body must not be empty.");
641 auto args = getBody().getArguments();
642 return KernelDim3{args[3], args[4], args[5]};
643}
644
645KernelDim3 LaunchOp::getGridSize() {
646 assert(!getBody().empty() && "LaunchOp body must not be empty.");
647 auto args = getBody().getArguments();
648 return KernelDim3{args[6], args[7], args[8]};
649}
650
651KernelDim3 LaunchOp::getBlockSize() {
652 assert(!getBody().empty() && "LaunchOp body must not be empty.");
653 auto args = getBody().getArguments();
654 return KernelDim3{args[9], args[10], args[11]};
655}
656
657std::optional<KernelDim3> LaunchOp::getClusterIds() {
658 assert(!getBody().empty() && "LaunchOp body must not be empty.");
659 if (!hasClusterSize())
660 return std::nullopt;
661 auto args = getBody().getArguments();
662 return KernelDim3{args[12], args[13], args[14]};
663}
664
665std::optional<KernelDim3> LaunchOp::getClusterSize() {
666 assert(!getBody().empty() && "LaunchOp body must not be empty.");
667 if (!hasClusterSize())
668 return std::nullopt;
669 auto args = getBody().getArguments();
670 return KernelDim3{args[15], args[16], args[17]};
671}
672
673KernelDim3 LaunchOp::getGridSizeOperandValues() {
674 auto operands = getOperands().drop_front(getAsyncDependencies().size());
675 return KernelDim3{operands[0], operands[1], operands[2]};
676}
677
678KernelDim3 LaunchOp::getBlockSizeOperandValues() {
679 auto operands = getOperands().drop_front(getAsyncDependencies().size());
680 return KernelDim3{operands[3], operands[4], operands[5]};
681}
682
683std::optional<KernelDim3> LaunchOp::getClusterSizeOperandValues() {
684 auto operands = getOperands().drop_front(getAsyncDependencies().size());
685 if (!hasClusterSize())
686 return std::nullopt;
687 return KernelDim3{operands[6], operands[7], operands[8]};
688}
689
690template <typename OpTy>
691static LogicalResult verifyLaunchAsyncModel(OpTy op) {
692 if (!op.getAsyncDependencies().empty() && !op.getAsyncToken())
693 return op.emitOpError("dependency operands require the dependency-based "
694 "async model i.e. returning a token");
695 if (op.getAsyncToken() && op.getAsyncObject())
696 return op.emitOpError("stream-based and dependency-based async models are "
697 "mutually exclusive");
698 if (op.getNumResults() == 0 && op.getAsyncToken())
699 return op.emitOpError("needs to be named when async keyword is specified");
700 return success();
701}
702
703LogicalResult LaunchOp::verify() {
704 if (verifyLaunchAsyncModel(*this).failed())
705 return failure();
706
707 if (!(hasClusterSize()) &&
708 (getClusterSizeX() || getClusterSizeY() || getClusterSizeZ()))
709 return emitOpError() << "cluster size must be all present";
710 return success();
711}
712
713LogicalResult LaunchOp::verifyRegions() {
714 // Kernel launch takes kNumConfigOperands leading operands for grid/block
715 // sizes and transforms them into kNumConfigRegionAttributes region arguments
716 // for block/thread identifiers and grid/block sizes.
717 if (getBody().empty()) {
718 return emitOpError("body region is empty");
719 }
720 unsigned actualNumRegionArgs = getBody().getNumArguments();
721 unsigned expectedNumRegionArgs =
722 getNumConfigRegionAttributes() + getNumWorkgroupAttributions();
723 if (actualNumRegionArgs < expectedNumRegionArgs) {
724 return emitOpError("expected at least ")
725 << expectedNumRegionArgs << " region arguments, but got "
726 << actualNumRegionArgs;
727 }
728
729 // Verify Attributions Address Spaces.
730 if (failed(verifyAttributions(getOperation(), getWorkgroupAttributionBBArgs(),
731 GPUDialect::getWorkgroupAddressSpace())) ||
732 failed(verifyAttributions(getOperation(), getPrivateAttributions(),
733 GPUDialect::getPrivateAddressSpace())))
734 return failure();
735
736 // Block terminators without successors are expected to exit the kernel region
737 // and must be `gpu.terminator`.
738 for (Block &block : getBody()) {
739 if (block.empty())
740 continue;
741 if (block.back().getNumSuccessors() != 0)
742 continue;
743 if (!isa<gpu::TerminatorOp>(&block.back())) {
744 return block.back()
745 .emitError()
746 .append("expected '", gpu::TerminatorOp::getOperationName(),
747 "' or a terminator with successors")
748 .attachNote(getLoc())
749 .append("in '", LaunchOp::getOperationName(), "' body region");
750 }
751 }
752
753 return success();
754}
755
756// Pretty-print the kernel grid/block size assignment as
757// (%iter-x, %iter-y, %iter-z) in
758// (%size-x = %ssa-use, %size-y = %ssa-use, %size-z = %ssa-use)
759// where %size-* and %iter-* will correspond to the body region arguments.
761 KernelDim3 operands, KernelDim3 ids) {
762 p << '(' << ids.x << ", " << ids.y << ", " << ids.z << ") in (";
763 p << size.x << " = " << operands.x << ", ";
764 p << size.y << " = " << operands.y << ", ";
765 p << size.z << " = " << operands.z << ')';
766}
767
768void LaunchOp::print(OpAsmPrinter &p) {
769 if (auto asyncObject = getAsyncObject()) {
770 p << " <" << asyncObject << " : " << asyncObject.getType() << ">";
771 }
772 if (getAsyncToken()) {
773 p << " async";
774 if (!getAsyncDependencies().empty())
775 p << " [" << getAsyncDependencies() << ']';
776 }
777 // Print the launch configuration.
778 if (hasClusterSize()) {
779 p << ' ' << getClustersKeyword();
780 printSizeAssignment(p, getClusterSize().value(),
781 getClusterSizeOperandValues().value(),
782 getClusterIds().value());
783 }
784 p << ' ' << getBlocksKeyword();
785 printSizeAssignment(p, getGridSize(), getGridSizeOperandValues(),
786 getBlockIds());
787 p << ' ' << getThreadsKeyword();
788 printSizeAssignment(p, getBlockSize(), getBlockSizeOperandValues(),
789 getThreadIds());
790 if (getDynamicSharedMemorySize())
791 p << ' ' << getDynamicSharedMemorySizeKeyword() << ' '
792 << getDynamicSharedMemorySize();
793
794 // Print optional module attribute.
795 StringRef moduleAttrName = getModuleAttrName();
796 if (auto module = getModule()) {
797 p << ' ' << moduleAttrName << '(';
798 p.printSymbolName(*module);
799 p << ')';
800 }
801 // Print optional function attribute.
802 StringRef functionAttrName = getFunctionAttrName();
803 if (auto function = getFunction()) {
804 p << ' ' << functionAttrName << '(';
805 p.printSymbolName(*function);
806 p << ')';
807 }
808
809 if (getCooperative())
810 p << " cooperative";
811
812 printAttributions(p, getWorkgroupKeyword(), getWorkgroupAttributionBBArgs());
813 printAttributions(p, getPrivateKeyword(), getPrivateAttributions());
814
815 p << ' ';
816
817 p.printRegion(getBody(), /*printEntryBlockArgs=*/false);
819 (*this)->getDiscardableAttrDictionary().getValue(), /*elidedAttrs=*/{
820 LaunchOp::getOperandSegmentSizeAttr(),
821 getWorkgroupAttributionsAttrName(), getCooperativeAttrName(),
822 moduleAttrName, functionAttrName});
823}
824
825// Parse the size assignment blocks for blocks and threads. These have the form
826// (%region_arg, %region_arg, %region_arg) in
827// (%region_arg = %operand, %region_arg = %operand, %region_arg = %operand)
828// where %region_arg are percent-identifiers for the region arguments to be
829// introduced further (SSA defs), and %operand are percent-identifiers for the
830// SSA value uses.
831static ParseResult
836 StringRef keyword) {
837 assert(indices.size() == 3 && "space for three indices expected");
840 /*allowResultNumber=*/false) ||
841 parser.parseKeyword("in") || parser.parseLParen())
842 return failure();
843
844 if (args.size() != 3) {
845 return parser.emitError(parser.getNameLoc())
846 << keyword << " expects 3 arguments, but got " << args.size();
847 }
848 std::move(args.begin(), args.end(), indices.begin());
849
850 for (int i = 0; i < 3; ++i) {
851 if (i != 0 && parser.parseComma())
852 return failure();
853 if (parser.parseOperand(regionSizes[i], /*allowResultNumber=*/false) ||
854 parser.parseEqual() || parser.parseOperand(sizes[i]))
855 return failure();
856 }
857
858 return parser.parseRParen();
859}
860
861/// Parses a Launch operation.
862/// operation ::= `gpu.launch` (`<` ssa-use `:` type `>`)?
863/// (`async` `[` ssa-id-list `]`)?
864/// (`clusters` `(` ssa-id-list `)` `in` ssa-reassignment)?
865/// `blocks` `(` ssa-id-list `)` `in` ssa-reassignment
866/// `threads` `(` ssa-id-list `)` `in` ssa-reassignment
867/// (`dynamic_shared_memory_size` ssa-use)?
868/// (`module(` symbol-ref-id `)`)?
869/// (`function(` symbol-ref-id `)`)?
870/// memory-attribution
871/// region attr-dict?
872/// ssa-reassignment ::= `(` ssa-id `=` ssa-use (`,` ssa-id `=` ssa-use)* `)`
873ParseResult LaunchOp::parse(OpAsmParser &parser, OperationState &result) {
874 // Sizes of the grid and block.
876 sizes(LaunchOp::kNumConfigOperands);
877
878 // Region arguments to be created.
880 LaunchOp::kNumConfigRegionAttributes);
881
882 // Parse optional asyncObject: < value : type >
883 OpAsmParser::UnresolvedOperand asyncObjectOperand;
884 Type asyncObjectType;
885 bool hasAsyncObject = false;
886 if (succeeded(parser.parseOptionalLess())) {
887 hasAsyncObject = true;
888 if (parser.parseOperand(asyncObjectOperand) || parser.parseColon() ||
889 parser.parseType(asyncObjectType) || parser.parseGreater())
890 return failure();
891 }
892
893 // Parse optional async dependencies.
895 Type asyncTokenType;
896 if (failed(
897 parseAsyncDependencies(parser, asyncTokenType, asyncDependencies)) ||
898 parser.resolveOperands(asyncDependencies, asyncTokenType,
899 result.operands))
900 return failure();
901 if (parser.getNumResults() > 0) {
902 if (!asyncTokenType)
903 return parser.emitError(
904 parser.getNameLoc(),
905 "gpu.launch requires 'async' keyword to return a value");
906 result.types.push_back(asyncTokenType);
907 }
908
909 bool hasCluster = false;
910 if (succeeded(parser.parseOptionalKeyword(LaunchOp::getClustersKeyword()))) {
911 hasCluster = true;
912 sizes.resize(9);
913 regionArgs.resize(18);
914 }
916 MutableArrayRef<OpAsmParser::UnresolvedOperand> regionArgsRef(regionArgs);
917
918 // Last three segment assigns the cluster size. In the region argument
919 // list, this is last 6 arguments.
920 if (hasCluster) {
922 parser, sizesRef.drop_front(6), regionArgsRef.slice(15, 3),
923 regionArgsRef.slice(12, 3), LaunchOp::getClustersKeyword()))
924 return failure();
925 }
926 // Parse the size assignment segments: the first segment assigns grid sizes
927 // and defines values for block identifiers; the second segment assigns block
928 // sizes and defines values for thread identifiers. In the region argument
929 // list, identifiers precede sizes, and block-related values precede
930 // thread-related values.
931 if (parser.parseKeyword(LaunchOp::getBlocksKeyword()) ||
932 parseSizeAssignment(parser, sizesRef.take_front(3),
933 regionArgsRef.slice(6, 3), regionArgsRef.slice(0, 3),
934 LaunchOp::getBlocksKeyword()) ||
935 parser.parseKeyword(LaunchOp::getThreadsKeyword()) ||
936 parseSizeAssignment(parser, sizesRef.drop_front(3),
937 regionArgsRef.slice(9, 3), regionArgsRef.slice(3, 3),
938 LaunchOp::getThreadsKeyword()) ||
939 parser.resolveOperands(sizes, parser.getBuilder().getIndexType(),
940 result.operands))
941 return failure();
942
943 OpAsmParser::UnresolvedOperand dynamicSharedMemorySize;
944 bool hasDynamicSharedMemorySize = false;
945 if (!parser.parseOptionalKeyword(
946 LaunchOp::getDynamicSharedMemorySizeKeyword())) {
947 hasDynamicSharedMemorySize = true;
948 if (parser.parseOperand(dynamicSharedMemorySize) ||
949 parser.resolveOperand(dynamicSharedMemorySize,
950 parser.getBuilder().getI32Type(),
951 result.operands))
952 return failure();
953 }
954
955 // Resolve the asyncObject operand
956 if (hasAsyncObject && parser.resolveOperand(asyncObjectOperand,
957 asyncObjectType, result.operands))
958 return failure();
959
960 // Parse optional module attribute.
961 StringRef moduleAttrName = getModuleAttrName(result.name);
962 if (succeeded(parser.parseOptionalKeyword(moduleAttrName))) {
963 FlatSymbolRefAttr moduleSymbol;
964 if (parser.parseLParen() ||
965 parser.parseAttribute(moduleSymbol, Type(), moduleAttrName,
966 result.attributes) ||
967 parser.parseRParen())
968 return failure();
969 }
970 // Parse optional function attribute.
971 StringRef functionAttrName = getFunctionAttrName(result.name);
972 if (succeeded(parser.parseOptionalKeyword(functionAttrName))) {
973 FlatSymbolRefAttr funcSymbol;
974 if (parser.parseLParen() ||
975 parser.parseAttribute(funcSymbol, Type(), functionAttrName,
976 result.attributes) ||
977 parser.parseRParen())
978 return failure();
979 }
980
981 // Parse optional cooperative keyword.
982 if (succeeded(parser.parseOptionalKeyword("cooperative")))
983 result.addAttribute("cooperative", parser.getBuilder().getUnitAttr());
984
985 // Create the region arguments: fixed launch-config args (`index`), then
986 // workgroup / private attribution args. The workgroup count is stored in the
987 // inherent `workgroup_attributions` attribute when non-zero.
988 Type index = parser.getBuilder().getIndexType();
990 LaunchOp::kNumConfigRegionAttributes + 6, index);
991
993 for (auto ssaValueAndType : llvm::zip(regionArgs, dataTypes)) {
995 arg.ssaName = std::get<0>(ssaValueAndType);
996 arg.type = std::get<1>(ssaValueAndType);
997 regionArguments.push_back(arg);
998 }
999
1000 Builder &builder = parser.getBuilder();
1001 // Parse workgroup memory attributions.
1002 if (failed(parseAttributions(parser, LaunchOp::getWorkgroupKeyword(),
1003 regionArguments)))
1004 return failure();
1005
1006 // Store the number of operands we just parsed as the number of workgroup
1007 // memory attributions.
1008 unsigned numWorkgroupAttrs = regionArguments.size() -
1009 LaunchOp::kNumConfigRegionAttributes -
1010 (hasCluster ? 6 : 0);
1011 if (numWorkgroupAttrs != 0)
1012 result.addAttribute(LaunchOp::getWorkgroupAttributionsAttrName(result.name),
1013 builder.getI64IntegerAttr(numWorkgroupAttrs));
1014
1015 // Parse private memory attributions.
1016 if (failed(parseAttributions(parser, LaunchOp::getPrivateKeyword(),
1017 regionArguments)))
1018 return failure();
1019
1020 // Introduce the body region and parse it. The region has
1021 // kNumConfigRegionAttributes arguments that correspond to
1022 // block/thread identifiers and grid/block sizes, all having `index` type.
1023 Region *body = result.addRegion();
1024 if (parser.parseRegion(*body, regionArguments) ||
1025 parser.parseOptionalAttrDict(result.attributes))
1026 return failure();
1027
1028 SmallVector<int32_t, 12> segmentSizes(12, 1);
1029 segmentSizes.front() = asyncDependencies.size();
1030
1031 if (!hasCluster) {
1032 segmentSizes[7] = 0;
1033 segmentSizes[8] = 0;
1034 segmentSizes[9] = 0;
1035 }
1036 segmentSizes[10] = hasDynamicSharedMemorySize ? 1 : 0;
1037 segmentSizes[11] = hasAsyncObject ? 1 : 0;
1038 result.addAttribute(LaunchOp::getOperandSegmentSizeAttr(),
1039 parser.getBuilder().getDenseI32ArrayAttr(segmentSizes));
1040 return success();
1041}
1042
1043/// Simplify the gpu.launch when the range of a thread or block ID is
1044/// trivially known to be one.
1045struct FoldLaunchArguments : public OpRewritePattern<LaunchOp> {
1046 using OpRewritePattern<LaunchOp>::OpRewritePattern;
1047 LogicalResult matchAndRewrite(LaunchOp op,
1048 PatternRewriter &rewriter) const override {
1049 // If the range implies a single value for `id`, replace `id`'s uses by
1050 // zero.
1051 Value zero;
1052 bool simplified = false;
1053 auto constPropIdUses = [&](Value id, Value size) {
1054 // Check if size is trivially one.
1055 if (!matchPattern(size, m_One()))
1056 return;
1057 if (id.getUses().empty())
1058 return;
1059 if (!simplified) {
1060 // Create a zero value the first time.
1061 OpBuilder::InsertionGuard guard(rewriter);
1062 rewriter.setInsertionPointToStart(&op.getBody().front());
1063 zero =
1064 arith::ConstantIndexOp::create(rewriter, op.getLoc(), /*value=*/0);
1065 }
1066 rewriter.replaceAllUsesWith(id, zero);
1067 simplified = true;
1068 };
1069 constPropIdUses(op.getBlockIds().x, op.getGridSizeX());
1070 constPropIdUses(op.getBlockIds().y, op.getGridSizeY());
1071 constPropIdUses(op.getBlockIds().z, op.getGridSizeZ());
1072 constPropIdUses(op.getThreadIds().x, op.getBlockSizeX());
1073 constPropIdUses(op.getThreadIds().y, op.getBlockSizeY());
1074 constPropIdUses(op.getThreadIds().z, op.getBlockSizeZ());
1075
1076 return success(simplified);
1077 }
1078};
1079
1080void LaunchOp::getCanonicalizationPatterns(RewritePatternSet &rewrites,
1081 MLIRContext *context) {
1082 rewrites.add<FoldLaunchArguments>(context);
1083}
1084
1085/// Adds a new block argument that corresponds to buffers located in
1086/// workgroup memory.
1087BlockArgument LaunchOp::addWorkgroupAttribution(Type type, Location loc) {
1088 int64_t cur = getWorkgroupAttributions().value_or(0);
1089 setWorkgroupAttributions(std::optional<int64_t>(cur + 1));
1090 return getBody().insertArgument(
1091 getNumConfigRegionAttributes() + static_cast<unsigned>(cur), type, loc);
1092}
1093
1094/// Adds a new block argument that corresponds to buffers located in
1095/// private memory.
1096BlockArgument LaunchOp::addPrivateAttribution(Type type, Location loc) {
1097 // Buffers on the private memory always come after buffers on the workgroup
1098 // memory.
1099 return getBody().addArgument(type, loc);
1100}
1101
1102//===----------------------------------------------------------------------===//
1103// LaunchFuncOp
1104//===----------------------------------------------------------------------===//
1105
1106void LaunchFuncOp::build(OpBuilder &builder, OperationState &result,
1107 SymbolRefAttr kernelSymbol, KernelDim3 gridSize,
1108 KernelDim3 getBlockSize, Value dynamicSharedMemorySize,
1109 ValueRange kernelOperands, Type asyncTokenType,
1110 ValueRange asyncDependencies, Value asyncObject,
1111 std::optional<KernelDim3> clusterSize) {
1112 assert(kernelSymbol.getNestedReferences().size() == 1 &&
1113 "expected a symbol reference with a single nested reference");
1114 result.addOperands(asyncDependencies);
1115 if (asyncTokenType)
1116 result.types.push_back(builder.getType<AsyncTokenType>());
1117
1118 // Add grid and block sizes as op operands, followed by the data operands.
1119 result.addOperands({gridSize.x, gridSize.y, gridSize.z, getBlockSize.x,
1121 if (clusterSize.has_value())
1122 result.addOperands({clusterSize->x, clusterSize->y, clusterSize->z});
1123 if (dynamicSharedMemorySize)
1124 result.addOperands(dynamicSharedMemorySize);
1125 result.addOperands(kernelOperands);
1126 if (asyncObject)
1127 result.addOperands(asyncObject);
1128
1129 Properties &prop = result.getOrAddProperties<Properties>();
1130 prop.kernel = kernelSymbol;
1131 size_t segmentSizesLen = std::size(prop.operandSegmentSizes);
1132 // Initialize the segment sizes to 1.
1133 llvm::fill(prop.operandSegmentSizes, 1);
1134 prop.operandSegmentSizes[0] = asyncDependencies.size();
1135 if (!clusterSize.has_value()) {
1136 prop.operandSegmentSizes[segmentSizesLen - 4] = 0;
1137 prop.operandSegmentSizes[segmentSizesLen - 5] = 0;
1138 prop.operandSegmentSizes[segmentSizesLen - 6] = 0;
1139 }
1140 prop.operandSegmentSizes[segmentSizesLen - 3] =
1141 dynamicSharedMemorySize ? 1 : 0;
1142 prop.operandSegmentSizes[segmentSizesLen - 2] =
1143 static_cast<int32_t>(kernelOperands.size());
1144 prop.operandSegmentSizes[segmentSizesLen - 1] = asyncObject ? 1 : 0;
1145}
1146
1147void LaunchFuncOp::build(OpBuilder &builder, OperationState &result,
1148 GPUFuncOp kernelFunc, KernelDim3 gridSize,
1149 KernelDim3 getBlockSize, Value dynamicSharedMemorySize,
1150 ValueRange kernelOperands, Type asyncTokenType,
1151 ValueRange asyncDependencies, Value asyncObject,
1152 std::optional<KernelDim3> clusterSize) {
1153 auto kernelModule = kernelFunc->getParentOfType<GPUModuleOp>();
1154 auto kernelSymbol =
1155 SymbolRefAttr::get(kernelModule.getNameAttr(),
1156 {SymbolRefAttr::get(kernelFunc.getNameAttr())});
1157 build(builder, result, kernelSymbol, gridSize, getBlockSize,
1158 dynamicSharedMemorySize, kernelOperands, asyncTokenType,
1159 asyncDependencies, asyncObject, clusterSize);
1160}
1161
1162StringAttr LaunchFuncOp::getKernelModuleName() {
1163 return getKernel().getRootReference();
1164}
1165
1166StringAttr LaunchFuncOp::getKernelName() {
1167 return getKernel().getLeafReference();
1168}
1169
1170unsigned LaunchFuncOp::getNumKernelOperands() {
1171 return getKernelOperands().size();
1172}
1173
1174Value LaunchFuncOp::getKernelOperand(unsigned i) {
1175 return getKernelOperands()[i];
1176}
1177
1178KernelDim3 LaunchFuncOp::getGridSizeOperandValues() {
1179 auto operands = getOperands().drop_front(getAsyncDependencies().size());
1180 return KernelDim3{operands[0], operands[1], operands[2]};
1181}
1182
1183KernelDim3 LaunchFuncOp::getBlockSizeOperandValues() {
1184 auto operands = getOperands().drop_front(getAsyncDependencies().size());
1185 return KernelDim3{operands[3], operands[4], operands[5]};
1186}
1187
1188KernelDim3 LaunchFuncOp::getClusterSizeOperandValues() {
1189 assert(hasClusterSize() &&
1190 "cluster size is not set, check hasClusterSize() first");
1191 auto operands = getOperands().drop_front(getAsyncDependencies().size());
1192 return KernelDim3{operands[6], operands[7], operands[8]};
1193}
1194
1195LogicalResult LaunchFuncOp::verify() {
1196 if (verifyLaunchAsyncModel(*this).failed())
1197 return failure();
1198
1199 auto module = (*this)->getParentOfType<ModuleOp>();
1200 if (!module)
1201 return emitOpError("expected to belong to a module");
1202
1203 if (!module->getDiscardableAttrOfType<UnitAttr>(
1204 GPUDialect::getContainerModuleAttrName()))
1205 return emitOpError("expected the closest surrounding module to have the '" +
1206 GPUDialect::getContainerModuleAttrName() +
1207 "' attribute");
1208
1209 if (hasClusterSize()) {
1210 if (getClusterSizeY().getType() != getClusterSizeX().getType() ||
1211 getClusterSizeZ().getType() != getClusterSizeX().getType())
1212 return emitOpError()
1213 << "expects types of the cluster dimensions must be the same";
1214 }
1215
1216 return success();
1217}
1218
1219LogicalResult
1220LaunchFuncOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
1221 LaunchFuncOp launchOp = *this;
1223 // GPU modules cannot be nested within each other, escape to resolve the name.
1224 if (isa<GPUModuleOp>(table))
1226
1227 // Ignore launches that are nested more or less deep than functions in the
1228 // module we are currently checking.
1229 if (!launchOp->getParentOp() ||
1230 launchOp->getParentOp()->getParentOp() != table)
1231 return success();
1232
1233 // Ignore launch ops with missing attributes here. The errors will be
1234 // reported by the verifiers of those ops.
1235 if (!launchOp.getKernelAttr())
1236 return success();
1237
1238 // Check that `launch_func` refers to a well-formed GPU kernel container.
1239 StringAttr kernelContainerName = launchOp.getKernelModuleName();
1240 Operation *kernelContainer =
1241 symbolTable.lookupNearestSymbolFrom(table, kernelContainerName);
1242 if (!kernelContainer)
1243 return launchOp.emitOpError()
1244 << "kernel container '" << kernelContainerName.getValue()
1245 << "' is undefined";
1246
1247 // If the container is a GPU binary op return success.
1248 if (isa<BinaryOp>(kernelContainer))
1249 return success();
1250
1251 auto kernelModule = dyn_cast<GPUModuleOp>(kernelContainer);
1252 if (!kernelModule)
1253 return launchOp.emitOpError()
1254 << "kernel module '" << kernelContainerName.getValue()
1255 << "' is undefined";
1256
1257 // Check that `launch_func` refers to a well-formed kernel function.
1258 Operation *kernelFunc = symbolTable.lookupNearestSymbolFrom(
1259 kernelModule, launchOp.getKernelName());
1260 if (!kernelFunc)
1261 return launchOp.emitOpError("kernel function '")
1262 << launchOp.getKernel() << "' is undefined";
1263 auto kernelConvertedFunction = dyn_cast<FunctionOpInterface>(kernelFunc);
1264 if (!kernelConvertedFunction) {
1265 InFlightDiagnostic diag = launchOp.emitOpError()
1266 << "referenced kernel '" << launchOp.getKernel()
1267 << "' is not a function";
1268 diag.attachNote(kernelFunc->getLoc()) << "see the kernel definition here";
1269 return diag;
1270 }
1271
1272 if (!GPUDialect::isKernel(kernelFunc))
1273 return launchOp.emitOpError("kernel function is missing the '")
1274 << GPUDialect::getKernelFuncAttrName() << "' attribute";
1275
1276 // TODO: If the kernel isn't a GPU function (which happens during separate
1277 // compilation), do not check type correspondence as it would require the
1278 // verifier to be aware of the type conversion.
1279 auto kernelGPUFunction = dyn_cast<gpu::GPUFuncOp>(kernelFunc);
1280 if (!kernelGPUFunction)
1281 return success();
1282
1283 unsigned actualNumArguments = launchOp.getNumKernelOperands();
1284 unsigned expectedNumArguments = kernelGPUFunction.getNumArguments();
1285 if (expectedNumArguments != actualNumArguments)
1286 return launchOp.emitOpError("got ")
1287 << actualNumArguments << " kernel operands but expected "
1288 << expectedNumArguments;
1289
1290 FunctionType functionType = kernelGPUFunction.getFunctionType();
1291 for (unsigned i = 0; i < expectedNumArguments; ++i) {
1292 if (launchOp.getKernelOperand(i).getType() != functionType.getInput(i)) {
1293 return launchOp.emitOpError("type of function argument ")
1294 << i << " does not match";
1295 }
1296 }
1297
1298 return success();
1299}
1300
1301static ParseResult
1303 std::optional<OpAsmParser::UnresolvedOperand> clusterValue,
1304 Type &clusterXTy, Type &clusterYTy, Type &clusterZTy) {
1305 if (succeeded(parser.parseOptionalColon())) {
1306 if (parser.parseType(dimTy))
1307 return failure();
1308 } else {
1309 dimTy = IndexType::get(parser.getContext());
1310 }
1311 if (clusterValue.has_value()) {
1312 clusterXTy = clusterYTy = clusterZTy = dimTy;
1313 }
1314 return success();
1315}
1316
1317static void printLaunchDimType(OpAsmPrinter &printer, Operation *op, Type dimTy,
1318 Value clusterValue, Type clusterXTy,
1319 Type clusterYTy, Type clusterZTy) {
1320 if (!dimTy.isIndex())
1321 printer << ": " << dimTy;
1322}
1323
1324static ParseResult parseLaunchFuncOperands(
1325 OpAsmParser &parser,
1327 SmallVectorImpl<Type> &argTypes) {
1328 if (parser.parseOptionalKeyword("args"))
1329 return success();
1330
1331 auto parseElement = [&]() -> ParseResult {
1332 return failure(parser.parseOperand(argNames.emplace_back()) ||
1333 parser.parseColonType(argTypes.emplace_back()));
1334 };
1335
1337 parseElement, " in argument list");
1338}
1339
1341 OperandRange operands, TypeRange types) {
1342 if (operands.empty())
1343 return;
1344 printer << "args(";
1345 llvm::interleaveComma(llvm::zip_equal(operands, types), printer,
1346 [&](const auto &pair) {
1347 auto [operand, type] = pair;
1348 printer << operand << " : " << type;
1349 });
1350 printer << ")";
1351}
1352
1353//===----------------------------------------------------------------------===//
1354// ShuffleOp
1355//===----------------------------------------------------------------------===//
1356
1357void ShuffleOp::build(OpBuilder &builder, OperationState &result, Value value,
1358 int32_t offset, int32_t width, ShuffleMode mode) {
1359 build(builder, result, value,
1360 arith::ConstantOp::create(builder, result.location,
1361 builder.getI32IntegerAttr(offset)),
1362 arith::ConstantOp::create(builder, result.location,
1363 builder.getI32IntegerAttr(width)),
1364 mode);
1365}
1366
1367//===----------------------------------------------------------------------===//
1368// RotateOp
1369//===----------------------------------------------------------------------===//
1370
1371LogicalResult RotateOp::verify() {
1372 uint32_t offset = getOffset();
1373 uint32_t width = getWidth();
1374
1375 if (offset >= width) {
1376 return emitOpError() << "offset must be in the range [0, " << width << ")";
1377 }
1378
1379 return success();
1380}
1381
1382//===----------------------------------------------------------------------===//
1383// BarrierOp
1384//===----------------------------------------------------------------------===//
1385
1386LogicalResult BarrierOp::verify() {
1387 BarrierScope scope = getScope();
1388
1389 if (getNamedBarrier() && scope != BarrierScope::Workgroup)
1390 return emitOpError("named barriers require workgroup scope");
1391
1392 return success();
1393}
1394
1395/// Remove gpu.barrier after gpu.barrier, the threads are already synchronized!
1396static LogicalResult eraseRedundantGpuBarrierOps(BarrierOp op,
1397 PatternRewriter &rewriter) {
1398 auto nextOp = dyn_cast_or_null<BarrierOp>(op->getNextNode());
1399 if (!nextOp)
1400 return failure();
1401
1402 // Cannot merge barriers of different scopes.
1403 if (op.getScope() != nextOp.getScope())
1404 return failure();
1405
1406 // Cannot merge named barriers unless both refer to the same handle.
1407 if (op.getNamedBarrier() != nextOp.getNamedBarrier())
1408 return failure();
1409
1410 std::optional<ArrayAttr> thisMemfence = op.getAddressSpaces();
1411 std::optional<ArrayAttr> nextMemfence = nextOp.getAddressSpaces();
1412
1413 if (thisMemfence) {
1414 rewriter.modifyOpInPlace(op, [&]() {
1415 if (!nextMemfence) {
1416 op.removeAddressSpacesAttr();
1417 return;
1418 }
1419 // Fast path - merge where the two barriers fence the same spaces.
1420 if (*thisMemfence == *nextMemfence) {
1421 return;
1422 }
1423
1424 llvm::SmallSetVector<Attribute, 4> mergedSpaces;
1425 for (Attribute attr : *thisMemfence)
1426 mergedSpaces.insert(attr);
1427 for (Attribute attr : *nextMemfence)
1428 mergedSpaces.insert(attr);
1429 op.setAddressSpacesAttr(rewriter.getArrayAttr(mergedSpaces.takeVector()));
1430 });
1431 }
1432
1433 rewriter.eraseOp(nextOp);
1434 return success();
1435}
1436
1437void BarrierOp::getCanonicalizationPatterns(RewritePatternSet &results,
1438 MLIRContext *context) {
1440}
1441
1442void BarrierOp::build(mlir::OpBuilder &odsBuilder,
1443 mlir::OperationState &odsState,
1444 std::optional<AddressSpace> addressSpace) {
1445 ArrayAttr addressSpacesAttr;
1446 if (addressSpace)
1447 addressSpacesAttr = odsBuilder.getArrayAttr(
1448 AddressSpaceAttr::get(odsBuilder.getContext(), addressSpace.value()));
1449 build(
1450 odsBuilder, odsState, addressSpacesAttr, /*named_barrier=*/Value{},
1451 BarrierScopeAttr::get(odsBuilder.getContext(), BarrierScope::Workgroup));
1452}
1453
1454/// Builds a barrier that causes memory operations affecting `memrefToFence` to
1455/// be completed after the barrier is concluded. Currently, this means setting
1456/// the fenced address spaces to those of the given memref if it is a gpu
1457/// address space.
1458void BarrierOp::build(OpBuilder &builder, OperationState &odsState,
1459 Value memrefToFence) {
1460 std::optional<AddressSpace> addrSpaceToFence;
1461 if (auto memrefType = dyn_cast<BaseMemRefType>(memrefToFence.getType()))
1462 if (auto addrSpaceAttr = dyn_cast_if_present<gpu::AddressSpaceAttr>(
1463 memrefType.getMemorySpace()))
1464 addrSpaceToFence = addrSpaceAttr.getValue();
1465 return build(builder, odsState, addrSpaceToFence);
1466}
1467
1468//===----------------------------------------------------------------------===//
1469// GPUFuncOp
1470//===----------------------------------------------------------------------===//
1471
1472/// Adds a new block argument that corresponds to buffers located in
1473/// workgroup memory.
1474BlockArgument GPUFuncOp::addWorkgroupAttribution(Type type, Location loc) {
1475 int64_t cur = getWorkgroupAttributions().value_or(0);
1476 setWorkgroupAttributions(std::optional<int64_t>(cur + 1));
1477 return getBody().insertArgument(
1478 getFunctionType().getNumInputs() + static_cast<unsigned>(cur), type, loc);
1479}
1480
1481/// Adds a new block argument that corresponds to buffers located in
1482/// private memory.
1483BlockArgument GPUFuncOp::addPrivateAttribution(Type type, Location loc) {
1484 // Buffers on the private memory always come after buffers on the workgroup
1485 // memory.
1486 return getBody().addArgument(type, loc);
1487}
1488
1489void GPUFuncOp::build(OpBuilder &builder, OperationState &result,
1490 StringRef name, FunctionType type,
1491 TypeRange workgroupAttributions,
1492 TypeRange privateAttributions,
1494 OpBuilder::InsertionGuard g(builder);
1495
1496 result.getOrAddProperties<Properties>().sym_name =
1497 builder.getStringAttr(name);
1498 result.addAttribute(getFunctionTypeAttrName(result.name),
1499 TypeAttr::get(type));
1500 result.addAttribute(getWorkgroupAttributionsAttrName(result.name),
1501 builder.getI64IntegerAttr(workgroupAttributions.size()));
1502 result.addAttributes(attrs);
1503 Region *body = result.addRegion();
1504 Block *entryBlock = builder.createBlock(body);
1505
1506 // TODO: Allow passing in proper locations here.
1507 for (Type argTy : type.getInputs())
1508 entryBlock->addArgument(argTy, result.location);
1509 for (Type argTy : workgroupAttributions)
1510 entryBlock->addArgument(argTy, result.location);
1511 for (Type argTy : privateAttributions)
1512 entryBlock->addArgument(argTy, result.location);
1513}
1514
1515/// Parses a GPU function memory attribution.
1516///
1517/// memory-attribution ::= (`workgroup` `(` ssa-id-and-type-list `)`)?
1518/// (`private` `(` ssa-id-and-type-list `)`)?
1519///
1520/// Note that this function parses only one of the two similar parts, with the
1521/// keyword provided as argument.
1522static ParseResult
1523parseAttributions(OpAsmParser &parser, StringRef keyword,
1525 Attribute &attributionAttrs) {
1526 // If we could not parse the keyword, just assume empty list and succeed.
1527 if (failed(parser.parseOptionalKeyword(keyword)))
1528 return success();
1529
1530 size_t existingArgs = args.size();
1531 ParseResult result =
1533 /*allowType=*/true, /*allowAttrs=*/true);
1534 if (failed(result))
1535 return result;
1536
1537 bool hadAttrs = llvm::any_of(ArrayRef(args).drop_front(existingArgs),
1538 [](const OpAsmParser::Argument &arg) -> bool {
1539 return arg.attrs && !arg.attrs.empty();
1540 });
1541 if (!hadAttrs) {
1542 attributionAttrs = nullptr;
1543 return result;
1544 }
1545
1546 Builder &builder = parser.getBuilder();
1547 SmallVector<Attribute> attributionAttrsVec;
1548 for (const auto &argument : ArrayRef(args).drop_front(existingArgs)) {
1549 if (!argument.attrs)
1550 attributionAttrsVec.push_back(builder.getDictionaryAttr({}));
1551 else
1552 attributionAttrsVec.push_back(argument.attrs);
1553 }
1554 attributionAttrs = builder.getArrayAttr(attributionAttrsVec);
1555 return result;
1556}
1557
1558/// Parses a GPU function.
1559///
1560/// <operation> ::= `gpu.func` symbol-ref-id `(` argument-list `)`
1561/// (`->` function-result-list)? memory-attribution `kernel`?
1562/// function-attributes? region
1563ParseResult GPUFuncOp::parse(OpAsmParser &parser, OperationState &result) {
1565 SmallVector<DictionaryAttr> resultAttrs;
1566 SmallVector<Type> resultTypes;
1567 bool isVariadic;
1568
1569 // Parse the function name.
1570 StringAttr nameAttr;
1571 if (parser.parseSymbolName(nameAttr, getSymNameAttrName(result.name),
1572 result.attributes))
1573 return failure();
1574
1575 auto signatureLocation = parser.getCurrentLocation();
1577 parser, /*allowVariadic=*/false, entryArgs, isVariadic, resultTypes,
1578 resultAttrs)))
1579 return failure();
1580
1581 if (!entryArgs.empty() && entryArgs[0].ssaName.name.empty())
1582 return parser.emitError(signatureLocation)
1583 << "gpu.func requires named arguments";
1584
1585 // Construct the function type. More types will be added to the region, but
1586 // not to the function type.
1587 Builder &builder = parser.getBuilder();
1588
1589 SmallVector<Type> argTypes;
1590 for (auto &arg : entryArgs)
1591 argTypes.push_back(arg.type);
1592 auto type = builder.getFunctionType(argTypes, resultTypes);
1593 result.addAttribute(getFunctionTypeAttrName(result.name),
1594 TypeAttr::get(type));
1595
1597 builder, result, entryArgs, resultAttrs, getArgAttrsAttrName(result.name),
1598 getResAttrsAttrName(result.name));
1599
1600 Attribute workgroupAttributionAttrs;
1601 // Parse workgroup memory attributions.
1602 if (failed(parseAttributions(parser, GPUFuncOp::getWorkgroupKeyword(),
1603 entryArgs, workgroupAttributionAttrs)))
1604 return failure();
1605
1606 // Store the number of operands we just parsed as the number of workgroup
1607 // memory attributions.
1608 unsigned numWorkgroupAttrs = entryArgs.size() - type.getNumInputs();
1609 if (numWorkgroupAttrs != 0)
1610 result.addAttribute(
1611 GPUFuncOp::getWorkgroupAttributionsAttrName(result.name),
1612 builder.getI64IntegerAttr(numWorkgroupAttrs));
1613 if (workgroupAttributionAttrs)
1614 result.addAttribute(GPUFuncOp::getWorkgroupAttribAttrsAttrName(result.name),
1615 workgroupAttributionAttrs);
1616
1617 Attribute privateAttributionAttrs;
1618 // Parse private memory attributions.
1619 if (failed(parseAttributions(parser, GPUFuncOp::getPrivateKeyword(),
1620 entryArgs, privateAttributionAttrs)))
1621 return failure();
1622 if (privateAttributionAttrs)
1623 result.addAttribute(GPUFuncOp::getPrivateAttribAttrsAttrName(result.name),
1624 privateAttributionAttrs);
1625
1626 // Parse the kernel attribute if present.
1627 if (succeeded(parser.parseOptionalKeyword(GPUFuncOp::getKernelKeyword())))
1628 result.addAttribute(GPUFuncOp::getKernelAttrName(result.name),
1629 builder.getUnitAttr());
1630
1631 // Parse attributes.
1632 if (failed(parser.parseOptionalAttrDictWithKeyword(result.attributes)))
1633 return failure();
1634
1635 // Parse the region. If no argument names were provided, take all names
1636 // (including those of attributions) from the entry block.
1637 auto *body = result.addRegion();
1638 return parser.parseRegion(*body, entryArgs);
1639}
1640
1641void GPUFuncOp::print(OpAsmPrinter &p) {
1642 p << ' ';
1643 p.printSymbolName(getName());
1644
1645 FunctionType type = getFunctionType();
1646 function_interface_impl::printFunctionSignature(p, *this, type.getInputs(),
1647 /*isVariadic=*/false,
1648 type.getResults());
1649
1650 printAttributions(p, getWorkgroupKeyword(), getWorkgroupAttributionBBArgs(),
1651 getWorkgroupAttribAttrs().value_or(nullptr));
1652 printAttributions(p, getPrivateKeyword(), getPrivateAttributions(),
1653 getPrivateAttribAttrs().value_or(nullptr));
1654 if (isKernel())
1655 p << ' ' << getKernelKeyword();
1656
1658 p, *this,
1659 {getWorkgroupAttributionsAttrName(), getKernelAttrName(),
1660 GPUDialect::getKernelFuncAttrName(), getFunctionTypeAttrName(),
1661 getArgAttrsAttrName(), getResAttrsAttrName(),
1662 getWorkgroupAttribAttrsAttrName(), getPrivateAttribAttrsAttrName()});
1663 p << ' ';
1664 p.printRegion(getBody(), /*printEntryBlockArgs=*/false);
1665}
1666
1667static DictionaryAttr getAttributionAttrs(GPUFuncOp op, unsigned index,
1668 StringAttr attrName) {
1669 ArrayAttr allAttrs = attrName == op.getWorkgroupAttribAttrsAttrName()
1670 ? op.getWorkgroupAttribAttrsAttr()
1671 : op.getPrivateAttribAttrsAttr();
1672 if (!allAttrs || index >= allAttrs.size())
1673 return DictionaryAttr();
1674 return llvm::cast<DictionaryAttr>(allAttrs[index]);
1675}
1676
1677DictionaryAttr GPUFuncOp::getworkgroupAttributionAttrs(unsigned index) {
1678 return getAttributionAttrs(*this, index, getWorkgroupAttribAttrsAttrName());
1679}
1680
1681DictionaryAttr GPUFuncOp::getPrivateAttributionAttrs(unsigned index) {
1682 return getAttributionAttrs(*this, index, getPrivateAttribAttrsAttrName());
1683}
1684
1685static void setAttributionAttrs(GPUFuncOp op, unsigned index,
1686 DictionaryAttr value, StringAttr attrName) {
1687 MLIRContext *ctx = op.getContext();
1688 ArrayAttr allAttrs = attrName == op.getWorkgroupAttribAttrsAttrName()
1689 ? op.getWorkgroupAttribAttrsAttr()
1690 : op.getPrivateAttribAttrsAttr();
1691 SmallVector<Attribute> elements;
1692 if (allAttrs)
1693 elements.append(allAttrs.begin(), allAttrs.end());
1694 while (elements.size() <= index)
1695 elements.push_back(DictionaryAttr::get(ctx));
1696 if (!value)
1697 elements[index] = DictionaryAttr::get(ctx);
1698 else
1699 elements[index] = value;
1700 ArrayAttr newValue = ArrayAttr::get(ctx, elements);
1701 if (attrName == op.getWorkgroupAttribAttrsAttrName())
1702 op.setWorkgroupAttribAttrsAttr(newValue);
1703 else
1704 op.setPrivateAttribAttrsAttr(newValue);
1705}
1706
1707void GPUFuncOp::setworkgroupAttributionAttrs(unsigned index,
1708 DictionaryAttr value) {
1709 setAttributionAttrs(*this, index, value, getWorkgroupAttribAttrsAttrName());
1710}
1711
1712void GPUFuncOp::setPrivateAttributionAttrs(unsigned int index,
1713 DictionaryAttr value) {
1714 setAttributionAttrs(*this, index, value, getPrivateAttribAttrsAttrName());
1715}
1716
1717static Attribute getAttributionAttr(GPUFuncOp op, unsigned index,
1718 StringAttr name, StringAttr attrsName) {
1719 DictionaryAttr dict = getAttributionAttrs(op, index, attrsName);
1720 if (!dict)
1721 return Attribute();
1722 return dict.get(name);
1723}
1724
1725Attribute GPUFuncOp::getWorkgroupAttributionAttr(unsigned index,
1726 StringAttr name) {
1727 assert(index < getNumWorkgroupAttributions() &&
1728 "index must map to a workgroup attribution");
1729 return getAttributionAttr(*this, index, name,
1730 getWorkgroupAttribAttrsAttrName());
1731}
1732
1733Attribute GPUFuncOp::getPrivateAttributionAttr(unsigned index,
1734 StringAttr name) {
1735 assert(index < getNumPrivateAttributions() &&
1736 "index must map to a private attribution");
1737 return getAttributionAttr(*this, index, name,
1738 getPrivateAttribAttrsAttrName());
1739}
1740
1741static void setAttributionAttr(GPUFuncOp op, unsigned index, StringAttr name,
1742 Attribute value, StringAttr attrsName) {
1743 MLIRContext *ctx = op.getContext();
1745 DictionaryAttr oldDict = getAttributionAttrs(op, index, attrsName);
1746 if (oldDict)
1747 elems.append(oldDict.getValue().begin(), oldDict.getValue().end());
1748
1749 bool found = false;
1750 bool mustSort = true;
1751 for (unsigned i = 0, e = elems.size(); i < e; ++i) {
1752 if (elems[i].getName() == name) {
1753 found = true;
1754 if (!value) {
1755 std::swap(elems[i], elems[elems.size() - 1]);
1756 elems.pop_back();
1757 } else {
1758 mustSort = false;
1759 elems[i] = NamedAttribute(elems[i].getName(), value);
1760 }
1761 break;
1762 }
1763 }
1764 if (!found) {
1765 if (!value)
1766 return;
1767 elems.emplace_back(name, value);
1768 }
1769 if (mustSort) {
1770 DictionaryAttr::sortInPlace(elems);
1771 }
1772 auto newDict = DictionaryAttr::getWithSorted(ctx, elems);
1773 setAttributionAttrs(op, index, newDict, attrsName);
1774}
1775
1776void GPUFuncOp::setWorkgroupAttributionAttr(unsigned index, StringAttr name,
1777 Attribute value) {
1778 assert(index < getNumWorkgroupAttributions() &&
1779 "index must map to a workgroup attribution");
1780 setAttributionAttr(*this, index, name, value,
1781 getWorkgroupAttribAttrsAttrName());
1782}
1783
1784void GPUFuncOp::setPrivateAttributionAttr(unsigned index, StringAttr name,
1785 Attribute value) {
1786 assert(index < getNumPrivateAttributions() &&
1787 "index must map to a private attribution");
1788 setAttributionAttr(*this, index, name, value,
1789 getPrivateAttribAttrsAttrName());
1790}
1791
1792LogicalResult GPUFuncOp::verifyType() {
1793 if (isKernel() && getFunctionType().getNumResults() != 0)
1794 return emitOpError() << "expected void return type for kernel function";
1795
1796 return success();
1797}
1798
1799/// Verifies the body of the function.
1800LogicalResult GPUFuncOp::verifyBody() {
1801 if (empty())
1802 return emitOpError() << "expected body with at least one block";
1803 unsigned numFuncArguments = getNumArguments();
1804 unsigned numWorkgroupAttributions = getNumWorkgroupAttributions();
1805 unsigned numBlockArguments = front().getNumArguments();
1806 if (numBlockArguments < numFuncArguments + numWorkgroupAttributions)
1807 return emitOpError() << "expected at least "
1808 << numFuncArguments + numWorkgroupAttributions
1809 << " arguments to body region";
1810
1811 ArrayRef<Type> funcArgTypes = getFunctionType().getInputs();
1812 for (unsigned i = 0; i < numFuncArguments; ++i) {
1813 Type blockArgType = front().getArgument(i).getType();
1814 if (funcArgTypes[i] != blockArgType)
1815 return emitOpError() << "expected body region argument #" << i
1816 << " to be of type " << funcArgTypes[i] << ", got "
1817 << blockArgType;
1818 }
1819
1820 if (failed(verifyAttributions(getOperation(), getWorkgroupAttributionBBArgs(),
1821 GPUDialect::getWorkgroupAddressSpace())) ||
1822 failed(verifyAttributions(getOperation(), getPrivateAttributions(),
1823 GPUDialect::getPrivateAddressSpace())))
1824 return failure();
1825
1826 return success();
1827}
1828
1829//===----------------------------------------------------------------------===//
1830// ReturnOp
1831//===----------------------------------------------------------------------===//
1832
1833LogicalResult gpu::ReturnOp::verify() {
1834 GPUFuncOp function = (*this)->getParentOfType<GPUFuncOp>();
1835
1836 FunctionType funType = function.getFunctionType();
1837
1838 if (funType.getNumResults() != getOperands().size())
1839 return emitOpError()
1840 .append("expected ", funType.getNumResults(), " result operands")
1841 .attachNote(function.getLoc())
1842 .append("return type declared here");
1843
1844 for (const auto &pair : llvm::enumerate(
1845 llvm::zip(function.getFunctionType().getResults(), getOperands()))) {
1846 auto [type, operand] = pair.value();
1847 if (type != operand.getType())
1848 return emitOpError() << "unexpected type `" << operand.getType()
1849 << "' for operand #" << pair.index();
1850 }
1851 return success();
1852}
1853
1854//===----------------------------------------------------------------------===//
1855// GPUModuleOp
1856//===----------------------------------------------------------------------===//
1857
1858LogicalResult GPUModuleOp::readProperties(DialectBytecodeReader &reader,
1859 OperationState &state) {
1860 auto &properties = state.getOrAddProperties<Properties>();
1861 // Preserve the pre-DLTI field order for unversioned bytecode.
1862 if (failed(reader.readOptionalAttribute(properties.offloadingHandler)) ||
1863 failed(reader.readAttribute(properties.sym_name)) ||
1864 failed(reader.readOptionalAttribute(properties.sym_visibility)) ||
1865 failed(reader.readOptionalAttribute(properties.targets)))
1866 return failure();
1867
1868 if (failed(reader.getDialectVersion<GPUDialect>()))
1869 return success();
1870 return reader.readOptionalAttribute(properties.dlti);
1871}
1872
1873void GPUModuleOp::writeProperties(DialectBytecodeWriter &writer) {
1874 Properties &properties = getProperties();
1875 writer.writeOptionalAttribute(properties.offloadingHandler);
1876 writer.writeAttribute(properties.sym_name);
1877 writer.writeOptionalAttribute(properties.sym_visibility);
1878 writer.writeOptionalAttribute(properties.targets);
1880 writer.writeOptionalAttribute(properties.dlti);
1881}
1882
1883void GPUModuleOp::build(OpBuilder &builder, OperationState &result,
1884 StringRef name, ArrayAttr targets,
1885 Attribute offloadingHandler) {
1886 result.addRegion()->emplaceBlock();
1887 Properties &props = result.getOrAddProperties<Properties>();
1888 if (targets)
1889 props.targets = targets;
1890 props.setSymName(builder.getStringAttr(name));
1891 props.offloadingHandler = offloadingHandler;
1892}
1893
1894void GPUModuleOp::build(OpBuilder &builder, OperationState &result,
1895 StringRef name, ArrayRef<Attribute> targets,
1896 Attribute offloadingHandler) {
1897 build(builder, result, name,
1898 targets.empty() ? ArrayAttr() : builder.getArrayAttr(targets),
1899 offloadingHandler);
1900}
1901
1902bool GPUModuleOp::hasTarget(Attribute target) {
1903 if (ArrayAttr targets = getTargetsAttr())
1904 return llvm::count(targets.getValue(), target);
1905 return false;
1906}
1907
1908void GPUModuleOp::setTargets(ArrayRef<TargetAttrInterface> targets) {
1909 ArrayAttr &targetsAttr = getProperties().targets;
1910 SmallVector<Attribute> targetsVector(targets);
1911 targetsAttr = ArrayAttr::get(getContext(), targetsVector);
1912}
1913
1914LogicalResult GPUModuleOp::verify() {
1915 auto targets = getTargetsAttr();
1916
1917 if (!targets)
1918 return success();
1919
1920 for (auto target : targets) {
1921 if (auto verifyTargetAttr =
1922 llvm::dyn_cast<TargetAttrVerifyInterface>(target)) {
1923 if (verifyTargetAttr.verifyTarget(getOperation()).failed())
1924 return failure();
1925 }
1926 }
1927 return success();
1928}
1929
1930//===----------------------------------------------------------------------===//
1931// GPUBinaryOp
1932//===----------------------------------------------------------------------===//
1933void BinaryOp::build(OpBuilder &builder, OperationState &result, StringRef name,
1934 Attribute offloadingHandler, ArrayAttr objects) {
1935 auto &properties = result.getOrAddProperties<Properties>();
1936 result.attributes.push_back(builder.getNamedAttr(
1937 getSymNameAttrName(result.name), builder.getStringAttr(name)));
1938 properties.objects = objects;
1939 if (offloadingHandler)
1940 properties.offloadingHandler = offloadingHandler;
1941 else
1942 properties.offloadingHandler = builder.getAttr<SelectObjectAttr>(nullptr);
1943}
1944
1945void BinaryOp::build(OpBuilder &builder, OperationState &result, StringRef name,
1946 Attribute offloadingHandler, ArrayRef<Attribute> objects) {
1947 build(builder, result, name, offloadingHandler,
1948 objects.empty() ? ArrayAttr() : builder.getArrayAttr(objects));
1949}
1950
1951static ParseResult parseOffloadingHandler(OpAsmParser &parser,
1952 Attribute &offloadingHandler) {
1953 if (succeeded(parser.parseOptionalLess())) {
1954 if (parser.parseAttribute(offloadingHandler))
1955 return failure();
1956 if (parser.parseGreater())
1957 return failure();
1958 }
1959 if (!offloadingHandler)
1960 offloadingHandler = parser.getBuilder().getAttr<SelectObjectAttr>(nullptr);
1961 return success();
1962}
1963
1965 Attribute offloadingHandler) {
1966 if (offloadingHandler != SelectObjectAttr::get(op->getContext(), nullptr))
1967 printer << '<' << offloadingHandler << '>';
1968}
1969
1970//===----------------------------------------------------------------------===//
1971// GPUMemcpyOp
1972//===----------------------------------------------------------------------===//
1973
1974LogicalResult MemcpyOp::verify() {
1975 auto srcType = getSrc().getType();
1976 auto dstType = getDst().getType();
1977
1978 if (getElementTypeOrSelf(srcType) != getElementTypeOrSelf(dstType))
1979 return emitOpError("arguments have incompatible element type");
1980
1981 if (failed(verifyCompatibleShape(srcType, dstType)))
1982 return emitOpError("arguments have incompatible shape");
1983
1984 return success();
1985}
1986
1987namespace {
1988/// Erases a common case of copy ops where a destination value is used only by
1989/// the copy op, alloc and dealloc ops.
1990struct EraseTrivialCopyOp : public OpRewritePattern<MemcpyOp> {
1991 using OpRewritePattern<MemcpyOp>::OpRewritePattern;
1992
1993 LogicalResult matchAndRewrite(MemcpyOp op,
1994 PatternRewriter &rewriter) const override {
1995 Value dest = op.getDst();
1996 Operation *destDefOp = dest.getDefiningOp();
1997 // `dest` must be defined by an op having Allocate memory effect in order to
1998 // perform the folding.
1999 if (!destDefOp ||
2001 return failure();
2002 // We can erase `op` iff `dest` has no other use apart from its
2003 // use by `op` and dealloc ops.
2004 if (llvm::any_of(dest.getUsers(), [op, dest](Operation *user) {
2005 return user != op &&
2006 !hasSingleEffect<MemoryEffects::Free>(user, dest);
2007 }))
2008 return failure();
2009 // We can perform the folding if and only if op has a single async
2010 // dependency and produces an async token as result, or if it does not have
2011 // any async dependency and does not produce any async token result.
2012 if (op.getAsyncDependencies().size() > 1 ||
2013 ((op.getAsyncDependencies().empty() && op.getAsyncToken()) ||
2014 (!op.getAsyncDependencies().empty() && !op.getAsyncToken())))
2015 return failure();
2016 rewriter.replaceOp(op, op.getAsyncDependencies());
2017 return success();
2018 }
2019};
2020} // end anonymous namespace
2021
2022void MemcpyOp::getCanonicalizationPatterns(RewritePatternSet &results,
2023 MLIRContext *context) {
2024 results.add<EraseTrivialCopyOp>(context);
2025}
2026
2027//===----------------------------------------------------------------------===//
2028// GPU_SubgroupMmaLoadMatrixOp
2029//===----------------------------------------------------------------------===//
2030
2031LogicalResult SubgroupMmaLoadMatrixOp::verify() {
2032 auto srcType = getSrcMemref().getType();
2033 auto resType = getRes().getType();
2034 auto resMatrixType = llvm::cast<gpu::MMAMatrixType>(resType);
2035 auto operand = resMatrixType.getOperand();
2036 auto srcMemrefType = llvm::cast<MemRefType>(srcType);
2037
2038 if (!srcMemrefType.isLastDimUnitStride())
2039 return emitError(
2040 "expected source memref most minor dim must have unit stride");
2041
2042 if (operand != "AOp" && operand != "BOp" && operand != "COp")
2043 return emitError("only AOp, BOp and COp can be loaded");
2044
2045 return success();
2046}
2047
2048//===----------------------------------------------------------------------===//
2049// GPU_SubgroupMmaStoreMatrixOp
2050//===----------------------------------------------------------------------===//
2051
2052LogicalResult SubgroupMmaStoreMatrixOp::verify() {
2053 auto srcType = getSrc().getType();
2054 auto dstType = getDstMemref().getType();
2055 auto srcMatrixType = llvm::cast<gpu::MMAMatrixType>(srcType);
2056 auto dstMemrefType = llvm::cast<MemRefType>(dstType);
2057
2058 if (!dstMemrefType.isLastDimUnitStride())
2059 return emitError(
2060 "expected destination memref most minor dim must have unit stride");
2061
2062 if (srcMatrixType.getOperand() != "COp")
2063 return emitError(
2064 "expected the operand matrix being stored to have 'COp' operand type");
2065
2066 return success();
2067}
2068
2069//===----------------------------------------------------------------------===//
2070// GPU_SubgroupMmaComputeOp
2071//===----------------------------------------------------------------------===//
2072
2073LogicalResult SubgroupMmaComputeOp::verify() {
2074 enum OperandMap { A, B, C };
2076 opTypes.push_back(llvm::cast<MMAMatrixType>(getOpA().getType()));
2077 opTypes.push_back(llvm::cast<MMAMatrixType>(getOpB().getType()));
2078 opTypes.push_back(llvm::cast<MMAMatrixType>(getOpC().getType()));
2079
2080 if (opTypes[A].getOperand() != "AOp" || opTypes[B].getOperand() != "BOp" ||
2081 opTypes[C].getOperand() != "COp")
2082 return emitError("operands must be in the order AOp, BOp, COp");
2083
2084 ArrayRef<int64_t> aShape, bShape, cShape;
2085 aShape = opTypes[A].getShape();
2086 bShape = opTypes[B].getShape();
2087 cShape = opTypes[C].getShape();
2088
2089 if (aShape[1] != bShape[0] || aShape[0] != cShape[0] ||
2090 bShape[1] != cShape[1])
2091 return emitError("operand shapes do not satisfy matmul constraints");
2092
2093 return success();
2094}
2095
2096LogicalResult MemcpyOp::fold(FoldAdaptor adaptor,
2098 return memref::foldMemRefCast(*this);
2099}
2100
2101LogicalResult MemsetOp::fold(FoldAdaptor adaptor,
2103 return memref::foldMemRefCast(*this);
2104}
2105
2106//===----------------------------------------------------------------------===//
2107// GPU_WaitOp
2108//===----------------------------------------------------------------------===//
2109
2110namespace {
2111/// Remove gpu.wait op use of gpu.wait op def without async dependencies.
2112/// %t = gpu.wait async [] // No async dependencies.
2113/// ... gpu.wait ... [%t, ...] // %t can be removed.
2114struct EraseRedundantGpuWaitOpPairs : public OpRewritePattern<WaitOp> {
2115public:
2116 using Base::Base;
2117
2118 LogicalResult matchAndRewrite(WaitOp op,
2119 PatternRewriter &rewriter) const final {
2120 auto predicate = [](Value value) {
2121 auto waitOp = value.getDefiningOp<WaitOp>();
2122 return waitOp && waitOp->getNumOperands() == 0;
2123 };
2124 if (llvm::none_of(op.getAsyncDependencies(), predicate))
2125 return failure();
2126 SmallVector<Value> validOperands;
2127 for (Value operand : op->getOperands()) {
2128 if (predicate(operand))
2129 continue;
2130 validOperands.push_back(operand);
2131 }
2132 rewriter.modifyOpInPlace(op, [&]() { op->setOperands(validOperands); });
2133 return success();
2134 }
2135};
2136
2137/// Simplify trivial gpu.wait ops for the following patterns.
2138/// 1. %t = gpu.wait async ... ops, where %t has no uses (regardless of async
2139/// dependencies).
2140/// 2. %t1 = gpu.wait async [%t0], in this case, we can replace uses of %t1 with
2141/// %t0.
2142/// 3. gpu.wait [] ops, i.e gpu.wait ops that neither have any async
2143/// dependencies nor return any token.
2144struct SimplifyGpuWaitOp : public OpRewritePattern<WaitOp> {
2145public:
2146 using Base::Base;
2147
2148 LogicalResult matchAndRewrite(WaitOp op,
2149 PatternRewriter &rewriter) const final {
2150 // Erase gpu.wait ops that neither have any async dependencies nor return
2151 // any async token.
2152 if (op.getAsyncDependencies().empty() && !op.getAsyncToken()) {
2153 rewriter.eraseOp(op);
2154 return success();
2155 }
2156 // Replace uses of %t1 = gpu.wait async [%t0] ops with %t0 and erase the op.
2157 if (llvm::hasSingleElement(op.getAsyncDependencies()) &&
2158 op.getAsyncToken()) {
2159 rewriter.replaceOp(op, op.getAsyncDependencies());
2160 return success();
2161 }
2162 // Erase %t = gpu.wait async ... ops, where %t has no uses.
2163 if (op.getAsyncToken() && op.getAsyncToken().use_empty()) {
2164 rewriter.eraseOp(op);
2165 return success();
2166 }
2167 return failure();
2168 }
2169};
2170} // end anonymous namespace
2171
2172void WaitOp::getCanonicalizationPatterns(RewritePatternSet &results,
2173 MLIRContext *context) {
2174 results.add<EraseRedundantGpuWaitOpPairs, SimplifyGpuWaitOp>(context);
2175}
2176
2177//===----------------------------------------------------------------------===//
2178// GPU_AllocOp
2179//===----------------------------------------------------------------------===//
2180
2181LogicalResult AllocOp::verify() {
2182 auto memRefType = llvm::cast<MemRefType>(getMemref().getType());
2183
2184 if (failed(verifyDynamicDimensionCount(getOperation(), memRefType,
2185 getDynamicSizes())))
2186 return failure();
2187
2188 unsigned numSymbols = 0;
2189 if (!memRefType.getLayout().isIdentity())
2190 numSymbols = memRefType.getLayout().getAffineMap().getNumSymbols();
2191 if (getSymbolOperands().size() != numSymbols) {
2192 return emitOpError(
2193 "symbol operand count does not equal memref symbol count");
2194 }
2195
2196 return success();
2197}
2198
2199namespace {
2200/// Folding of memref.dim(gpu.alloc(%size), %idx) -> %size similar to
2201/// `memref::AllocOp`.
2202struct SimplifyDimOfAllocOp : public OpRewritePattern<memref::DimOp> {
2203 using OpRewritePattern<memref::DimOp>::OpRewritePattern;
2204
2205 LogicalResult matchAndRewrite(memref::DimOp dimOp,
2206 PatternRewriter &rewriter) const override {
2207 std::optional<int64_t> index = dimOp.getConstantIndex();
2208 if (!index)
2209 return failure();
2210
2211 int64_t indexVal = index.value();
2212 auto memrefType = llvm::dyn_cast<MemRefType>(dimOp.getSource().getType());
2213 if (!memrefType || indexVal < 0 || indexVal >= memrefType.getRank() ||
2214 !memrefType.isDynamicDim(indexVal))
2215 return failure();
2216
2217 auto alloc = dimOp.getSource().getDefiningOp<AllocOp>();
2218 if (!alloc)
2219 return failure();
2220
2221 Value substituteOp = *(alloc.getDynamicSizes().begin() +
2222 memrefType.getDynamicDimIndex(indexVal));
2223 rewriter.replaceOp(dimOp, substituteOp);
2224 return success();
2225 }
2226};
2227} // namespace
2228
2229void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
2230 MLIRContext *context) {
2231 results.add<SimplifyDimOfAllocOp>(context);
2232}
2233
2234//===----------------------------------------------------------------------===//
2235// DynamicSharedMemoryOp
2236//===----------------------------------------------------------------------===//
2237
2238LogicalResult gpu::DynamicSharedMemoryOp::verify() {
2239 if (!getOperation()->getParentWithTrait<OpTrait::SymbolTable>())
2240 return emitOpError() << "must be inside an op with symbol table";
2241
2242 MemRefType memrefType = getResultMemref().getType();
2243 // Check address space
2244 if (!GPUDialect::hasWorkgroupMemoryAddressSpace(memrefType)) {
2245 return emitOpError() << "address space must be "
2246 << gpu::AddressSpaceAttr::getMnemonic() << "<"
2247 << stringifyEnum(gpu::AddressSpace::Workgroup) << ">";
2248 }
2249 if (memrefType.hasStaticShape()) {
2250 return emitOpError() << "result memref type must be memref<?xi8, "
2251 "#gpu.address_space<workgroup>>";
2252 }
2253 return success();
2254}
2255
2256//===----------------------------------------------------------------------===//
2257// GPU WarpExecuteOnLane0Op
2258//===----------------------------------------------------------------------===//
2259
2260void WarpExecuteOnLane0Op::print(OpAsmPrinter &p) {
2261 p << "(" << getLaneid() << ")";
2262
2263 SmallVector<StringRef> coreAttr = {getWarpSizeAttrName()};
2264 p << "[" << getWarpSize() << "]";
2265
2266 if (!getArgs().empty())
2267 p << " args(" << getArgs() << " : " << getArgs().getTypes() << ")";
2268 if (!getResults().empty())
2269 p << " -> (" << getResults().getTypes() << ')';
2270 p << " ";
2271 p.printRegion(getRegion(),
2272 /*printEntryBlockArgs=*/true,
2273 /*printBlockTerminators=*/!getResults().empty());
2275 getOperation()->getDiscardableAttrDictionary().getValue(), coreAttr);
2276}
2277
2278ParseResult WarpExecuteOnLane0Op::parse(OpAsmParser &parser,
2280 // Create the region.
2281 result.regions.reserve(1);
2282 Region *warpRegion = result.addRegion();
2283
2284 auto &builder = parser.getBuilder();
2286
2287 // Parse predicate operand.
2288 if (parser.parseLParen() ||
2289 parser.parseOperand(laneId, /*allowResultNumber=*/false) ||
2290 parser.parseRParen())
2291 return failure();
2292
2293 int64_t warpSize;
2294 if (parser.parseLSquare() || parser.parseInteger(warpSize) ||
2295 parser.parseRSquare())
2296 return failure();
2297 result.addAttribute(getWarpSizeAttrName(OperationName(getOperationName(),
2298 builder.getContext())),
2299 builder.getI64IntegerAttr(warpSize));
2300
2301 if (parser.resolveOperand(laneId, builder.getIndexType(), result.operands))
2302 return failure();
2303
2304 llvm::SMLoc inputsOperandsLoc;
2306 SmallVector<Type> inputTypes;
2307 if (succeeded(parser.parseOptionalKeyword("args"))) {
2308 if (parser.parseLParen())
2309 return failure();
2310
2311 inputsOperandsLoc = parser.getCurrentLocation();
2312 if (parser.parseOperandList(inputsOperands) ||
2313 parser.parseColonTypeList(inputTypes) || parser.parseRParen())
2314 return failure();
2315 }
2316 if (parser.resolveOperands(inputsOperands, inputTypes, inputsOperandsLoc,
2317 result.operands))
2318 return failure();
2319
2320 // Parse optional results type list.
2321 if (parser.parseOptionalArrowTypeList(result.types))
2322 return failure();
2323 // Parse the region.
2324 if (parser.parseRegion(*warpRegion, /*arguments=*/{},
2325 /*argTypes=*/{}))
2326 return failure();
2327 WarpExecuteOnLane0Op::ensureTerminator(*warpRegion, builder, result.location);
2328
2329 // Parse the optional attribute list.
2330 if (parser.parseOptionalAttrDict(result.attributes))
2331 return failure();
2332 return success();
2333}
2334
2335void WarpExecuteOnLane0Op::getSuccessorRegions(
2337 if (!point.isParent()) {
2338 regions.push_back(RegionSuccessor(getOperation()));
2339 return;
2340 }
2341
2342 // The warp region is always executed
2343 regions.push_back(RegionSuccessor(&getWarpRegion()));
2344}
2345
2346ValueRange WarpExecuteOnLane0Op::getSuccessorInputs(RegionSuccessor successor) {
2347 return successor.isOperation() ? ValueRange(getResults()) : ValueRange();
2348}
2349void WarpExecuteOnLane0Op::build(OpBuilder &builder, OperationState &result,
2350 TypeRange resultTypes, Value laneId,
2351 int64_t warpSize) {
2352 build(builder, result, resultTypes, laneId, warpSize,
2353 /*operands=*/{}, /*argTypes=*/{});
2354}
2355
2356void WarpExecuteOnLane0Op::build(OpBuilder &builder, OperationState &result,
2357 TypeRange resultTypes, Value laneId,
2358 int64_t warpSize, ValueRange args,
2359 TypeRange blockArgTypes) {
2360 result.addOperands(laneId);
2361 result.addAttribute(getAttributeNames()[0],
2362 builder.getI64IntegerAttr(warpSize));
2363 result.addTypes(resultTypes);
2364 result.addOperands(args);
2365 assert(args.size() == blockArgTypes.size());
2366 OpBuilder::InsertionGuard guard(builder);
2367 Region *warpRegion = result.addRegion();
2368 Block *block = builder.createBlock(warpRegion);
2369 for (auto [type, arg] : llvm::zip_equal(blockArgTypes, args))
2370 block->addArgument(type, arg.getLoc());
2371}
2372
2373/// Helper check if the distributed vector type is consistent with the expanded
2374/// type and distributed size.
2375static LogicalResult verifyDistributedType(Type expanded, Type distributed,
2376 int64_t warpSize, Operation *op) {
2377 // If the types matches there is no distribution.
2378 if (expanded == distributed)
2379 return success();
2380 auto expandedVecType = llvm::dyn_cast<VectorType>(expanded);
2381 auto distributedVecType = llvm::dyn_cast<VectorType>(distributed);
2382 if (!expandedVecType || !distributedVecType)
2383 return op->emitOpError("expected vector type for distributed operands.");
2384 if (expandedVecType.getRank() != distributedVecType.getRank() ||
2385 expandedVecType.getElementType() != distributedVecType.getElementType())
2386 return op->emitOpError(
2387 "expected distributed vectors to have same rank and element type.");
2388
2389 SmallVector<int64_t> scales(expandedVecType.getRank(), 1);
2390 for (int64_t i = 0, e = expandedVecType.getRank(); i < e; i++) {
2391 int64_t eDim = expandedVecType.getDimSize(i);
2392 int64_t dDim = distributedVecType.getDimSize(i);
2393 if (eDim == dDim)
2394 continue;
2395 if (eDim % dDim != 0)
2396 return op->emitOpError()
2397 << "expected expanded vector dimension #" << i << " (" << eDim
2398 << ") to be a multipler of the distributed vector dimension ("
2399 << dDim << ")";
2400 scales[i] = eDim / dDim;
2401 }
2402 if (llvm::product_of(scales) != warpSize)
2403 return op->emitOpError()
2404 << "incompatible distribution dimensions from " << expandedVecType
2405 << " to " << distributedVecType << " with warp size = " << warpSize;
2406
2407 return success();
2408}
2409
2410LogicalResult WarpExecuteOnLane0Op::verify() {
2411 if (getArgs().size() != getWarpRegion().getNumArguments())
2412 return emitOpError(
2413 "expected same number op arguments and block arguments.");
2414 auto yield = dyn_cast<gpu::YieldOp>(getBody()->getTerminator());
2415 if (!yield)
2416 return emitOpError("expected body to be terminated with 'gpu.yield'");
2417 if (yield.getNumOperands() != getNumResults())
2418 return emitOpError(
2419 "expected same number of yield operands and return values.");
2420 int64_t warpSize = getWarpSize();
2421 for (auto [regionArg, arg] :
2422 llvm::zip_equal(getWarpRegion().getArguments(), getArgs())) {
2423 if (failed(verifyDistributedType(regionArg.getType(), arg.getType(),
2424 warpSize, getOperation())))
2425 return failure();
2426 }
2427 for (auto [yieldOperand, result] :
2428 llvm::zip_equal(yield.getOperands(), getResults())) {
2429 if (failed(verifyDistributedType(yieldOperand.getType(), result.getType(),
2430 warpSize, getOperation())))
2431 return failure();
2432 }
2433 return success();
2434}
2435bool WarpExecuteOnLane0Op::areTypesCompatible(Type lhs, Type rhs) {
2436 return succeeded(
2437 verifyDistributedType(lhs, rhs, getWarpSize(), getOperation()));
2438}
2439
2440gpu::YieldOp WarpExecuteOnLane0Op::getTerminator() {
2441 return cast<gpu::YieldOp>(getBody()->getTerminator());
2442}
2443
2444//===----------------------------------------------------------------------===//
2445// GPU_SubgroupBroadcastOp
2446//===----------------------------------------------------------------------===//
2447
2448void gpu::SubgroupBroadcastOp::inferResultRanges(
2449 ArrayRef<ConstantIntRanges> argRanges, SetIntRangeFn setResultRange) {
2450 setResultRange(getResult(), argRanges.front());
2451}
2452
2453Speculation::Speculatability gpu::SubgroupBroadcastOp::getSpeculatability() {
2454 switch (getBroadcastType()) {
2455 case BroadcastType::first_active_lane:
2456 // Cannot speculate first_lane broadcast, because speculating it across
2457 // control flow can change the active lanes.
2459 case BroadcastType::specific_lane:
2460 // Speculation should be safe as long as we inside structured control flow.
2462 }
2463 llvm_unreachable("Unknown BroadcastType");
2464}
2465
2466LogicalResult gpu::SubgroupBroadcastOp::verify() {
2467 switch (getBroadcastType()) {
2468 case BroadcastType::first_active_lane:
2469 if (getLane())
2470 return emitOpError()
2471 << "lane can only be specified for `specific_lane` broadcast";
2472 return success();
2473 case BroadcastType::specific_lane:
2474 if (!getLane())
2475 return emitOpError()
2476 << "lane must be specified for `specific_lane` broadcast";
2477 return success();
2478 }
2479 llvm_unreachable("Unknown BroadcastType");
2480}
2481
2482OpFoldResult gpu::SubgroupBroadcastOp::fold(FoldAdaptor /*adaptor*/) {
2483 // Broadcast result is always uniform.
2484 if (auto prev = getSrc().getDefiningOp<SubgroupBroadcastOp>())
2485 return prev.getResult();
2486
2487 return nullptr;
2488}
2489
2490//===----------------------------------------------------------------------===//
2491// GPU_BallotOp
2492//===----------------------------------------------------------------------===//
2493
2494// No custom implementations needed; ballot uses default behavior from ODS.
2495
2496#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
2497
2498#define GET_OP_CLASSES
2499#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
return success()
static void printLaunchFuncOperands(OpAsmPrinter &printer, Operation *, OperandRange operands, TypeRange types)
static ParseResult parseAsyncDependencies(OpAsmParser &parser, Type &asyncTokenType, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &asyncDependencies)
Parses an optional list of async operands with an optional leading keyword.
static void setAttributionAttrs(GPUFuncOp op, unsigned index, DictionaryAttr value, StringAttr attrName)
static void printAttributions(OpAsmPrinter &p, StringRef keyword, ArrayRef< BlockArgument > values, ArrayAttr attributes={})
static LogicalResult verifyDistributedType(Type expanded, Type distributed, int64_t warpSize, Operation *op)
Helper check if the distributed vector type is consistent with the expanded type and distributed size...
static void printAsyncDependencies(OpAsmPrinter &printer, Operation *op, Type asyncTokenType, OperandRange asyncDependencies)
Prints optional async dependencies with its leading keyword.
static ParseResult parseSizeAssignment(OpAsmParser &parser, MutableArrayRef< OpAsmParser::UnresolvedOperand > sizes, MutableArrayRef< OpAsmParser::UnresolvedOperand > regionSizes, MutableArrayRef< OpAsmParser::UnresolvedOperand > indices, StringRef keyword)
static LogicalResult eraseRedundantGpuBarrierOps(BarrierOp op, PatternRewriter &rewriter)
Remove gpu.barrier after gpu.barrier, the threads are already synchronized!
static ParseResult parseOffloadingHandler(OpAsmParser &parser, Attribute &offloadingHandler)
static DictionaryAttr getAttributionAttrs(GPUFuncOp op, unsigned index, StringAttr attrName)
static void printLaunchDimType(OpAsmPrinter &printer, Operation *op, Type dimTy, Value clusterValue, Type clusterXTy, Type clusterYTy, Type clusterZTy)
static bool canMakeGroupOpUniform(Operation *op)
static std::string getSparseHandleKeyword(SparseHandleKind kind)
static LogicalResult verifyKnownLaunchSizeAttr(Operation *op, NamedAttribute attr)
static LogicalResult verifyLaunchAsyncModel(OpTy op)
static ParseResult parseAttributions(OpAsmParser &parser, StringRef keyword, SmallVectorImpl< OpAsmParser::Argument > &args)
Parses a GPU function memory attribution.
static ParseResult parseLaunchDimType(OpAsmParser &parser, Type &dimTy, std::optional< OpAsmParser::UnresolvedOperand > clusterValue, Type &clusterXTy, Type &clusterYTy, Type &clusterZTy)
static void setAttributionAttr(GPUFuncOp op, unsigned index, StringAttr name, Attribute value, StringAttr attrsName)
static ParseResult parseLaunchFuncOperands(OpAsmParser &parser, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &argNames, SmallVectorImpl< Type > &argTypes)
static void printOffloadingHandler(OpAsmPrinter &printer, Operation *op, Attribute offloadingHandler)
static LogicalResult verifyReduceOpAndType(gpu::AllReduceOperation opName, Type resType)
static void printSizeAssignment(OpAsmPrinter &p, KernelDim3 size, KernelDim3 operands, KernelDim3 ids)
static Attribute getAttributionAttr(GPUFuncOp op, unsigned index, StringAttr name, StringAttr attrsName)
static LogicalResult verifyAttributions(Operation *op, ArrayRef< BlockArgument > attributions, gpu::AddressSpace memorySpace)
Verifies a GPU function memory attribution.
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.
b
Return true if permutation is a valid permutation of the outer_dims_perm (case OuterOrInnerPerm::Oute...
ArrayAttr()
b getContext())
template bool mlir::hasSingleEffect< MemoryEffects::Allocate >(Operation *)
static void getDynamicSizes(RankedTensorType tp, ValueRange sizes, SmallVectorImpl< Value > &dynSizes)
Collects the dynamic dimension sizes for tp with the assumption that sizes are the dimension sizes fo...
static sycl::kernel * getKernel(ze_module_handle_t zeModule, const char *name)
This class contains all of the static information common to all instances of a registered Attribute.
static AbstractAttribute get(Dialect &dialect)
This method is used by Dialect objects when they register the list of attributes they contain.
ParseResult parseSymbolName(StringAttr &result)
Parse an -identifier and store it (without the '@' symbol) in a string attribute.
@ Paren
Parens surrounding zero or more operands.
@ OptionalSquare
Square brackets supporting zero or more ops, or nothing.
virtual ParseResult parseColonTypeList(SmallVectorImpl< Type > &result)=0
Parse a colon followed by a type list, which must have at least one type.
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.
virtual ParseResult parseOptionalKeyword(StringRef keyword)=0
Parse the given keyword if present.
MLIRContext * getContext() const
virtual Location getEncodedSourceLoc(SMLoc loc)=0
Re-encode the given source location as an MLIR location and return it.
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 ParseResult parseOptionalColon()=0
Parse a : token if present.
virtual ParseResult parseLSquare()=0
Parse a [ token.
virtual ParseResult parseRSquare()=0
Parse a ] token.
ParseResult parseInteger(IntT &result)
Parse an integer value from the stream.
virtual ParseResult parseLess()=0
Parse a '<' token.
virtual ParseResult parseDimensionList(SmallVectorImpl< int64_t > &dimensions, bool allowDynamic=true, bool withTrailingX=true)=0
Parse a dimension list of a tensor or memref type.
virtual ParseResult parseEqual()=0
Parse a = token.
virtual ParseResult parseOptionalAttrDictWithKeyword(NamedAttrList &result)=0
Parse a named dictionary into 'result' if the attributes keyword is present.
virtual ParseResult parseColonType(Type &result)=0
Parse a colon followed by a type.
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 parseOptionalString(std::string *string)=0
Parse a quoted string token if present.
virtual ParseResult parseOptionalLess()=0
Parse a '<' token if present.
virtual ParseResult parseGreater()=0
Parse a '>' token.
virtual ParseResult parseLParen()=0
Parse a ( token.
virtual ParseResult parseType(Type &result)=0
Parse a type.
virtual ParseResult parseComma()=0
Parse a , token.
virtual ParseResult parseOptionalArrowTypeList(SmallVectorImpl< Type > &result)=0
Parse an optional arrow followed by a type list.
ParseResult parseKeyword(StringRef keyword)
Parse a given keyword.
virtual ParseResult parseAttribute(Attribute &result, Type type={})=0
Parse an arbitrary attribute of a given type and return it in result.
virtual void printSymbolName(StringRef symbolRef)
Print the given string as a symbol reference, i.e.
Attributes are known-constant values of operations.
Definition Attributes.h:25
This class represents an argument of a Block.
Definition Value.h:306
Block represents an ordered list of Operations.
Definition Block.h:34
BlockArgument addArgument(Type type, Location loc)
Add one value to the argument list.
Definition Block.cpp:158
This class is a general helper class for creating context-global objects like types,...
Definition Builders.h:51
UnitAttr getUnitAttr()
Definition Builders.cpp:106
IntegerAttr getI32IntegerAttr(int32_t value)
Definition Builders.cpp:208
DenseI32ArrayAttr getDenseI32ArrayAttr(ArrayRef< int32_t > values)
Definition Builders.cpp:171
FunctionType getFunctionType(TypeRange inputs, TypeRange results)
Definition Builders.cpp:84
IntegerType getI32Type()
Definition Builders.cpp:71
IntegerAttr getI64IntegerAttr(int64_t value)
Definition Builders.cpp:120
Ty getType(Args &&...args)
Get or construct an instance of the type Ty with provided arguments.
Definition Builders.h:94
StringAttr getStringAttr(const Twine &bytes)
Definition Builders.cpp:271
ArrayAttr getArrayAttr(ArrayRef< Attribute > value)
Definition Builders.cpp:275
MLIRContext * getContext() const
Definition Builders.h:56
IndexType getIndexType()
Definition Builders.cpp:59
DictionaryAttr getDictionaryAttr(ArrayRef< NamedAttribute > value)
Definition Builders.cpp:112
NamedAttribute getNamedAttr(StringRef name, Attribute val)
Definition Builders.cpp:102
Attr getAttr(Args &&...args)
Get or construct an instance of the attribute Attr with provided arguments.
Definition Builders.h:101
The DialectAsmParser has methods for interacting with the asm parser when parsing attributes and type...
This is a pure-virtual base class that exposes the asmprinter hooks necessary to implement a custom p...
This class defines a virtual interface for reading a bytecode stream, providing hooks into the byteco...
virtual LogicalResult readVarInt(uint64_t &result)=0
Read a variable width integer.
virtual FailureOr< const DialectVersion * > getDialectVersion(StringRef dialectName) const =0
Retrieve the dialect version by name if available.
virtual LogicalResult readOptionalAttribute(Attribute &attr)=0
Read an optional reference to the given attribute.
virtual InFlightDiagnostic emitError(const Twine &msg={}) const =0
Emit an error to the reader.
virtual LogicalResult readAttribute(Attribute &result)=0
Read a reference to the given attribute.
This class defines a virtual interface for writing to a bytecode stream, providing hooks into the byt...
virtual void writeOptionalAttribute(Attribute attr)=0
virtual void writeVarInt(uint64_t value)=0
Write a variable width integer to the output stream.
virtual void writeAttribute(Attribute attr)=0
Write a reference to the given attribute.
virtual int64_t getBytecodeVersion() const =0
Return the bytecode version being emitted for.
Dialects are groups of MLIR operations, types and attributes, as well as behavior associated with the...
Definition Dialect.h:38
A symbol reference with a reference path containing a single element.
This class represents a diagnostic that is inflight and set to be reported.
This class defines the main interface for locations in MLIR and acts as a non-nullable wrapper around...
Definition Location.h:76
MLIRContext is the top-level object for a collection of MLIR operations.
Definition MLIRContext.h:64
NamedAttribute represents a combination of a name and an Attribute value.
Definition Attributes.h:164
StringAttr getName() const
Return the name of the attribute.
Attribute getValue() const
Return the value of the attribute.
Definition Attributes.h:179
The OpAsmParser has methods for interacting with the asm parser: parsing things from it,...
virtual size_t getNumResults() const =0
Return the number of declared SSA results.
virtual ParseResult parseRegion(Region &region, ArrayRef< Argument > arguments={}, bool enableNameShadowing=false)=0
Parses a region.
virtual ParseResult parseArgumentList(SmallVectorImpl< Argument > &result, Delimiter delimiter=Delimiter::None, bool allowType=false, bool allowAttrs=false)=0
Parse zero or more arguments with a specified surrounding delimiter.
virtual ParseResult resolveOperand(const UnresolvedOperand &operand, Type type, SmallVectorImpl< Value > &result)=0
Resolve an operand to an SSA value, emitting an error on failure.
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 parseOperand(UnresolvedOperand &result, bool allowResultNumber=true)=0
Parse a single SSA value operand name along with a result number if allowResultNumber is true.
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...
virtual void printOptionalAttrDict(ArrayRef< NamedAttribute > attrs, ArrayRef< StringRef > elidedAttrs={})=0
If the specified operation has attributes, print out an attribute dictionary with their values.
virtual void printRegion(Region &blocks, bool printEntryBlockArgs=true, bool printBlockTerminators=true, bool printEmptyBlock=false)=0
Prints a region.
RAII guard to reset the insertion point of the builder when destroyed.
Definition Builders.h:351
This class helps build Operations.
Definition Builders.h:210
Block * createBlock(Region *parent, Region::iterator insertPt={}, TypeRange argTypes={}, ArrayRef< Location > locs={})
Add new block with 'argTypes' arguments and set the insertion point to the end of it.
Definition Builders.cpp:439
void setInsertionPointToStart(Block *block)
Sets the insertion point to the start of the specified block.
Definition Builders.h:434
This class represents a single result from folding an operation.
This class implements the operand iterators for the Operation class.
Definition ValueRange.h:44
Operation is the basic unit of execution within MLIR.
Definition Operation.h:87
void setInherentAttr(StringAttr name, Attribute value)
Set an inherent attribute by name.
void insertOperands(unsigned index, ValueRange operands)
Insert the given operands into the operand list at the given 'index'.
Block * getBlock()
Returns the operation block that contains this operation.
Definition Operation.h:230
std::optional< Attribute > getInherentAttr(StringRef name)
Access an inherent attribute by name: returns an empty optional if there is no inherent attribute wit...
Location getLoc()
The source location the operation was defined or derived from.
Definition Operation.h:240
Operation * getParentOp()
Returns the closest surrounding operation that contains this operation or nullptr if this is a top-le...
Definition Operation.h:251
InFlightDiagnostic emitError(const Twine &message={})
Emit an error about fatal conditions with this operation, reporting up to any diagnostic handlers tha...
AttrClass getDiscardableAttrOfType(StringRef name)
Access a discardable attribute by name and cast it to AttrClass.
Definition Operation.h:493
MLIRContext * getContext()
Return the context this operation is associated with.
Definition Operation.h:233
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...
This class represents a point being branched from in the methods of the RegionBranchOpInterface.
bool isParent() const
Returns true if branching from the parent op.
This class represents a successor of a region.
bool isOperation() const
Return true if the successor is an operation.
This class contains a list of basic blocks and a link to the parent operation it is attached to.
Definition Region.h:26
Block & front()
Definition Region.h:65
bool empty()
Definition Region.h:60
RewritePatternSet & add(ConstructorArg &&arg, ConstructorArgs &&...args)
Add an instance of each of the pattern types 'Ts' to the pattern list with the given arguments.
virtual void replaceOp(Operation *op, ValueRange newValues)
Replace the results of the given (original) operation with the specified list of values (replacements...
virtual void eraseOp(Operation *op)
This method erases an operation that is known to have no uses.
void modifyOpInPlace(Operation *root, CallableT &&callable)
This method is a utility wrapper around an in-place modification of an operation.
virtual void replaceAllUsesWith(Value from, Value to)
Find uses of from and replace them with to.
This class represents a collection of SymbolTables.
virtual Operation * lookupNearestSymbolFrom(Operation *from, StringAttr symbol)
Returns the operation registered with the given symbol name within the closest parent operation of,...
static Operation * getNearestSymbolTable(Operation *from)
Returns the nearest symbol table from a given operation from.
This class provides an efficient unique identifier for a specific C++ type.
Definition TypeID.h:107
This class provides an abstraction over the various different ranges of value types.
Definition TypeRange.h:40
Instances of the Type class are uniqued, have an immutable identifier and an optional mutable compone...
Definition Types.h:74
bool isIndex() const
Definition Types.cpp:56
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
user_range getUsers() const
Definition Value.h:218
Location getLoc() const
Return the location of this value.
Definition Value.cpp:24
Operation * getDefiningOp() const
If this value is the result of an operation, return the operation that defines it.
Definition Value.cpp:18
static ConstantIndexOp create(OpBuilder &builder, Location location, int64_t value)
Definition ArithOps.cpp:434
MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations.
Definition GPUBase.h:132
ArrayRef< int64_t > getShape() const
Get shape of the matrix.
Definition GPUBase.cpp:129
Type getElementType() const
Get elementType of a single element.
Definition GPUBase.cpp:133
StringRef getOperand() const
The general form of operation this type supports is given by the equation C += A*B.
Definition GPUBase.cpp:135
static MMAMatrixType getChecked(function_ref< InFlightDiagnostic()> emitError, ArrayRef< int64_t > shape, Type elementType, StringRef operand)
Get MMAMatrixType at a particular location and verify construction Invariants.
Definition GPUBase.cpp:120
Speculatability
This enum is returned from the getSpeculatability method in the ConditionallySpeculatable op interfac...
constexpr auto Speculatable
constexpr auto NotSpeculatable
@ kDialectVersioning
Dialects versioning was added in version 1.
Definition Encoding.h:32
void addArgAndResultAttrs(Builder &builder, OperationState &result, ArrayRef< DictionaryAttr > argAttrs, ArrayRef< DictionaryAttr > resultAttrs, StringAttr argAttrsName, StringAttr resAttrsName)
Adds argument and result attributes, provided as argAttrs and resultAttrs arguments,...
llvm::unique_function< InFlightDiagnostic()> getDefaultDiagnosticEmitFn(MLIRContext *ctx)
Utility method to generate a callback that can be used to generate a diagnostic when checking the con...
ParseResult parseFunctionSignatureWithArguments(OpAsmParser &parser, bool allowVariadic, SmallVectorImpl< OpAsmParser::Argument > &arguments, bool &isVariadic, SmallVectorImpl< Type > &resultTypes, SmallVectorImpl< DictionaryAttr > &resultAttrs)
Parses a function signature using parser.
void printFunctionAttributes(OpAsmPrinter &p, Operation *op, ArrayRef< StringRef > elided={})
Prints the list of function prefixed with the "attributes" keyword.
void printFunctionSignature(OpAsmPrinter &p, FunctionOpInterface op, ArrayRef< Type > argTypes, bool isVariadic, ArrayRef< Type > resultTypes)
Prints the signature of the function-like operation op.
void registerAttributeStorage(MLIRContext *context)
Register storage whose definitions are private to MLIRGPUBase.
Definition GPUBase.cpp:502
void addAsyncDependency(Operation *op, Value token)
SparseHandleKind
Definition GPUBase.h:174
LogicalResult foldMemRefCast(Operation *op, Value inner=nullptr)
This is a common utility used for patterns of the form "someop(memref.cast) -> someop".
Definition MemRefOps.cpp:48
detail::InFlightRemark failed(Location loc, RemarkOpts opts)
Report an optimization remark that failed.
Definition Remarks.h:734
SmallVector< unsigned > getBlockSize(AffineMap dimToLvl)
Given the dimToLvl map, returns the block sizes in a vector.
Include the generated interface declarations.
bool matchPattern(Value value, const Pattern &pattern)
Entry point for matching a pattern over a Value.
Definition Matchers.h:490
llvm::function_ref< void(Value, const ConstantIntRanges &)> SetIntRangeFn
The type of the setResultRanges callback provided to ops implementing InferIntRangeInterface.
LogicalResult verifyDynamicDimensionCount(Operation *op, ShapedType type, ValueRange dynamicSizes)
Verify that the number of dynamic size operands matches the number of dynamic dimensions in the shape...
Type getType(OpFoldResult ofr)
Returns the int type of the integer in ofr.
Definition Utils.cpp:311
InFlightDiagnostic emitError(Location loc)
Utility method to emit an error message using this location.
Type getElementTypeOrSelf(Type type)
Return the element type or return the type itself.
detail::constant_int_predicate_matcher m_One()
Matches a constant scalar / vector splat / tensor splat integer one.
Definition Matchers.h:478
llvm::TypeSwitch< T, ResultT > TypeSwitch
Definition LLVM.h:147
LogicalResult verifyCompatibleShape(ArrayRef< int64_t > shape1, ArrayRef< int64_t > shape2)
Returns success if the given two shapes are compatible.
llvm::function_ref< Fn > function_ref
Definition LLVM.h:155
Simplify the gpu.launch when the range of a thread or block ID is trivially known to be one.
LogicalResult matchAndRewrite(LaunchOp op, PatternRewriter &rewriter) const override
This is the representation of an operand reference.
OpRewritePattern is a wrapper around RewritePattern that allows for matching and rewriting against an...
OpRewritePattern(MLIRContext *context, PatternBenefit benefit=1, ArrayRef< StringRef > generatedNames={})
This represents an operation in an abstracted form, suitable for use with the builder APIs.
T & getOrAddProperties()
Get (or create) the properties of the provided type to be set on the operation on creation.
Utility class for the GPU dialect to represent triples of Values accessible through ....
Definition GPUDialect.h:40