MLIR
21.0.0git
|
Builder for warpgroup ids used to map scf.forall to reindexed warpgroups. More...
#include "mlir/Dialect/GPU/TransformOps/Utils.h"
Public Member Functions | |
GpuWarpgroupIdBuilder (MLIRContext *ctx, int64_t warpSize, bool useLinearMapping=false) | |
![]() | |
GpuIdBuilder ()=default | |
GpuIdBuilder (MLIRContext *ctx, bool useLinearMapping, const MappingIdBuilderFnType &builder) | |
Public Attributes | |
int64_t | warpSize = 32 |
![]() | |
SmallVector< DeviceMappingAttrInterface > | mappingAttributes |
The mapping attributes targeted by this generator. More... | |
GpuIdBuilderFnType | idBuilder |
The constructor that builds the concrete IR for mapping ids. More... | |
Static Public Attributes | |
static constexpr int64_t | kNumWarpsPerGroup = 4 |
In the future this may be configured by the transformation. More... | |
Additional Inherited Members | |
![]() | |
using | MappingIdBuilderFnType = std::function< DeviceMappingAttrInterface(MLIRContext *, mlir::gpu::MappingId)> |
Builder for warpgroup ids used to map scf.forall to reindexed warpgroups.
If useLinearMapping
is false, the idBuilder
method returns 3D values used for indexing rewrites as well as 3D sizes for predicate generation. If useLinearMapping
is true, the idBuilder
method returns nD values used for indexing rewrites as well as 1D sizes for predicate generation.
mlir::transform::gpu::GpuWarpgroupIdBuilder::GpuWarpgroupIdBuilder | ( | MLIRContext * | ctx, |
int64_t | warpSize, | ||
bool | useLinearMapping = false |
||
) |
Definition at line 196 of file Utils.cpp.
References mlir::get().
|
staticconstexpr |
int64_t mlir::transform::gpu::GpuWarpgroupIdBuilder::warpSize = 32 |