GPUTransformOps.cpp File Reference

#include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/Dialect/Transform/IR/TransformDialect.h"
#include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
#include "mlir/Dialect/Utils/IndexingUtils.h"
#include "mlir/IR/AffineExpr.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/IRMapping.h"
#include "mlir/IR/MLIRContext.h"
#include "mlir/IR/OpDefinition.h"
#include "mlir/IR/Visitors.h"
#include "mlir/Support/LLVM.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.cpp.inc"

Go to the source code of this file.

Classes
struct	ForallRewriteResult
	Struct to return the result of the rewrite of a forall operation. More...

Macros
#define	DEBUG_TYPE   "gpu-transforms"
#define	DBGS()   (llvm::dbgs() << '[' << DEBUG_TYPE << "] ")
#define	LDBG(X)   LLVM_DEBUG(DBGS() << X << "\n")
#define	GET_OP_LIST
#define	GET_OP_CLASSES

Functions
static DiagnosedSilenceableFailure	definiteFailureHelper (std::optional< TransformOpInterface > transformOp, Operation *target, const Twine &message)
static DiagnosedSilenceableFailure	checkMappingAttributeTypes (std::optional< TransformOpInterface > transformOp, scf::ForallOp forallOp)
	Check if given mapping attributes are one of the desired attributes. More...
static DiagnosedSilenceableFailure	verifyGpuMapping (std::optional< TransformOpInterface > transformOp, scf::ForallOp forallOp)
static DiagnosedSilenceableFailure	checkGpuLimits (TransformOpInterface transformOp, std::optional< int64_t > gridDimX, std::optional< int64_t > gridDimY, std::optional< int64_t > gridDimZ, std::optional< int64_t > blockDimX, std::optional< int64_t > blockDimY, std::optional< int64_t > blockDimZ)
	Determines if the size of the kernel configuration is supported by the GPU architecture being used. More...
static DiagnosedSilenceableFailure	createGpuLaunch (RewriterBase &rewriter, Location loc, TransformOpInterface transformOp, LaunchOp &launchOp, std::optional< int64_t > gridDimX=std::nullopt, std::optional< int64_t > gridDimY=std::nullopt, std::optional< int64_t > gridDimZ=std::nullopt, std::optional< int64_t > blockDimX=std::nullopt, std::optional< int64_t > blockDimY=std::nullopt, std::optional< int64_t > blockDimZ=std::nullopt)
	Creates an empty-body gpu::LaunchOp using the provided kernel settings and put a terminator within. More...
static DiagnosedSilenceableFailure	alterGpuLaunch (IRRewriter &rewriter, LaunchOp gpuLaunch, TransformOpInterface transformOp, std::optional< int64_t > gridDimX=std::nullopt, std::optional< int64_t > gridDimY=std::nullopt, std::optional< int64_t > gridDimZ=std::nullopt, std::optional< int64_t > blockDimX=std::nullopt, std::optional< int64_t > blockDimY=std::nullopt, std::optional< int64_t > blockDimZ=std::nullopt)
	Alter kernel configuration of the given kernel. More...
template<typename OpTy , typename OperationOrBlock >
static void	replaceUnitMappingIdsHelper (RewriterBase &rewriter, Location loc, OperationOrBlock *parent, Value replacement, ArrayRef< int64_t > availableMappingSizes)
	Helper to replace ids of dimensions known to be 1 by 0 to simplify the IR. More...
static DiagnosedSilenceableFailure	rewriteOneForallCommonImpl (RewriterBase &rewriter, std::optional< TransformOpInterface > transformOp, scf::ForallOp forallOp, ForallRewriteResult &result, ArrayRef< int64_t > availableMappingSizes, const GpuIdBuilder &gpuIdBuilder)

Macro Definition Documentation

DBGS

#define DBGS

(

)

(llvm::dbgs() << '[' << DEBUG_TYPE << "] ")

Definition at line 40 of file GPUTransformOps.cpp.

DEBUG_TYPE

#define DEBUG_TYPE   "gpu-transforms"

Definition at line 38 of file GPUTransformOps.cpp.

GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 917 of file GPUTransformOps.cpp.

GET_OP_LIST

#define GET_OP_LIST

LDBG

#define LDBG

(

X

)

LLVM_DEBUG(DBGS() << X << "\n")

Definition at line 41 of file GPUTransformOps.cpp.

Function Documentation

alterGpuLaunch()

static DiagnosedSilenceableFailure alterGpuLaunch ( IRRewriter &  rewriter, LaunchOp  gpuLaunch, TransformOpInterface  transformOp, std::optional< int64_t >  gridDimX = std::nullopt, std::optional< int64_t >  gridDimY = std::nullopt, std::optional< int64_t >  gridDimZ = std::nullopt, std::optional< int64_t >  blockDimX = std::nullopt, std::optional< int64_t >  blockDimY = std::nullopt, std::optional< int64_t >  blockDimZ = std::nullopt  )

static

Alter kernel configuration of the given kernel.

Definition at line 401 of file GPUTransformOps.cpp.

References checkGpuLimits(), mlir::OpBuilder::create(), diag(), mlir::Value::getLoc(), mlir::OpBuilder::setInsertionPointAfterValue(), mlir::DiagnosedSilenceableFailure::success(), and mlir::gpu::KernelDim3::x.

