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"
49#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
55template <
typename... Attrs>
56void registerAttributeMetadata(
63struct GPUDialectBytecodeInterface : BytecodeDialectInterface {
64 using BytecodeDialectInterface::BytecodeDialectInterface;
66 void writeVersion(DialectBytecodeWriter &writer)
const override {
70 std::unique_ptr<DialectVersion>
71 readVersion(DialectBytecodeReader &reader)
const override {
76 reader.
emitError(
"unsupported gpu dialect bytecode version");
79 return std::make_unique<DialectVersion>();
84int64_t GPUMappingMaskAttr::getMaxNumPhysicalIds()
const {
return 64; }
96Value GPUMappingMaskAttr::createLogicalLinearMappingId(
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);
118Value GPUMappingMaskAttr::createIsActiveIdPredicate(
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,
131bool GPUDialect::isKernel(
Operation *op) {
132 if (
auto gpuFunc = dyn_cast<GPUFuncOp>(op))
133 return gpuFunc.isKernel();
134 return static_cast<bool>(
141struct GPUInlinerInterface :
public DialectInlinerInterface {
142 using DialectInlinerInterface::DialectInlinerInterface;
145 bool isLegalToInline(Operation *, Region *,
bool, IRMapping &)
const final {
151void GPUDialect::initialize() {
152 addInterfaces<GPUDialectBytecodeInterface>();
153 addTypes<AsyncTokenType>();
154 addTypes<MMAMatrixType>();
155 addTypes<NamedBarrierType>();
156 addTypes<SparseDnTensorHandleType>();
157 addTypes<SparseSpMatHandleType>();
158 addTypes<SparseSpGEMMOpHandleType>();
161#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
165 registerAttributeMetadata<
166#define GET_ATTRDEF_LIST
167#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
169 addAttribute(
id, std::move(attribute));
172 addInterfaces<GPUInlinerInterface>();
173 declarePromisedInterface<bufferization::BufferDeallocationOpInterface,
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>();
188 return "sparse.dntensor_handle";
190 return "sparse.spmat_handle";
192 return "sparse.spgemmop_handle";
194 llvm_unreachable(
"unknown sparse handle kind");
206 if (keyword ==
"async.token")
209 if (keyword ==
"mma_matrix") {
238 shape, elementType, operand);
241 if (keyword ==
"named_barrier")
259 .Case<NamedBarrierType>([&](
Type) { os <<
"named_barrier"; })
260 .Case<SparseDnTensorHandleType>([&](
Type) {
263 .Case<SparseSpMatHandleType>(
265 .Case<SparseSpGEMMOpHandleType>([&](
Type) {
271 for (
auto dim =
shape.begin(), e =
shape.end() - 1; dim != e; ++dim)
274 os <<
", \"" << fragTy.
getOperand() <<
"\"" <<
'>';
276 .DefaultUnreachable(
"unexpected 'gpu' type kind");
281 auto array = dyn_cast<DenseI32ArrayAttr>(attr.
getValue());
284 " must be a dense i32 array");
285 if (array.size() != 3)
287 " must contain exactly 3 elements");
291LogicalResult GPUDialect::verifyOperationAttribute(
Operation *op,
293 if (attr.
getName() == getKnownBlockSizeAttrHelper().getName())
295 if (attr.
getName() == getKnownGridSizeAttrHelper().getName())
297 if (attr.
getName() == getKnownClusterSizeAttrHelper().getName())
299 if (!llvm::isa<UnitAttr>(attr.
getValue()) ||
300 attr.
getName() != getContainerModuleAttrName())
303 auto module = dyn_cast<ModuleOp>(op);
306 << getContainerModuleAttrName() <<
"' attribute to be attached to '"
307 << ModuleOp::getOperationName() <<
'\'';
321 return parser.
emitError(loc,
"needs to be named when marked 'async'");
336 if (asyncDependencies.empty())
340 printer << llvm::interleaved_array(asyncDependencies);
368 p <<
' ' << keyword <<
'(';
369 llvm::interleaveComma(
370 llvm::enumerate(values), p, [&p, attributes](
auto pair) {
372 p << v <<
" : " << v.
getType();
374 size_t attributionIndex = pair.index();
375 DictionaryAttr attrs;
376 if (attributes && attributionIndex < attributes.size())
377 attrs = llvm::cast<DictionaryAttr>(attributes[attributionIndex]);
387 gpu::AddressSpace memorySpace) {
388 for (
Value v : attributions) {
389 auto type = llvm::dyn_cast<MemRefType>(v.
getType());
391 return op->
emitOpError() <<
"expected memref type in attribution";
396 llvm::dyn_cast_or_null<gpu::AddressSpaceAttr>(type.getMemorySpace());
399 if (addressSpace.getValue() != memorySpace)
401 <<
"expected memory space " << stringifyAddressSpace(memorySpace)
402 <<
" in attribution";
413 using Kind = gpu::AllReduceOperation;
414 if (llvm::is_contained({Kind::MINNUMF, Kind::MAXNUMF, Kind::MINIMUMF,
415 Kind::MAXIMUMF, Kind::MINIMUMNUMF, Kind::MAXIMUMNUMF},
417 if (!isa<FloatType>(resType))
421 if (llvm::is_contained({Kind::MINSI, Kind::MINUI, Kind::MAXSI, Kind::MAXUI,
422 Kind::AND, Kind::OR, Kind::XOR},
424 if (!isa<IntegerType>(resType))
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");
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");
452 return emitError(
"expected gpu.yield op in region");
454 gpu::AllReduceOperation opName = *getOp();
456 return emitError() <<
'`' << gpu::stringifyAllReduceOperation(opName)
457 <<
"` reduction operation is not compatible with type "
466 auto launchOp = dyn_cast<gpu::LaunchOp>(op->
getParentOp());
470 Region &body = launchOp.getBody();
471 assert(!body.
empty() &&
"Invalid region");
490LogicalResult gpu::SubgroupReduceOp::verify() {
492 if (
auto vecTy = dyn_cast<VectorType>(elemType)) {
493 if (vecTy.isScalable())
494 return emitOpError() <<
"is not compatible with scalable vector types";
496 elemType = vecTy.getElementType();
499 gpu::AllReduceOperation opName = getOp();
501 return emitError() <<
'`' << gpu::stringifyAllReduceOperation(opName)
502 <<
"` reduction operation is not compatible with type "
506 auto clusterSize = getClusterSize();
508 uint32_t size = *clusterSize;
509 if (!llvm::isPowerOf2_32(size)) {
510 return emitOpError() <<
"cluster size " << size
511 <<
" is not a power of two";
515 uint32_t stride = getClusterStride();
516 if (stride != 1 && !clusterSize) {
517 return emitOpError() <<
"cluster stride can only be specified if cluster "
520 if (!llvm::isPowerOf2_32(stride)) {
521 return emitOpError() <<
"cluster stride " << stride
522 <<
" is not a power of two";
529 if (getClusterSize() == 1)
546 if (!op->template hasTrait<OpTrait::AttrSizedOperandSegments>())
550 auto sizeAttr = dyn_cast_or_null<DenseI32ArrayAttr>(
570 Value getBlockSizeZ,
Value dynamicSharedMemorySize,
578 if (!workgroupAttributions.empty())
580 getWorkgroupAttributionsAttrName(
result.name),
584 result.addOperands(asyncDependencies);
589 result.addOperands({gridSizeX, gridSizeY, gridSizeZ, getBlockSizeX,
590 getBlockSizeY, getBlockSizeZ});
592 result.addOperands(clusterSizeX);
594 result.addOperands(clusterSizeY);
596 result.addOperands(clusterSizeZ);
597 if (dynamicSharedMemorySize)
598 result.addOperands(dynamicSharedMemorySize);
600 result.addOperands(asyncObject);
604 result.addAttribute(getModuleAttrName(
result.name), module);
606 result.addAttribute(getFunctionAttrName(
result.name), function);
614 for (
unsigned i = 0; i < kNumConfigRegionAttributes; ++i)
617 for (
Type argTy : workgroupAttributions)
619 for (
Type argTy : privateAttributions)
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(),
634 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
635 auto args = getBody().getArguments();
640 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
641 auto args = getBody().getArguments();
646 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
647 auto args = getBody().getArguments();
652 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
653 auto args = getBody().getArguments();
654 return KernelDim3{args[9], args[10], args[11]};
657std::optional<KernelDim3> LaunchOp::getClusterIds() {
658 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
659 if (!hasClusterSize())
661 auto args = getBody().getArguments();
662 return KernelDim3{args[12], args[13], args[14]};
665std::optional<KernelDim3> LaunchOp::getClusterSize() {
666 assert(!getBody().empty() &&
"LaunchOp body must not be empty.");
667 if (!hasClusterSize())
669 auto args = getBody().getArguments();
670 return KernelDim3{args[15], args[16], args[17]};
673KernelDim3 LaunchOp::getGridSizeOperandValues() {
674 auto operands = getOperands().drop_front(getAsyncDependencies().size());
675 return KernelDim3{operands[0], operands[1], operands[2]};
678KernelDim3 LaunchOp::getBlockSizeOperandValues() {
679 auto operands = getOperands().drop_front(getAsyncDependencies().size());
680 return KernelDim3{operands[3], operands[4], operands[5]};
683std::optional<KernelDim3> LaunchOp::getClusterSizeOperandValues() {
684 auto operands = getOperands().drop_front(getAsyncDependencies().size());
685 if (!hasClusterSize())
687 return KernelDim3{operands[6], operands[7], operands[8]};
690template <
typename OpTy>
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");
703LogicalResult LaunchOp::verify() {
707 if (!(hasClusterSize()) &&
708 (getClusterSizeX() || getClusterSizeY() || getClusterSizeZ()))
709 return emitOpError() <<
"cluster size must be all present";
713LogicalResult LaunchOp::verifyRegions() {
717 if (getBody().empty()) {
718 return emitOpError(
"body region is empty");
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;
731 GPUDialect::getWorkgroupAddressSpace())) ||
733 GPUDialect::getPrivateAddressSpace())))
738 for (
Block &block : getBody()) {
741 if (block.back().getNumSuccessors() != 0)
743 if (!isa<gpu::TerminatorOp>(&block.back())) {
746 .append(
"expected '", gpu::TerminatorOp::getOperationName(),
747 "' or a terminator with successors")
748 .attachNote(getLoc())
749 .append(
"in '", LaunchOp::getOperationName(),
"' body region");
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 <<
')';
769 if (
auto asyncObject = getAsyncObject()) {
770 p <<
" <" << asyncObject <<
" : " << asyncObject.
getType() <<
">";
772 if (getAsyncToken()) {
774 if (!getAsyncDependencies().empty())
775 p <<
" [" << getAsyncDependencies() <<
']';
778 if (hasClusterSize()) {
779 p <<
' ' << getClustersKeyword();
781 getClusterSizeOperandValues().value(),
782 getClusterIds().value());
784 p <<
' ' << getBlocksKeyword();
787 p <<
' ' << getThreadsKeyword();
790 if (getDynamicSharedMemorySize())
791 p <<
' ' << getDynamicSharedMemorySizeKeyword() <<
' '
792 << getDynamicSharedMemorySize();
795 StringRef moduleAttrName = getModuleAttrName();
796 if (
auto module = getModule()) {
797 p <<
' ' << moduleAttrName <<
'(';
802 StringRef functionAttrName = getFunctionAttrName();
803 if (
auto function = getFunction()) {
804 p <<
' ' << functionAttrName <<
'(';
809 if (getCooperative())
819 (*this)->getDiscardableAttrDictionary().getValue(), {
820 LaunchOp::getOperandSegmentSizeAttr(),
821 getWorkgroupAttributionsAttrName(), getCooperativeAttrName(),
822 moduleAttrName, functionAttrName});
837 assert(
indices.size() == 3 &&
"space for three indices expected");
844 if (args.size() != 3) {
846 << keyword <<
" expects 3 arguments, but got " << args.size();
848 std::move(args.begin(), args.end(),
indices.begin());
850 for (
int i = 0; i < 3; ++i) {
876 sizes(LaunchOp::kNumConfigOperands);
880 LaunchOp::kNumConfigRegionAttributes);
884 Type asyncObjectType;
885 bool hasAsyncObject =
false;
887 hasAsyncObject =
true;
905 "gpu.launch requires 'async' keyword to return a value");
906 result.types.push_back(asyncTokenType);
909 bool hasCluster =
false;
913 regionArgs.resize(18);
922 parser, sizesRef.drop_front(6), regionArgsRef.slice(15, 3),
923 regionArgsRef.slice(12, 3), LaunchOp::getClustersKeyword()))
931 if (parser.
parseKeyword(LaunchOp::getBlocksKeyword()) ||
933 regionArgsRef.slice(6, 3), regionArgsRef.slice(0, 3),
934 LaunchOp::getBlocksKeyword()) ||
937 regionArgsRef.slice(9, 3), regionArgsRef.slice(3, 3),
938 LaunchOp::getThreadsKeyword()) ||
944 bool hasDynamicSharedMemorySize =
false;
946 LaunchOp::getDynamicSharedMemorySizeKeyword())) {
947 hasDynamicSharedMemorySize =
true;
957 asyncObjectType,
result.operands))
961 StringRef moduleAttrName = getModuleAttrName(
result.name);
971 StringRef functionAttrName = getFunctionAttrName(
result.name);
990 LaunchOp::kNumConfigRegionAttributes + 6,
index);
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);
1008 unsigned numWorkgroupAttrs = regionArguments.size() -
1009 LaunchOp::kNumConfigRegionAttributes -
1010 (hasCluster ? 6 : 0);
1011 if (numWorkgroupAttrs != 0)
1012 result.addAttribute(LaunchOp::getWorkgroupAttributionsAttrName(
result.name),
1029 segmentSizes.front() = asyncDependencies.size();
1032 segmentSizes[7] = 0;
1033 segmentSizes[8] = 0;
1034 segmentSizes[9] = 0;
1036 segmentSizes[10] = hasDynamicSharedMemorySize ? 1 : 0;
1037 segmentSizes[11] = hasAsyncObject ? 1 : 0;
1038 result.addAttribute(LaunchOp::getOperandSegmentSizeAttr(),
1052 bool simplified =
false;
1053 auto constPropIdUses = [&](
Value id,
Value size) {
1057 if (
id.getUses().empty())
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());
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);
1099 return getBody().addArgument(type, loc);
1107 SymbolRefAttr kernelSymbol,
KernelDim3 gridSize,
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);
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);
1127 result.addOperands(asyncObject);
1129 Properties &prop =
result.getOrAddProperties<Properties>();
1130 prop.kernel = kernelSymbol;
1131 size_t segmentSizesLen = std::size(prop.operandSegmentSizes);
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;
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;
1152 std::optional<KernelDim3> clusterSize) {
1153 auto kernelModule = kernelFunc->getParentOfType<GPUModuleOp>();
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);
1162StringAttr LaunchFuncOp::getKernelModuleName() {
1166StringAttr LaunchFuncOp::getKernelName() {
1170unsigned LaunchFuncOp::getNumKernelOperands() {
1171 return getKernelOperands().size();
1174Value LaunchFuncOp::getKernelOperand(
unsigned i) {
1175 return getKernelOperands()[i];
1178KernelDim3 LaunchFuncOp::getGridSizeOperandValues() {
1179 auto operands = getOperands().drop_front(getAsyncDependencies().size());
1180 return KernelDim3{operands[0], operands[1], operands[2]};
1183KernelDim3 LaunchFuncOp::getBlockSizeOperandValues() {
1184 auto operands = getOperands().drop_front(getAsyncDependencies().size());
1185 return KernelDim3{operands[3], operands[4], operands[5]};
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]};
1195LogicalResult LaunchFuncOp::verify() {
1199 auto module = (*this)->getParentOfType<ModuleOp>();
1201 return emitOpError(
"expected to belong to a module");
1203 if (!module->getDiscardableAttrOfType<UnitAttr>(
1204 GPUDialect::getContainerModuleAttrName()))
1205 return emitOpError(
"expected the closest surrounding module to have the '" +
1206 GPUDialect::getContainerModuleAttrName() +
1209 if (hasClusterSize()) {
1210 if (getClusterSizeY().
getType() != getClusterSizeX().
getType() ||
1212 return emitOpError()
1213 <<
"expects types of the cluster dimensions must be the same";
1221 LaunchFuncOp launchOp = *
this;
1224 if (isa<GPUModuleOp>(table))
1229 if (!launchOp->getParentOp() ||
1230 launchOp->getParentOp()->getParentOp() != table)
1235 if (!launchOp.getKernelAttr())
1239 StringAttr kernelContainerName = launchOp.getKernelModuleName();
1242 if (!kernelContainer)
1244 <<
"kernel container '" << kernelContainerName.getValue()
1245 <<
"' is undefined";
1248 if (isa<BinaryOp>(kernelContainer))
1251 auto kernelModule = dyn_cast<GPUModuleOp>(kernelContainer);
1253 return launchOp.emitOpError()
1254 <<
"kernel module '" << kernelContainerName.getValue()
1255 <<
"' is undefined";
1259 kernelModule, launchOp.getKernelName());
1262 << launchOp.getKernel() <<
"' is undefined";
1263 auto kernelConvertedFunction = dyn_cast<FunctionOpInterface>(kernelFunc);
1264 if (!kernelConvertedFunction) {
1266 <<
"referenced kernel '" << launchOp.getKernel()
1267 <<
"' is not a function";
1268 diag.attachNote(kernelFunc->
getLoc()) <<
"see the kernel definition here";
1272 if (!GPUDialect::isKernel(kernelFunc))
1273 return launchOp.emitOpError(
"kernel function is missing the '")
1274 << GPUDialect::getKernelFuncAttrName() <<
"' attribute";
1279 auto kernelGPUFunction = dyn_cast<gpu::GPUFuncOp>(kernelFunc);
1280 if (!kernelGPUFunction)
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;
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";
1303 std::optional<OpAsmParser::UnresolvedOperand> clusterValue,
1304 Type &clusterXTy,
Type &clusterYTy,
Type &clusterZTy) {
1311 if (clusterValue.has_value()) {
1312 clusterXTy = clusterYTy = clusterZTy = dimTy;
1319 Type clusterYTy,
Type clusterZTy) {
1321 printer <<
": " << dimTy;
1331 auto parseElement = [&]() -> ParseResult {
1332 return failure(parser.
parseOperand(argNames.emplace_back()) ||
1337 parseElement,
" in argument list");
1342 if (operands.empty())
1345 llvm::interleaveComma(llvm::zip_equal(operands, types), printer,
1346 [&](
const auto &pair) {
1347 auto [operand, type] = pair;
1348 printer << operand <<
" : " << type;
1358 int32_t offset, int32_t width, ShuffleMode mode) {
1359 build(builder,
result, value,
1360 arith::ConstantOp::create(builder,
result.location,
1362 arith::ConstantOp::create(builder,
result.location,
1371LogicalResult RotateOp::verify() {
1372 uint32_t offset = getOffset();
1373 uint32_t width = getWidth();
1375 if (offset >= width) {
1376 return emitOpError() <<
"offset must be in the range [0, " << width <<
")";
1386LogicalResult BarrierOp::verify() {
1387 BarrierScope scope = getScope();
1389 if (getNamedBarrier() && scope != BarrierScope::Workgroup)
1390 return emitOpError(
"named barriers require workgroup scope");
1398 auto nextOp = dyn_cast_or_null<BarrierOp>(op->getNextNode());
1403 if (op.getScope() != nextOp.getScope())
1407 if (op.getNamedBarrier() != nextOp.getNamedBarrier())
1410 std::optional<ArrayAttr> thisMemfence = op.getAddressSpaces();
1411 std::optional<ArrayAttr> nextMemfence = nextOp.getAddressSpaces();
1415 if (!nextMemfence) {
1416 op.removeAddressSpacesAttr();
1420 if (*thisMemfence == *nextMemfence) {
1424 llvm::SmallSetVector<Attribute, 4> mergedSpaces;
1426 mergedSpaces.insert(attr);
1428 mergedSpaces.insert(attr);
1429 op.setAddressSpacesAttr(rewriter.
getArrayAttr(mergedSpaces.takeVector()));
1444 std::optional<AddressSpace> addressSpace) {
1448 AddressSpaceAttr::get(odsBuilder.
getContext(), addressSpace.value()));
1450 odsBuilder, odsState, addressSpacesAttr,
Value{},
1451 BarrierScopeAttr::get(odsBuilder.
getContext(), BarrierScope::Workgroup));
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);
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);
1486 return getBody().addArgument(type, loc);
1490 StringRef name, FunctionType type,
1496 result.getOrAddProperties<Properties>().sym_name =
1498 result.addAttribute(getFunctionTypeAttrName(
result.name),
1499 TypeAttr::get(type));
1500 result.addAttribute(getWorkgroupAttributionsAttrName(
result.name),
1502 result.addAttributes(attrs);
1507 for (
Type argTy : type.getInputs())
1509 for (
Type argTy : workgroupAttributions)
1511 for (
Type argTy : privateAttributions)
1530 size_t existingArgs = args.size();
1537 bool hadAttrs = llvm::any_of(
ArrayRef(args).drop_front(existingArgs),
1542 attributionAttrs =
nullptr;
1548 for (
const auto &argument :
ArrayRef(args).drop_front(existingArgs)) {
1549 if (!argument.attrs)
1552 attributionAttrsVec.push_back(argument.attrs);
1554 attributionAttrs = builder.
getArrayAttr(attributionAttrsVec);
1570 StringAttr nameAttr;
1577 parser,
false, entryArgs, isVariadic, resultTypes,
1581 if (!entryArgs.empty() && entryArgs[0].ssaName.name.empty())
1582 return parser.
emitError(signatureLocation)
1583 <<
"gpu.func requires named arguments";
1590 for (
auto &arg : entryArgs)
1591 argTypes.push_back(arg.
type);
1593 result.addAttribute(getFunctionTypeAttrName(
result.name),
1594 TypeAttr::get(type));
1597 builder,
result, entryArgs, resultAttrs, getArgAttrsAttrName(
result.name),
1598 getResAttrsAttrName(
result.name));
1603 entryArgs, workgroupAttributionAttrs)))
1608 unsigned numWorkgroupAttrs = entryArgs.size() - type.getNumInputs();
1609 if (numWorkgroupAttrs != 0)
1611 GPUFuncOp::getWorkgroupAttributionsAttrName(
result.name),
1613 if (workgroupAttributionAttrs)
1614 result.addAttribute(GPUFuncOp::getWorkgroupAttribAttrsAttrName(
result.name),
1615 workgroupAttributionAttrs);
1620 entryArgs, privateAttributionAttrs)))
1622 if (privateAttributionAttrs)
1623 result.addAttribute(GPUFuncOp::getPrivateAttribAttrsAttrName(
result.name),
1624 privateAttributionAttrs);
1628 result.addAttribute(GPUFuncOp::getKernelAttrName(
result.name),
1637 auto *body =
result.addRegion();
1645 FunctionType type = getFunctionType();
1651 getWorkgroupAttribAttrs().value_or(
nullptr));
1653 getPrivateAttribAttrs().value_or(
nullptr));
1655 p <<
' ' << getKernelKeyword();
1659 {getWorkgroupAttributionsAttrName(), getKernelAttrName(),
1660 GPUDialect::getKernelFuncAttrName(), getFunctionTypeAttrName(),
1661 getArgAttrsAttrName(), getResAttrsAttrName(),
1662 getWorkgroupAttribAttrsAttrName(), getPrivateAttribAttrsAttrName()});
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]);
1677DictionaryAttr GPUFuncOp::getworkgroupAttributionAttrs(
unsigned index) {
1681DictionaryAttr GPUFuncOp::getPrivateAttributionAttrs(
unsigned index) {
1686 DictionaryAttr value, StringAttr attrName) {
1688 ArrayAttr allAttrs = attrName == op.getWorkgroupAttribAttrsAttrName()
1689 ? op.getWorkgroupAttribAttrsAttr()
1690 : op.getPrivateAttribAttrsAttr();
1693 elements.append(allAttrs.begin(), allAttrs.end());
1694 while (elements.size() <=
index)
1695 elements.push_back(DictionaryAttr::get(ctx));
1697 elements[
index] = DictionaryAttr::get(ctx);
1699 elements[
index] = value;
1700 ArrayAttr newValue = ArrayAttr::get(ctx, elements);
1701 if (attrName == op.getWorkgroupAttribAttrsAttrName())
1702 op.setWorkgroupAttribAttrsAttr(newValue);
1704 op.setPrivateAttribAttrsAttr(newValue);
1707void GPUFuncOp::setworkgroupAttributionAttrs(
unsigned index,
1708 DictionaryAttr value) {
1712void GPUFuncOp::setPrivateAttributionAttrs(
unsigned int index,
1713 DictionaryAttr value) {
1718 StringAttr name, StringAttr attrsName) {
1722 return dict.get(name);
1727 assert(
index < getNumWorkgroupAttributions() &&
1728 "index must map to a workgroup attribution");
1730 getWorkgroupAttribAttrsAttrName());
1735 assert(
index < getNumPrivateAttributions() &&
1736 "index must map to a private attribution");
1738 getPrivateAttribAttrsAttrName());
1742 Attribute value, StringAttr attrsName) {
1747 elems.append(oldDict.getValue().begin(), oldDict.getValue().end());
1750 bool mustSort =
true;
1751 for (
unsigned i = 0, e = elems.size(); i < e; ++i) {
1752 if (elems[i].getName() == name) {
1755 std::swap(elems[i], elems[elems.size() - 1]);
1767 elems.emplace_back(name, value);
1770 DictionaryAttr::sortInPlace(elems);
1772 auto newDict = DictionaryAttr::getWithSorted(ctx, elems);
1776void GPUFuncOp::setWorkgroupAttributionAttr(
unsigned index, StringAttr name,
1778 assert(
index < getNumWorkgroupAttributions() &&
1779 "index must map to a workgroup attribution");
1781 getWorkgroupAttribAttrsAttrName());
1784void GPUFuncOp::setPrivateAttributionAttr(
unsigned index, StringAttr name,
1786 assert(
index < getNumPrivateAttributions() &&
1787 "index must map to a private attribution");
1789 getPrivateAttribAttrsAttrName());
1792LogicalResult GPUFuncOp::verifyType() {
1793 if (isKernel() && getFunctionType().getNumResults() != 0)
1794 return emitOpError() <<
"expected void return type for kernel function";
1800LogicalResult GPUFuncOp::verifyBody() {
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";
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 "
1821 GPUDialect::getWorkgroupAddressSpace())) ||
1823 GPUDialect::getPrivateAddressSpace())))
1833LogicalResult gpu::ReturnOp::verify() {
1834 GPUFuncOp function = (*this)->getParentOfType<GPUFuncOp>();
1836 FunctionType funType = function.getFunctionType();
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");
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();
1874 Properties &properties = getProperties();
1886 result.addRegion()->emplaceBlock();
1887 Properties &props =
result.getOrAddProperties<Properties>();
1889 props.targets = targets;
1891 props.offloadingHandler = offloadingHandler;
1897 build(builder,
result, name,
1903 if (
ArrayAttr targets = getTargetsAttr())
1904 return llvm::count(targets.getValue(),
target);
1909 ArrayAttr &targetsAttr = getProperties().targets;
1911 targetsAttr = ArrayAttr::get(
getContext(), targetsVector);
1914LogicalResult GPUModuleOp::verify() {
1915 auto targets = getTargetsAttr();
1920 for (
auto target : targets) {
1921 if (
auto verifyTargetAttr =
1922 llvm::dyn_cast<TargetAttrVerifyInterface>(
target)) {
1923 if (verifyTargetAttr.verifyTarget(getOperation()).
failed())
1935 auto &properties =
result.getOrAddProperties<Properties>();
1938 properties.objects = objects;
1939 if (offloadingHandler)
1940 properties.offloadingHandler = offloadingHandler;
1942 properties.offloadingHandler = builder.
getAttr<SelectObjectAttr>(
nullptr);
1947 build(builder,
result, name, offloadingHandler,
1959 if (!offloadingHandler)
1966 if (offloadingHandler != SelectObjectAttr::get(op->
getContext(),
nullptr))
1967 printer <<
'<' << offloadingHandler <<
'>';
1974LogicalResult MemcpyOp::verify() {
1975 auto srcType = getSrc().getType();
1976 auto dstType = getDst().getType();
1979 return emitOpError(
"arguments have incompatible element type");
1982 return emitOpError(
"arguments have incompatible shape");
1991 using OpRewritePattern<MemcpyOp>::OpRewritePattern;
1993 LogicalResult matchAndRewrite(MemcpyOp op,
1994 PatternRewriter &rewriter)
const override {
1995 Value dest = op.getDst();
2004 if (llvm::any_of(dest.
getUsers(), [op, dest](Operation *user) {
2005 return user != op &&
2006 !hasSingleEffect<MemoryEffects::Free>(user, dest);
2012 if (op.getAsyncDependencies().size() > 1 ||
2013 ((op.getAsyncDependencies().empty() && op.getAsyncToken()) ||
2014 (!op.getAsyncDependencies().empty() && !op.getAsyncToken())))
2016 rewriter.
replaceOp(op, op.getAsyncDependencies());
2024 results.
add<EraseTrivialCopyOp>(context);
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);
2038 if (!srcMemrefType.isLastDimUnitStride())
2040 "expected source memref most minor dim must have unit stride");
2042 if (operand !=
"AOp" && operand !=
"BOp" && operand !=
"COp")
2043 return emitError(
"only AOp, BOp and COp can be loaded");
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);
2058 if (!dstMemrefType.isLastDimUnitStride())
2060 "expected destination memref most minor dim must have unit stride");
2062 if (srcMatrixType.getOperand() !=
"COp")
2064 "expected the operand matrix being stored to have 'COp' operand type");
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()));
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");
2085 aShape = opTypes[
A].getShape();
2086 bShape = opTypes[
B].getShape();
2087 cShape = opTypes[
C].getShape();
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");
2096LogicalResult MemcpyOp::fold(FoldAdaptor adaptor,
2101LogicalResult MemsetOp::fold(FoldAdaptor adaptor,
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;
2124 if (llvm::none_of(op.getAsyncDependencies(), predicate))
2126 SmallVector<Value> validOperands;
2127 for (Value operand : op->getOperands()) {
2128 if (predicate(operand))
2130 validOperands.push_back(operand);
2132 rewriter.
modifyOpInPlace(op, [&]() { op->setOperands(validOperands); });
2148 LogicalResult matchAndRewrite(WaitOp op,
2149 PatternRewriter &rewriter)
const final {
2152 if (op.getAsyncDependencies().empty() && !op.getAsyncToken()) {
2157 if (llvm::hasSingleElement(op.getAsyncDependencies()) &&
2158 op.getAsyncToken()) {
2159 rewriter.
replaceOp(op, op.getAsyncDependencies());
2163 if (op.getAsyncToken() && op.getAsyncToken().use_empty()) {
2174 results.
add<EraseRedundantGpuWaitOpPairs, SimplifyGpuWaitOp>(context);
2181LogicalResult AllocOp::verify() {
2182 auto memRefType = llvm::cast<MemRefType>(getMemref().
getType());
2188 unsigned numSymbols = 0;
2189 if (!memRefType.getLayout().isIdentity())
2190 numSymbols = memRefType.getLayout().getAffineMap().getNumSymbols();
2191 if (getSymbolOperands().size() != numSymbols) {
2193 "symbol operand count does not equal memref symbol count");
2203 using OpRewritePattern<memref::DimOp>::OpRewritePattern;
2205 LogicalResult matchAndRewrite(memref::DimOp dimOp,
2206 PatternRewriter &rewriter)
const override {
2207 std::optional<int64_t> index = dimOp.getConstantIndex();
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))
2217 auto alloc = dimOp.getSource().getDefiningOp<AllocOp>();
2221 Value substituteOp = *(alloc.getDynamicSizes().begin() +
2222 memrefType.getDynamicDimIndex(indexVal));
2223 rewriter.
replaceOp(dimOp, substituteOp);
2231 results.
add<SimplifyDimOfAllocOp>(context);
2238LogicalResult gpu::DynamicSharedMemoryOp::verify() {
2239 if (!getOperation()->getParentWithTrait<OpTrait::SymbolTable>())
2240 return emitOpError() <<
"must be inside an op with symbol table";
2242 MemRefType memrefType = getResultMemref().getType();
2244 if (!GPUDialect::hasWorkgroupMemoryAddressSpace(memrefType)) {
2245 return emitOpError() <<
"address space must be "
2246 << gpu::AddressSpaceAttr::getMnemonic() <<
"<"
2247 << stringifyEnum(gpu::AddressSpace::Workgroup) <<
">";
2249 if (memrefType.hasStaticShape()) {
2250 return emitOpError() <<
"result memref type must be memref<?xi8, "
2251 "#gpu.address_space<workgroup>>";
2261 p <<
"(" << getLaneid() <<
")";
2264 p <<
"[" << getWarpSize() <<
"]";
2266 if (!getArgs().empty())
2267 p <<
" args(" << getArgs() <<
" : " << getArgs().getTypes() <<
")";
2268 if (!getResults().empty())
2269 p <<
" -> (" << getResults().getTypes() <<
')';
2273 !getResults().empty());
2275 getOperation()->getDiscardableAttrDictionary().getValue(), coreAttr);
2278ParseResult WarpExecuteOnLane0Op::parse(
OpAsmParser &parser,
2281 result.regions.reserve(1);
2304 llvm::SMLoc inputsOperandsLoc;
2316 if (parser.
resolveOperands(inputsOperands, inputTypes, inputsOperandsLoc,
2327 WarpExecuteOnLane0Op::ensureTerminator(*warpRegion, builder,
result.location);
2335void WarpExecuteOnLane0Op::getSuccessorRegions(
2352 build(builder,
result, resultTypes, laneId, warpSize,
2360 result.addOperands(laneId);
2361 result.addAttribute(getAttributeNames()[0],
2363 result.addTypes(resultTypes);
2364 result.addOperands(args);
2365 assert(args.size() == blockArgTypes.size());
2369 for (
auto [type, arg] : llvm::zip_equal(blockArgTypes, args))
2378 if (expanded == distributed)
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())
2387 "expected distributed vectors to have same rank and element type.");
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);
2395 if (eDim % dDim != 0)
2397 <<
"expected expanded vector dimension #" << i <<
" (" << eDim
2398 <<
") to be a multipler of the distributed vector dimension ("
2400 scales[i] = eDim / dDim;
2402 if (llvm::product_of(scales) != warpSize)
2404 <<
"incompatible distribution dimensions from " << expandedVecType
2405 <<
" to " << distributedVecType <<
" with warp size = " << warpSize;
2410LogicalResult WarpExecuteOnLane0Op::verify() {
2411 if (getArgs().size() != getWarpRegion().getNumArguments())
2413 "expected same number op arguments and block arguments.");
2414 auto yield = dyn_cast<gpu::YieldOp>(getBody()->getTerminator());
2416 return emitOpError(
"expected body to be terminated with 'gpu.yield'");
2417 if (yield.getNumOperands() != getNumResults())
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())) {
2424 warpSize, getOperation())))
2427 for (
auto [yieldOperand,
result] :
2428 llvm::zip_equal(yield.getOperands(), getResults())) {
2430 warpSize, getOperation())))
2435bool WarpExecuteOnLane0Op::areTypesCompatible(
Type lhs,
Type rhs) {
2440gpu::YieldOp WarpExecuteOnLane0Op::getTerminator() {
2441 return cast<gpu::YieldOp>(getBody()->getTerminator());
2448void gpu::SubgroupBroadcastOp::inferResultRanges(
2450 setResultRange(getResult(), argRanges.front());
2454 switch (getBroadcastType()) {
2455 case BroadcastType::first_active_lane:
2459 case BroadcastType::specific_lane:
2463 llvm_unreachable(
"Unknown BroadcastType");
2466LogicalResult gpu::SubgroupBroadcastOp::verify() {
2467 switch (getBroadcastType()) {
2468 case BroadcastType::first_active_lane:
2470 return emitOpError()
2471 <<
"lane can only be specified for `specific_lane` broadcast";
2473 case BroadcastType::specific_lane:
2475 return emitOpError()
2476 <<
"lane must be specified for `specific_lane` broadcast";
2479 llvm_unreachable(
"Unknown BroadcastType");
2482OpFoldResult gpu::SubgroupBroadcastOp::fold(FoldAdaptor ) {
2484 if (
auto prev = getSrc().getDefiningOp<SubgroupBroadcastOp>())
2485 return prev.getResult();
2496#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
2498#define GET_OP_CLASSES
2499#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
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.
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.
This class represents an argument of a Block.
Block represents an ordered list of Operations.
BlockArgument addArgument(Type type, Location loc)
Add one value to the argument list.
This class is a general helper class for creating context-global objects like types,...
IntegerAttr getI32IntegerAttr(int32_t value)
DenseI32ArrayAttr getDenseI32ArrayAttr(ArrayRef< int32_t > values)
FunctionType getFunctionType(TypeRange inputs, TypeRange results)
IntegerAttr getI64IntegerAttr(int64_t value)
Ty getType(Args &&...args)
Get or construct an instance of the type Ty with provided arguments.
StringAttr getStringAttr(const Twine &bytes)
ArrayAttr getArrayAttr(ArrayRef< Attribute > value)
MLIRContext * getContext() const
DictionaryAttr getDictionaryAttr(ArrayRef< NamedAttribute > value)
NamedAttribute getNamedAttr(StringRef name, Attribute val)
Attr getAttr(Args &&...args)
Get or construct an instance of the attribute Attr with provided arguments.
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...
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...
MLIRContext is the top-level object for a collection of MLIR operations.
NamedAttribute represents a combination of a name and an Attribute value.
StringAttr getName() const
Return the name of the attribute.
Attribute getValue() const
Return the value of the attribute.
The OpAsmParser has methods for interacting with the asm parser: parsing things from it,...
virtual size_t getNumResults() const =0
Return the number of declared SSA results.
virtual ParseResult parseRegion(Region ®ion, 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.
This class helps build Operations.
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.
void setInsertionPointToStart(Block *block)
Sets the insertion point to the start of the specified block.
This class represents a single result from folding an operation.
static StringRef getOperandSegmentSizeAttr()
This class implements the operand iterators for the Operation class.
Operation is the basic unit of execution within MLIR.
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.
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.
Operation * getParentOp()
Returns the closest surrounding operation that contains this operation or nullptr if this is a top-le...
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.
MLIRContext * getContext()
Return the context this operation is associated with.
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.
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.
This class provides an abstraction over the various different ranges of value types.
Instances of the Type class are uniqued, have an immutable identifier and an optional mutable compone...
This class provides an abstraction over the different types of ranges over Values.
This class represents an instance of an SSA value in the MLIR system, representing a computable value...
Type getType() const
Return the type of this value.
user_range getUsers() const
Location getLoc() const
Return the location of this value.
Operation * getDefiningOp() const
If this value is the result of an operation, return the operation that defines it.
static ConstantIndexOp create(OpBuilder &builder, Location location, int64_t value)
static ConcreteType get(MLIRContext *ctx, Args &&...args)
MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply accumulate operations.
ArrayRef< int64_t > getShape() const
Get shape of the matrix.
Type getElementType() const
Get elementType of a single element.
StringRef getOperand() const
The general form of operation this type supports is given by the equation C += A*B.
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.
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.
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.
void addAsyncDependency(Operation *op, Value token)
LogicalResult foldMemRefCast(Operation *op, Value inner=nullptr)
This is a common utility used for patterns of the form "someop(memref.cast) -> someop".
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.
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.
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.
llvm::TypeSwitch< T, ResultT > TypeSwitch
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
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
UnresolvedOperand ssaName
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 ....