MLIR  19.0.0git
Functions | Variables
InferIntRangeInterfaceImpls.cpp File Reference
#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>

Go to the source code of this file.

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

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

◆ getKnownLaunchDim()

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

Definition at line 88 of file InferIntRangeInterfaceImpls.cpp.

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

◆ zext()

static uint64_t zext ( uint32_t  arg)
static

Definition at line 56 of file InferIntRangeInterfaceImpls.cpp.

Referenced by getKnownLaunchAttr().

Variable Documentation

◆ kMaxClusterDim

constexpr uint64_t kMaxClusterDim = 8
staticconstexpr

Definition at line 24 of file InferIntRangeInterfaceImpls.cpp.

◆ kMaxDim

constexpr uint64_t kMaxDim = std::numeric_limits<uint32_t>::max()
staticconstexpr

Definition at line 22 of file InferIntRangeInterfaceImpls.cpp.

◆ kMaxSubgroupSize

constexpr uint64_t kMaxSubgroupSize = 128
staticconstexpr

Definition at line 26 of file InferIntRangeInterfaceImpls.cpp.