#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/IR/Matchers.h"
#include "mlir/Interfaces/FunctionInterfaces.h"
#include "mlir/Interfaces/InferIntRangeInterface.h"
#include "llvm/Support/ErrorHandling.h"
#include <optional>

Functions
static ConstantIntRanges	getIndexRange (uint64_t umin, uint64_t umax)
static Value	valueByDim (KernelDim3 dims, Dimension dim)
	If the operation op is in a context that is annotated with maximum launch dimensions (a launch op with constant block or grid sizes or a launch_func op with the appropriate dimensions), return the bound on the maximum size of the dimension that the op is querying.
static uint64_t	zext (uint32_t arg)
static std::optional< uint64_t >	getKnownLaunchAttr (GPUFuncOp func, LaunchDims dims, Dimension dim)
static std::optional< uint64_t >	getKnownLaunchAttr (FunctionOpInterface func, StringRef attrName, Dimension dim)
template<typename Op>
static std::optional< uint64_t >	getKnownLaunchDim (Op op, LaunchDims type)

Variables
static constexpr uint64_t	kMaxDim = std::numeric_limits<uint32_t>::max()
static constexpr uint64_t	kMaxClusterDim = 8
static constexpr uint64_t	kMaxSubgroupSize = 128

Function Documentation

◆ getIndexRange()

ConstantIntRanges getIndexRange	(	uint64_t	umin,
		uint64_t	umax )

static

Definition at line 26 of file InferIntRangeInterfaceImpls.cpp.

References mlir::ConstantIntRanges::fromUnsigned().

◆ getKnownLaunchAttr() [1/2]

std::optional< uint64_t > getKnownLaunchAttr	(	FunctionOpInterface	func,
		StringRef	attrName,
		Dimension	dim )

static

Definition at line 74 of file InferIntRangeInterfaceImpls.cpp.

References zext().

◆ getKnownLaunchAttr() [2/2]

std::optional< uint64_t > getKnownLaunchAttr	(	GPUFuncOp	func,
		LaunchDims	dims,
		Dimension	dim )

static

Definition at line 57 of file InferIntRangeInterfaceImpls.cpp.

References zext().

Referenced by getKnownLaunchDim().

◆ getKnownLaunchDim()

template<typename Op>

std::optional< uint64_t > getKnownLaunchDim	(	Op	op,
		LaunchDims	type )

static

Definition at line 86 of file InferIntRangeInterfaceImpls.cpp.

References getKnownLaunchAttr(), mlir::m_ConstantInt(), mlir::matchPattern(), and valueByDim().

◆ valueByDim()

Value valueByDim	(	KernelDim3	dims,
		Dimension	dim )

static

If the operation op is in a context that is annotated with maximum launch dimensions (a launch op with constant block or grid sizes or a launch_func op with the appropriate dimensions), return the bound on the maximum size of the dimension that the op is querying.

IDs will be one less than this bound.

Definition at line 42 of file InferIntRangeInterfaceImpls.cpp.

References mlir::gpu::KernelDim3::x, mlir::gpu::KernelDim3::y, and mlir::gpu::KernelDim3::z.

Referenced by getKnownLaunchDim().