checkGpuLimits()

static DiagnosedSilenceableFailure checkGpuLimits ( TransformOpInterface  transformOp, std::optional< int64_t >  gridDimX, std::optional< int64_t >  gridDimY, std::optional< int64_t >  gridDimZ, std::optional< int64_t >  blockDimX, std::optional< int64_t >  blockDimY, std::optional< int64_t >  blockDimZ  )

static

Determines if the size of the kernel configuration is supported by the GPU architecture being used.

It presently makes use of CUDA limitations, however that aspect may be enhanced for other GPUs.

Definition at line 329 of file GPUTransformOps.cpp.

References mlir::DiagnosedSilenceableFailure::success().

Referenced by alterGpuLaunch(), and createGpuLaunch().

checkMappingAttributeTypes()

static DiagnosedSilenceableFailure checkMappingAttributeTypes ( std::optional< TransformOpInterface >  transformOp, scf::ForallOp  forallOp  )

static

Check if given mapping attributes are one of the desired attributes.

Definition at line 251 of file GPUTransformOps.cpp.

References definiteFailureHelper(), and mlir::DiagnosedSilenceableFailure::success().

Referenced by verifyGpuMapping().

createGpuLaunch()

static DiagnosedSilenceableFailure createGpuLaunch ( RewriterBase &  rewriter, Location  loc, TransformOpInterface  transformOp, LaunchOp &  launchOp, std::optional< int64_t >  gridDimX = std::nullopt, std::optional< int64_t >  gridDimY = std::nullopt, std::optional< int64_t >  gridDimZ = std::nullopt, std::optional< int64_t >  blockDimX = std::nullopt, std::optional< int64_t >  blockDimY = std::nullopt, std::optional< int64_t >  blockDimZ = std::nullopt  )

static

Creates an empty-body gpu::LaunchOp using the provided kernel settings and put a terminator within.

Definition at line 367 of file GPUTransformOps.cpp.

References checkGpuLimits(), mlir::OpBuilder::create(), createConst(), diag(), mlir::OpBuilder::setInsertionPointToEnd(), and mlir::DiagnosedSilenceableFailure::success().

definiteFailureHelper()

static DiagnosedSilenceableFailure definiteFailureHelper ( std::optional< TransformOpInterface >  transformOp, Operation *  target, const Twine &  message  )

static

Definition at line 242 of file GPUTransformOps.cpp.

References mlir::emitDefiniteFailure().

Referenced by checkMappingAttributeTypes(), mlir::transform::gpu::mapNestedForallToThreadsImpl(), rewriteOneForallCommonImpl(), and verifyGpuMapping().

replaceUnitMappingIdsHelper()

template<typename OpTy , typename OperationOrBlock >

static void replaceUnitMappingIdsHelper ( RewriterBase &  rewriter, Location  loc, OperationOrBlock *  parent, Value  replacement, ArrayRef< int64_t >  availableMappingSizes  )

static

Helper to replace ids of dimensions known to be 1 by 0 to simplify the IR.

Definition at line 450 of file GPUTransformOps.cpp.

References mlir::RewriterBase::replaceAllUsesWith().

rewriteOneForallCommonImpl()

static DiagnosedSilenceableFailure rewriteOneForallCommonImpl ( RewriterBase &  rewriter, std::optional< TransformOpInterface >  transformOp, scf::ForallOp  forallOp, ForallRewriteResult &  result, ArrayRef< int64_t >  availableMappingSizes, const GpuIdBuilder &  gpuIdBuilder  )

static

Definition at line 459 of file GPUTransformOps.cpp.

References mlir::Block::begin(), mlir::OpBuilder::create(), DBGS, definiteFailureHelper(), diag(), mlir::RewriterBase::eraseOp(), mlir::Block::front(), mlir::OpBuilder::getInsertionPoint(), mlir::Block::getOperations(), mlir::getValuesSortedByKey(), mlir::transform::gpu::GpuIdBuilder::idBuilder, LDBG, mlir::IRMapping::lookup(), mlir::IRMapping::map(), mlir::transform::gpu::GpuIdBuilder::mappingAttributes, mlir::transform::gpu::IdBuilderResult::mappingIdOps, mlir::transform::gpu::IdBuilderResult::predicateIdOps, mlir::transform::gpu::IdBuilderResult::predicateMappingSizes, mlir::RewriterBase::replaceAllUsesWith(), mlir::OpBuilder::setInsertionPoint(), mlir::DiagnosedSilenceableFailure::success(), and verifyGpuMapping().

Referenced by mlir::transform::gpu::mapForallToBlocksImpl(), and mlir::transform::gpu::mapOneForallToThreadsImpl().

verifyGpuMapping()

static DiagnosedSilenceableFailure verifyGpuMapping ( std::optional< TransformOpInterface >  transformOp, scf::ForallOp  forallOp  )

static

Definition at line 299 of file GPUTransformOps.cpp.

References checkMappingAttributeTypes(), definiteFailureHelper(), mlir::DiagnosedSilenceableFailure::succeeded(), and mlir::DiagnosedSilenceableFailure::success().

Referenced by rewriteOneForallCommonImpl().