#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/IR/Matchers.h"
#include "mlir/Interfaces/FunctionInterfaces.h"
#include "mlir/Interfaces/InferIntRangeInterface.h"
#include "llvm/ADT/STLForwardCompat.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/MathExtras.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. More...

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()

static ConstantIntRanges getIndexRange	(	uint64_t	umin,
		uint64_t	umax
	)

static

Definition at line 28 of file InferIntRangeInterfaceImpls.cpp.

References mlir::ConstantIntRanges::fromUnsigned().

◆ getKnownLaunchAttr() [1/2]

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

static

Definition at line 76 of file InferIntRangeInterfaceImpls.cpp.

References zext().

◆ getKnownLaunchAttr() [2/2]

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

static

Definition at line 59 of file InferIntRangeInterfaceImpls.cpp.

References zext().

Referenced by getKnownLaunchDim().

◆ getKnownLaunchDim()

template<typename Op >

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

static

Definition at line 88 of file InferIntRangeInterfaceImpls.cpp.

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

◆ valueByDim()

static 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 44 of file InferIntRangeInterfaceImpls.cpp.

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

Referenced by getKnownLaunchDim().