MLIR
17.0.0git
|
#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) |
#define DBGS | ( | ) | (llvm::dbgs() << '[' << DEBUG_TYPE << "] ") |
Definition at line 40 of file GPUTransformOps.cpp.
#define DEBUG_TYPE "gpu-transforms" |
Definition at line 38 of file GPUTransformOps.cpp.
#define GET_OP_CLASSES |
Definition at line 917 of file GPUTransformOps.cpp.
#define GET_OP_LIST |
#define LDBG | ( | X | ) | LLVM_DEBUG(DBGS() << X << "\n") |
Definition at line 41 of file GPUTransformOps.cpp.
|
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.
|
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().
|
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().
|
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().
|
static |
Definition at line 242 of file GPUTransformOps.cpp.
References mlir::emitDefiniteFailure().
Referenced by checkMappingAttributeTypes(), mlir::transform::gpu::mapNestedForallToThreadsImpl(), rewriteOneForallCommonImpl(), and verifyGpuMapping().
|
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().
|
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().
|
static |
Definition at line 299 of file GPUTransformOps.cpp.
References checkMappingAttributeTypes(), definiteFailureHelper(), mlir::DiagnosedSilenceableFailure::succeeded(), and mlir::DiagnosedSilenceableFailure::success().
Referenced by rewriteOneForallCommonImpl().