Builder for warpgroup ids used to map scf.forall to reindexed warpgroups. More...

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

Inheritance diagram for mlir::transform::gpu::GpuWarpgroupIdBuilder:

Public Member Functions
	GpuWarpgroupIdBuilder (MLIRContext *ctx, int64_t warpSize, bool useLinearMapping=false, DeviceMaskingAttrInterface mask=nullptr)

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...

Static Public Attributes
static constexpr int64_t	kNumWarpsPerGroup = 4
	In the future this may be configured by the transformation. 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 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. If mask is provided, it will be used to filter the active warpgroups.

Definition at line 92 of file Utils.h.

Constructor & Destructor Documentation

◆ GpuWarpgroupIdBuilder()

mlir::transform::gpu::GpuWarpgroupIdBuilder::GpuWarpgroupIdBuilder	(	MLIRContext *	ctx,
		int64_t	warpSize,
		bool	useLinearMapping = `false`,
		DeviceMaskingAttrInterface	mask = `nullptr`
	)

Definition at line 318 of file Utils.cpp.

References mlir::get().

Member Data Documentation

◆ kNumWarpsPerGroup

constexpr int64_t mlir::transform::gpu::GpuWarpgroupIdBuilder::kNumWarpsPerGroup = 4

staticconstexpr

In the future this may be configured by the transformation.

Definition at line 98 of file Utils.h.

◆ warpSize

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

Definition at line 96 of file Utils.h.

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

include/mlir/Dialect/GPU/TransformOps/Utils.h
lib/Dialect/GPU/TransformOps/Utils.cpp

Public Member Functions

Public Attributes