mlir::gpu::index_lowering Namespace Reference

Classes
struct	OpLowering

Typedefs
using	IndexKind = gpu::DimensionKind

Enumerations
enum class	IntrType : uint32_t { None = 0 , Id = 1 , Dim = 2 }

Functions
LLVM::ConstantRangeAttr	getIndexOpRange (Operation *op, gpu::Dimension dim, std::optional< uint32_t > opUpperBound, IndexKind indexKind, IntrType intrType, unsigned bitWidth)
	Returns a ConstantRangeAttr for a GPU index op, or nullptr if no bounds are found.

Typedef Documentation

IndexKind

using mlir::gpu::index_lowering::IndexKind = gpu::DimensionKind

Definition at line 20 of file IndexIntrinsicsOpLowering.h.

Enumeration Type Documentation

IntrType

enum class mlir::gpu::index_lowering::IntrType : uint32_t

strong

Enumerator
None
Id
Dim

Definition at line 22 of file IndexIntrinsicsOpLowering.h.

Function Documentation

getIndexOpRange()

LLVM::ConstantRangeAttr mlir::gpu::index_lowering::getIndexOpRange	(	Operation *	op,
		gpu::Dimension	dim,
		std::optional< uint32_t >	opUpperBound,
		IndexKind	indexKind,
		IntrType	intrType,
		unsigned	bitWidth )

Returns a ConstantRangeAttr for a GPU index op, or nullptr if no bounds are found.

bitWidth controls the width of the returned range. Checks the provided upper_bound from the op (highest priority), inherent attrs on enclosing gpu.funcs, and discardable attributes on other enclosing function ops (lowest priority). However, in the case where a dimension is known to have a constant value, returns a range indicating that value.

Definition at line 17 of file IndexIntrinsicsOpLowering.cpp.

References Dim, mlir::Operation::getContext(), mlir::gpu::getKnownDimensionSizeAround(), Id, max(), min(), and None.

Referenced by mlir::gpu::index_lowering::OpLowering< Op, XOp, YOp, ZOp >::matchAndRewrite().