MLIR 24.0.0git
Utils.cpp
Go to the documentation of this file.
1//===- Utils.cpp - Utils for GPU transform ops ----------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8
10
20#include "mlir/IR/AffineExpr.h"
21#include "mlir/IR/Builders.h"
23#include "mlir/IR/MLIRContext.h"
25#include "mlir/IR/Value.h"
26#include "mlir/IR/Visitors.h"
27#include "mlir/Support/LLVM.h"
28#include "llvm/ADT/STLExtras.h"
29#include "llvm/ADT/SmallVector.h"
30#include "llvm/Support/DebugLog.h"
31#include "llvm/Support/InterleavedRange.h"
32
33#include <utility>
34
35using namespace mlir;
36using namespace mlir::gpu;
37using namespace mlir::transform;
38using namespace mlir::transform::gpu;
39
40#define DEBUG_TYPE "gpu-transforms"
41
42/// Build predicates to filter execution by only the activeIds. Along each
43/// dimension, 3 cases appear:
44/// 1. activeMappingSize > availableMappingSize: this is an unsupported case
45/// as this requires additional looping. An error message is produced to
46/// advise the user to tile more or to use more threads.
47/// 2. activeMappingSize == availableMappingSize: no predication is needed.
48/// 3. activeMappingSize < availableMappingSize: only a subset of threads
49/// should be active and we produce the boolean `id < activeMappingSize`
50/// for further use in building predicated execution.
51static FailureOr<SmallVector<Value>>
53 ArrayRef<int64_t> activeMappingSizes,
54 ArrayRef<int64_t> availableMappingSizes,
55 std::string &errorMsg) {
56 LDBG() << "----activeMappingSizes: " << llvm::interleaved(activeMappingSizes);
57 LDBG() << "----availableMappingSizes: "
58 << llvm::interleaved(availableMappingSizes);
59
60 SmallVector<Value> predicateOps;
61 for (auto [activeId, activeMappingSize, availableMappingSize] :
62 llvm::zip_equal(activeIds, activeMappingSizes, availableMappingSizes)) {
63 if (activeMappingSize > availableMappingSize) {
64 errorMsg = "Trying to map to fewer GPU threads than loop iterations but "
65 "overprovisioning is not yet supported. Try additional tiling "
66 "before mapping or map to more threads.";
67 return failure();
68 }
69 if (activeMappingSize == availableMappingSize)
70 continue;
71 Value idx =
72 arith::ConstantIndexOp::create(rewriter, loc, activeMappingSize);
73 Value pred = arith::CmpIOp::create(rewriter, loc, arith::CmpIPredicate::ult,
74 activeId, idx);
75 predicateOps.push_back(pred);
76 }
77 return predicateOps;
78}
79
80/// Return a flattened thread id for the workgroup with given sizes.
81template <typename ThreadOrBlockIdOp>
83 ArrayRef<OpFoldResult> originalBasisOfr) {
84 LDBG() << "----buildLinearId with originalBasisOfr: "
85 << llvm::interleaved(originalBasisOfr);
86 assert(originalBasisOfr.size() == 3 && "expected 3 sizes");
87 IndexType indexType = rewriter.getIndexType();
88 AffineExpr tx, ty, tz, bdx, bdy;
89 bindDims(rewriter.getContext(), tx, ty, tz);
90 bindSymbols(rewriter.getContext(), bdx, bdy);
92 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::x)
93 .getResult(),
94 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::y)
95 .getResult(),
96 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::z)
97 .getResult(),
98 originalBasisOfr[0], originalBasisOfr[1]};
100 rewriter, loc, tx + ty * bdx + tz * bdx * bdy, vals);
101 return getValueOrCreateConstantIndexOp(rewriter, loc, ofr);
102}
103
104/// Create a linear id builder that takes the `originalBasisOfr` and decompose
105/// it in the basis of `forallMappingSizes`. The linear id builder returns an
106/// n-D vector of ids for indexing and 1-D size + id for predicate generation.
107template <typename ThreadOrBlockIdOp>
110 DeviceMaskingAttrInterface mask = nullptr) {
111 auto res = [multiplicity, mask](RewriterBase &rewriter, Location loc,
112 ArrayRef<int64_t> forallMappingSizes,
113 ArrayRef<int64_t> originalBasis) {
114 // 0. Early-exit mask case.
115 if (mask) {
116 if (computeProduct(originalBasis) >
117 mask.getMaxNumPhysicalIds() * multiplicity) {
118 return IdBuilderResult{
119 /*errorMsg=*/std::string(
120 "mask representation too short to capture all physical ids: ") +
121 std::to_string(mask.getMaxNumPhysicalIds()),
122 /*mappingIdOps=*/{},
123 /*predicateOps=*/{}};
124 }
125 }
126
127 // 1. Compute linearId.
128 SmallVector<OpFoldResult> originalBasisOfr =
129 getAsIndexOpFoldResult(rewriter.getContext(), originalBasis);
130 Value physicalLinearId =
131 buildLinearId<ThreadOrBlockIdOp>(rewriter, loc, originalBasisOfr);
132
133 // 2. Compute scaledLinearId.
134 AffineExpr d0 = getAffineDimExpr(0, rewriter.getContext());
136 rewriter, loc, d0.floorDiv(multiplicity), {physicalLinearId});
137
138 // 2.b. Adjust with mask if needed.
139 Value scaledLinearIdI64;
140 Value scaledLinearId =
141 getValueOrCreateConstantIndexOp(rewriter, loc, scaledLinearIdOfr);
142 if (mask) {
143 scaledLinearId =
144 getValueOrCreateConstantIndexOp(rewriter, loc, scaledLinearIdOfr);
145 scaledLinearIdI64 = arith::IndexCastUIOp::create(
146 rewriter, loc, rewriter.getI64Type(), scaledLinearId);
147 Value logicalLinearIdI64 =
148 mask.createLogicalLinearMappingId(rewriter, scaledLinearIdI64);
149 scaledLinearId = arith::IndexCastUIOp::create(
150 rewriter, loc, rewriter.getIndexType(), logicalLinearIdI64);
151 LDBG() << "------adjusting linearId with mask: " << scaledLinearId;
152 }
153
154 // 3. Compute remapped indices.
156 // Sizes in [0 .. n] -> [n .. 0] order to properly compute strides in
157 // "row-major" order.
158 SmallVector<int64_t> reverseBasisSizes(llvm::reverse(forallMappingSizes));
159 SmallVector<int64_t> strides = computeStrides(reverseBasisSizes);
160 SmallVector<AffineExpr> delinearizingExprs = delinearize(d0, strides);
161 // Reverse back to be in [0 .. n] order.
162 for (AffineExpr e : llvm::reverse(delinearizingExprs)) {
163 ids.push_back(
164 affine::makeComposedAffineApply(rewriter, loc, e, {scaledLinearId}));
165 }
166
167 std::string errorMsg;
168 SmallVector<Value> predicateOps;
169 // 4. If mask present, it takes precedence to determine predication.
170 if (mask) {
171 Value isActiveIdPredicate =
172 mask.createIsActiveIdPredicate(rewriter, scaledLinearIdI64);
173 LDBG() << "------adjusting predicate with mask: " << isActiveIdPredicate;
174 predicateOps.push_back(isActiveIdPredicate);
175 } else {
176 // 4.b. Otherwise, handle predicates using physicalLinearId.
177 FailureOr<SmallVector<Value>> maybePredicateOps =
178 buildPredicates(rewriter, loc, physicalLinearId,
179 computeProduct(forallMappingSizes) * multiplicity,
180 computeProduct(originalBasis), errorMsg);
181 if (succeeded(maybePredicateOps))
182 predicateOps = std::move(*maybePredicateOps);
183 }
184
185 return IdBuilderResult{/*errorMsg=*/std::move(errorMsg),
186 /*mappingIdOps=*/std::move(ids),
187 /*predicateOps=*/std::move(predicateOps)};
188 };
189
190 return res;
191}
192
193/// Create a simple 3-D id builder that takes the `originalBasisOfr`
194/// The 3-D id builder returns a 3-D vector of ids for indexing and 3-D sizes
195/// + ids for predicate generation.
196template <typename ThreadOrBlockIdOp>
198 auto res = [multiplicity](RewriterBase &rewriter, Location loc,
199 ArrayRef<int64_t> forallMappingSizes,
200 ArrayRef<int64_t> originalBasis) {
201 IndexType indexType = rewriter.getIndexType();
203 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::x),
204 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::y),
205 ThreadOrBlockIdOp::create(rewriter, loc, indexType, Dimension::z)};
206 // In the 3-D mapping case, scale the first dimension by the multiplicity.
207 SmallVector<Value> scaledIds = ids;
208 AffineExpr d0 = getAffineDimExpr(0, rewriter.getContext());
209 scaledIds[0] = cast<Value>(affine::makeComposedFoldedAffineApply(
210 rewriter, loc, d0.floorDiv(multiplicity), {scaledIds[0]}));
211 // In the 3-D mapping case, unscale the first dimension by the multiplicity.
212 SmallVector<int64_t> forallMappingSizeInOriginalBasis(forallMappingSizes);
213 forallMappingSizeInOriginalBasis[0] *= multiplicity;
214
215 std::string errorMsg;
216 SmallVector<Value> predicateOps;
217 FailureOr<SmallVector<Value>> maybePredicateOps =
218 buildPredicates(rewriter, loc, ids, forallMappingSizeInOriginalBasis,
219 originalBasis, errorMsg);
220 if (succeeded(maybePredicateOps))
221 predicateOps = std::move(*maybePredicateOps);
222
223 return IdBuilderResult{/*errorMsg=*/std::move(errorMsg),
224 /*mappingIdOps=*/std::move(scaledIds),
225 /*predicateOps=*/std::move(predicateOps)};
226 };
227 return res;
228}
229
230/// Create a lane id builder that takes the `originalBasis` and decompose
231/// it in the basis of `forallMappingSizes`. The linear id builder returns an
232/// n-D vector of ids for indexing and 1-D size + id for predicate generation.
234 auto res = [warpSize](RewriterBase &rewriter, Location loc,
235 ArrayRef<int64_t> forallMappingSizes,
236 ArrayRef<int64_t> originalBasis) {
237 // 1. Compute linearId.
238 SmallVector<OpFoldResult> originalBasisOfr =
239 getAsIndexOpFoldResult(rewriter.getContext(), originalBasis);
240 Value physicalLinearId =
241 buildLinearId<ThreadIdOp>(rewriter, loc, originalBasisOfr);
242
243 // 2. Compute laneId.
244 AffineExpr d0 = getAffineDimExpr(0, rewriter.getContext());
246 rewriter, loc, d0 % warpSize, {physicalLinearId});
247
248 // 3. Compute remapped indices.
250 // Sizes in [0 .. n] -> [n .. 0] order to properly compute strides in
251 // "row-major" order.
252 SmallVector<int64_t> reverseBasisSizes(llvm::reverse(forallMappingSizes));
253 SmallVector<int64_t> strides = computeStrides(reverseBasisSizes);
254 SmallVector<AffineExpr> delinearizingExprs = delinearize(d0, strides);
255 // Reverse back to be in [0 .. n] order.
256 for (AffineExpr e : llvm::reverse(delinearizingExprs)) {
257 ids.push_back(
258 affine::makeComposedAffineApply(rewriter, loc, e, {laneId}));
259 }
260
261 // 4. Handle predicates using laneId.
262 std::string errorMsg;
263 SmallVector<Value> predicateOps;
264 FailureOr<SmallVector<Value>> maybePredicateOps = buildPredicates(
265 rewriter, loc, cast<Value>(laneId), computeProduct(forallMappingSizes),
266 computeProduct(originalBasis), errorMsg);
267 if (succeeded(maybePredicateOps))
268 predicateOps = std::move(*maybePredicateOps);
269
270 return IdBuilderResult{/*errorMsg=*/std::move(errorMsg),
271 /*mappingIdOps=*/std::move(ids),
272 /*predicateOps=*/std::move(predicateOps)};
273 };
274
275 return res;
276}
277
278namespace mlir {
279namespace transform {
280namespace gpu {
281
282GpuIdBuilder::GpuIdBuilder(MLIRContext *ctx, bool useLinearMapping,
283 const MappingIdBuilderFnType &fn)
285 if (useLinearMapping) {
286 for (uint64_t d = static_cast<uint64_t>(MappingId::LinearDim0),
287 e = getMaxEnumValForMappingId();
288 d <= e; ++d)
289 mappingAttributes.push_back(fn(ctx, symbolizeMappingId(d).value()));
290 } else {
291 for (uint64_t d = static_cast<uint64_t>(MappingId::DimX),
292 e = static_cast<uint64_t>(MappingId::DimZ);
293 d <= e; ++d)
294 mappingAttributes.push_back(fn(ctx, symbolizeMappingId(d).value()));
295 }
296}
297
299 DeviceMaskingAttrInterface mask)
300 : GpuIdBuilder(ctx, useLinearMapping, [](MLIRContext *ctx, MappingId id) {
301 return GPUBlockMappingAttr::get(ctx, id);
302 }) {
303 assert((!mask || useLinearMapping) && "mask requires linear mapping");
304 idBuilder = useLinearMapping
305 ? commonLinearIdBuilderFn<BlockIdOp>(/*multiplicity=*/1, mask)
306 : common3DIdBuilderFn<BlockIdOp>(/*multiplicity=*/1);
307}
308
310 bool useLinearMapping,
311 DeviceMaskingAttrInterface mask)
312 : GpuIdBuilder(ctx, useLinearMapping,
313 [](MLIRContext *ctx, MappingId id) {
314 return GPUWarpgroupMappingAttr::get(ctx, id);
315 }),
316 warpSize(warpSize) {
317 assert((!mask || useLinearMapping) && "mask requires linear mapping");
318 idBuilder = useLinearMapping
320 /*multiplicity=*/kNumWarpsPerGroup * warpSize, mask)
321 : common3DIdBuilderFn<ThreadIdOp>(
322 /*multiplicity=*/kNumWarpsPerGroup * warpSize);
323}
324
326 bool useLinearMapping,
327 DeviceMaskingAttrInterface mask)
328 : GpuIdBuilder(ctx, useLinearMapping,
329 [](MLIRContext *ctx, MappingId id) {
330 return GPUWarpMappingAttr::get(ctx, id);
331 }),
332 warpSize(warpSize) {
333 assert((!mask || useLinearMapping) && "mask requires linear mapping");
334 idBuilder = useLinearMapping
336 /*multiplicity=*/warpSize, mask)
337 : common3DIdBuilderFn<ThreadIdOp>(/*multiplicity=*/warpSize);
338}
339
341 DeviceMaskingAttrInterface mask)
342 : GpuIdBuilder(ctx, useLinearMapping, [](MLIRContext *ctx, MappingId id) {
343 return GPUThreadMappingAttr::get(ctx, id);
344 }) {
345 idBuilder =
346 useLinearMapping
347 ? commonLinearIdBuilderFn<ThreadIdOp>(/*multiplicity=*/1, mask)
348 : common3DIdBuilderFn<ThreadIdOp>(/*multiplicity=*/1);
349}
350
352 bool unused, DeviceMaskingAttrInterface mask)
353 : GpuIdBuilder(ctx, /*useLinearMapping=*/true,
354 [](MLIRContext *ctx, MappingId id) {
355 return GPULaneMappingAttr::get(ctx, id);
356 }),
357 warpSize(warpSize) {
358 assert(!mask && "mask NYI for lanes, unclear it should be at all");
359 idBuilder = laneIdBuilderFn(/*periodicity=*/warpSize);
360}
361
362DiagnosedSilenceableFailure checkGpuLimits(TransformOpInterface transformOp,
363 std::optional<int64_t> gridDimX,
364 std::optional<int64_t> gridDimY,
365 std::optional<int64_t> gridDimZ,
366 std::optional<int64_t> blockDimX,
367 std::optional<int64_t> blockDimY,
368 std::optional<int64_t> blockDimZ) {
369
370 // TODO: pass a configuration object to set the limits properly.
371
372 if ((blockDimX.value_or(1) * blockDimY.value_or(1) * blockDimZ.value_or(1)) >
374 (gridDimX.value_or(1) * gridDimY.value_or(1) * gridDimZ.value_or(1)) >
376 blockDimX.value_or(1) > kMaxBlockdimx ||
377 blockDimY.value_or(1) > kMaxBlockdimy ||
378 blockDimZ.value_or(1) > kMaxBlockdimz ||
379 gridDimY.value_or(1) > kMaxGriddimy ||
380 gridDimZ.value_or(1) > kMaxGriddimz ||
381 gridDimX.value_or(1) > kMaxGriddimx) {
382 return transformOp.emitSilenceableError()
383 << "Trying to launch a GPU kernel with grid_dims = ("
384 << gridDimX.value_or(1) << ", " << gridDimY.value_or(1) << ", "
385 << gridDimZ.value_or(1) << ") block_dims = ("
386 << blockDimX.value_or(1) << ", " << blockDimY.value_or(1) << ", "
387 << blockDimZ.value_or(1) << "). It is larger than the limits.";
388 }
390}
391
393 RewriterBase &rewriter, Location loc, TransformOpInterface transformOp,
394 LaunchOp &launchOp, std::optional<int64_t> gridDimX,
395 std::optional<int64_t> gridDimY, std::optional<int64_t> gridDimZ,
396 std::optional<int64_t> blockDimX, std::optional<int64_t> blockDimY,
397 std::optional<int64_t> blockDimZ) {
399 checkGpuLimits(transformOp, gridDimX, gridDimY, gridDimZ, blockDimX,
400 blockDimY, blockDimZ);
401 if (!diag.succeeded())
402 return diag;
403
404 auto createConst = [&](int dim) {
405 return arith::ConstantIndexOp::create(rewriter, loc, dim);
406 };
407 OpBuilder::InsertionGuard guard(rewriter);
408 Value one = createConst(1);
409 Value gridSizeX = gridDimX.has_value() ? createConst(gridDimX.value()) : one;
410 Value gridSizeY = gridDimY.has_value() ? createConst(gridDimY.value()) : one;
411 Value gridSizeZ = gridDimZ.has_value() ? createConst(gridDimZ.value()) : one;
412 Value blkSizeX = blockDimX.has_value() ? createConst(blockDimX.value()) : one;
413 Value blkSizeY = blockDimY.has_value() ? createConst(blockDimY.value()) : one;
414 Value blkSizeZ = blockDimZ.has_value() ? createConst(blockDimZ.value()) : one;
415 launchOp = LaunchOp::create(rewriter, loc, gridSizeX, gridSizeY, gridSizeZ,
416 blkSizeX, blkSizeY, blkSizeZ);
417 rewriter.setInsertionPointToEnd(&launchOp.getBody().front());
418 TerminatorOp::create(rewriter, loc);
420}
421
422/// Alter kernel configuration of the given kernel.
424 RewriterBase &rewriter, LaunchOp gpuLaunch,
425 TransformOpInterface transformOp, std::optional<int64_t> gridDimX,
426 std::optional<int64_t> gridDimY, std::optional<int64_t> gridDimZ,
427 std::optional<int64_t> blockDimX, std::optional<int64_t> blockDimY,
428 std::optional<int64_t> blockDimZ) {
430 checkGpuLimits(transformOp, gridDimX, gridDimY, gridDimZ, blockDimX,
431 blockDimY, blockDimZ);
432 if (!diag.succeeded())
433 return diag;
434
435 KernelDim3 currentBlockdim = gpuLaunch.getBlockSizeOperandValues();
436 OpBuilder::InsertionGuard guard(rewriter);
437 rewriter.setInsertionPointAfterValue(currentBlockdim.x);
438 auto createConstValue = [&](int dim) {
439 return arith::ConstantIndexOp::create(rewriter, currentBlockdim.x.getLoc(),
440 dim);
441 };
442
443 if (gridDimX.has_value())
444 gpuLaunch.getGridSizeXMutable().assign(createConstValue(gridDimX.value()));
445 if (gridDimY.has_value())
446 gpuLaunch.getGridSizeYMutable().assign(createConstValue(gridDimY.value()));
447 if (gridDimZ.has_value())
448 gpuLaunch.getGridSizeZMutable().assign(createConstValue(gridDimZ.value()));
449 if (blockDimX.has_value())
450 gpuLaunch.getBlockSizeXMutable().assign(
451 createConstValue(blockDimX.value()));
452 if (blockDimY.has_value())
453 gpuLaunch.getBlockSizeYMutable().assign(
454 createConstValue(blockDimY.value()));
455 if (blockDimZ.has_value())
456 gpuLaunch.getBlockSizeZMutable().assign(
457 createConstValue(blockDimZ.value()));
459}
460
461} // namespace gpu
462} // namespace transform
463} // namespace mlir
static Value createConst(Location loc, Type type, int value, PatternRewriter &rewriter)
Create an integer or index constant.
Definition ExpandOps.cpp:27
static FailureOr< SmallVector< Value > > buildPredicates(RewriterBase &rewriter, Location loc, ArrayRef< Value > activeIds, ArrayRef< int64_t > activeMappingSizes, ArrayRef< int64_t > availableMappingSizes, std::string &errorMsg)
Build predicates to filter execution by only the activeIds.
Definition Utils.cpp:52
static Value buildLinearId(RewriterBase &rewriter, Location loc, ArrayRef< OpFoldResult > originalBasisOfr)
Return a flattened thread id for the workgroup with given sizes.
Definition Utils.cpp:82
static GpuIdBuilderFnType common3DIdBuilderFn(int64_t multiplicity=1)
Create a simple 3-D id builder that takes the originalBasisOfr The 3-D id builder returns a 3-D vecto...
Definition Utils.cpp:197
static GpuIdBuilderFnType commonLinearIdBuilderFn(int64_t multiplicity=1, DeviceMaskingAttrInterface mask=nullptr)
Create a linear id builder that takes the originalBasisOfr and decompose it in the basis of forallMap...
Definition Utils.cpp:109
static GpuIdBuilderFnType laneIdBuilderFn(int64_t warpSize)
Create a lane id builder that takes the originalBasis and decompose it in the basis of forallMappingS...
Definition Utils.cpp:233
true
Given two iterators into the same block, return "true" if a is before `b.
static std::string diag(const llvm::Value &value)
constexpr int kMaxGriddimz
constexpr int kMaxTotalBlockdim
constexpr int kMaxGriddimy
constexpr int kMaxBlockdimx
constexpr int kMaxBlockdimz
constexpr int kMaxGriddimx
constexpr int kMaxBlockdimy
constexpr int kMaxTotalGriddim
Base type for affine expression.
Definition AffineExpr.h:68
AffineExpr floorDiv(uint64_t v) const
MLIRContext * getContext() const
Definition Builders.h:56
IndexType getIndexType()
Definition Builders.cpp:59
The result of a transform IR operation application.
static DiagnosedSilenceableFailure success()
Constructs a DiagnosedSilenceableFailure in the success state.
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:63
RAII guard to reset the insertion point of the builder when destroyed.
Definition Builders.h:351
void setInsertionPointToEnd(Block *block)
Sets the insertion point to the end of the specified block.
Definition Builders.h:439
void setInsertionPointAfterValue(Value val)
Sets the insertion point to the node after the specified value.
Definition Builders.h:424
This class represents a single result from folding an operation.
This class coordinates the application of a rewrite on a set of IR, providing a way for clients to tr...
This class represents an instance of an SSA value in the MLIR system, representing a computable value...
Definition Value.h:96
Location getLoc() const
Return the location of this value.
Definition Value.cpp:24
static ConstantIndexOp create(OpBuilder &builder, Location location, int64_t value)
Definition ArithOps.cpp:398
AffineApplyOp makeComposedAffineApply(OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
Returns a composed AffineApplyOp by composing map and operands with other AffineApplyOps supplying th...
OpFoldResult makeComposedFoldedAffineApply(OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
Constructs an AffineApplyOp that applies map to operands after composing the map with the maps of any...
std::function< IdBuilderResult( RewriterBase &, Location, ArrayRef< int64_t >, ArrayRef< int64_t >)> GpuIdBuilderFnType
Common gpu id builder type, allows the configuration of lowering for various mapping schemes.
Definition Utils.h:54
DiagnosedSilenceableFailure alterGpuLaunch(RewriterBase &rewriter, mlir::gpu::LaunchOp gpuLaunch, TransformOpInterface transformOp, std::optional< int64_t > gridDimX=std::nullopt, std::optional< int64_t > gridDimY=std::nullopt, std::optional< int64_t > gridDimZ=std::nullopt, std::optional< int64_t > blockDimX=std::nullopt, std::optional< int64_t > blockDimY=std::nullopt, std::optional< int64_t > blockDimZ=std::nullopt)
Alter kernel configuration of the given kernel.
DiagnosedSilenceableFailure createGpuLaunch(RewriterBase &rewriter, Location loc, TransformOpInterface transformOp, mlir::gpu::LaunchOp &launchOp, std::optional< int64_t > gridDimX=std::nullopt, std::optional< int64_t > gridDimY=std::nullopt, std::optional< int64_t > gridDimZ=std::nullopt, std::optional< int64_t > blockDimX=std::nullopt, std::optional< int64_t > blockDimY=std::nullopt, std::optional< int64_t > blockDimZ=std::nullopt)
Create an empty-body gpu::LaunchOp using the provided kernel settings and put a terminator within.
DiagnosedSilenceableFailure checkGpuLimits(TransformOpInterface transformOp, std::optional< int64_t > gridDimX, std::optional< int64_t > gridDimY, std::optional< int64_t > gridDimZ, std::optional< int64_t > blockDimX, std::optional< int64_t > blockDimY, std::optional< int64_t > blockDimZ)
Determine if the size of the kernel configuration is supported by the GPU architecture being used.
Definition Utils.cpp:362
Include the generated interface declarations.
OpFoldResult getAsIndexOpFoldResult(MLIRContext *ctx, int64_t val)
Convert int64_t to integer attributes of index type and return them as OpFoldResult.
SmallVector< int64_t > computeStrides(ArrayRef< int64_t > sizes)
void bindDims(MLIRContext *ctx, AffineExprTy &...exprs)
Bind a list of AffineExpr references to DimExpr at positions: [0 .
Definition AffineExpr.h:311
SmallVector< int64_t > delinearize(int64_t linearIndex, ArrayRef< int64_t > strides)
Given the strides together with a linear index in the dimension space, return the vector-space offset...
int64_t computeProduct(ArrayRef< int64_t > basis)
Self-explicit.
void bindSymbols(MLIRContext *ctx, AffineExprTy &...exprs)
Bind a list of AffineExpr references to SymbolExpr at positions: [0 .
Definition AffineExpr.h:325
Value getValueOrCreateConstantIndexOp(OpBuilder &b, Location loc, OpFoldResult ofr)
Converts an OpFoldResult to a Value.
Definition Utils.cpp:114
AffineExpr getAffineDimExpr(unsigned position, MLIRContext *context)
These free functions allow clients of the API to not use classes in detail.
Utility class for the GPU dialect to represent triples of Values accessible through ....
Definition GPUDialect.h:39
GpuBlockIdBuilder(MLIRContext *ctx, bool useLinearMapping=false, DeviceMaskingAttrInterface mask=nullptr)
Definition Utils.cpp:298
std::function< DeviceMappingAttrInterface( MLIRContext *, mlir::gpu::MappingId)> MappingIdBuilderFnType
Definition Utils.h:60
SmallVector< DeviceMappingAttrInterface > mappingAttributes
The mapping attributes targeted by this generator.
Definition Utils.h:68
GpuIdBuilderFnType idBuilder
The constructor that builds the concrete IR for mapping ids.
Definition Utils.h:71
GpuLaneIdBuilder(MLIRContext *ctx, int64_t warpSize, bool unused, DeviceMaskingAttrInterface mask=nullptr)
Definition Utils.cpp:351
GpuThreadIdBuilder(MLIRContext *ctx, bool useLinearMapping=false, DeviceMaskingAttrInterface mask=nullptr)
Definition Utils.cpp:340
GpuWarpIdBuilder(MLIRContext *ctx, int64_t warpSize, bool useLinearMapping=false, DeviceMaskingAttrInterface mask=nullptr)
Definition Utils.cpp:325
GpuWarpgroupIdBuilder(MLIRContext *ctx, int64_t warpSize, bool useLinearMapping=false, DeviceMaskingAttrInterface mask=nullptr)
Definition Utils.cpp:309
Helper type for functions that generate ids for the mapping of a scf.forall.
Definition Utils.h:30