MLIR  19.0.0git
Public Member Functions | Public Attributes | List of all members
mlir::transform::gpu::GpuWarpIdBuilder Struct Reference

Builder for warp ids used to map scf.forall to reindexed warps. More...

#include "mlir/Dialect/GPU/TransformOps/Utils.h"

+ Inheritance diagram for mlir::transform::gpu::GpuWarpIdBuilder:

Public Member Functions

 GpuWarpIdBuilder (MLIRContext *ctx, int64_t warpSize, bool useLinearMapping=false)
- Public Member Functions inherited from mlir::transform::gpu::GpuIdBuilder
 GpuIdBuilder ()=default
 GpuIdBuilder (MLIRContext *ctx, bool useLinearMapping, const MappingIdBuilderFnType &builder)

Public Attributes

int64_t warpSize = 32
- Public Attributes inherited from mlir::transform::gpu::GpuIdBuilder
SmallVector< DeviceMappingAttrInterface > mappingAttributes
 The mapping attributes targeted by this generator. More...
GpuIdBuilderFnType idBuilder
 The constructor that builds the concrete IR for mapping ids. More...

Additional Inherited Members

- Public Types inherited from mlir::transform::gpu::GpuIdBuilder
using MappingIdBuilderFnType = std::function< DeviceMappingAttrInterface(MLIRContext *, mlir::gpu::MappingId)>

Detailed Description

Builder for warp ids used to map scf.forall to reindexed warps.

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.

Definition at line 105 of file Utils.h.

Constructor & Destructor Documentation

◆ GpuWarpIdBuilder()

mlir::transform::gpu::GpuWarpIdBuilder::GpuWarpIdBuilder ( MLIRContext ctx,
int64_t  warpSize,
bool  useLinearMapping = false 

Definition at line 210 of file Utils.cpp.

References mlir::get().

Member Data Documentation

◆ warpSize

int64_t mlir::transform::gpu::GpuWarpIdBuilder::warpSize = 32

Definition at line 108 of file Utils.h.

The documentation for this struct was generated from the following files: