MLIR  16.0.0git
Functions
NVGPUToNVVM.cpp File Reference
#include "mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h"
#include "../PassDetail.h"
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
+ Include dependency graph for NVGPUToNVVM.cpp:

Go to the source code of this file.

Functions

static Type inferIntrinsicResultType (Type vectorResultType)
 Returns the type for the intrinsic given the vectorResultType of the gpu.mma.sync operation. More...
 
static Value convertIntrinsicResult (Location loc, Type intrinsicResultType, Type resultType, Value intrinsicResult, RewriterBase &rewriter)
 Convert the SSA result of the NVVM intrinsic nvvm.mma.sync (which is always an LLVM struct) into a fragment that is compatible with the vector type of this operation. More...
 
static SmallVector< ValueunpackOperandVector (RewriterBase &rewriter, Location loc, Value operand, NVVM::MMATypes operandPtxType)
 The gpu.mma.sync converter below expects matrix fragment operands to be given as 2D vectors where the rows are 32b or 64b wide. More...
 

Function Documentation

◆ convertIntrinsicResult()

static Value convertIntrinsicResult ( Location  loc,
Type  intrinsicResultType,
Type  resultType,
Value  intrinsicResult,
RewriterBase rewriter 
)
static

Convert the SSA result of the NVVM intrinsic nvvm.mma.sync (which is always an LLVM struct) into a fragment that is compatible with the vector type of this operation.

This involves extracting elements from the struct and inserting them into an LLVM array. These extra data-movement operations should be canonicalized away by the LLVM backend.

Definition at line 60 of file NVGPUToNVVM.cpp.

References mlir::OpBuilder::create(), mlir::OpBuilder::createOrFold(), mlir::Type::dyn_cast(), mlir::detail::enumerate(), mlir::Builder::getContext(), mlir::Builder::getF16Type(), mlir::Builder::getF32Type(), mlir::Builder::getF64Type(), mlir::LLVM::getFixedVectorType(), mlir::Builder::getI32IntegerAttr(), and mlir::Builder::getI32Type().

Referenced by unpackOperandVector().

◆ inferIntrinsicResultType()

static Type inferIntrinsicResultType ( Type  vectorResultType)
static

Returns the type for the intrinsic given the vectorResultType of the gpu.mma.sync operation.

Definition at line 21 of file NVGPUToNVVM.cpp.

References mlir::Type::cast(), mlir::Type::getContext(), mlir::LLVM::getFixedVectorType(), and mlir::LLVM::LLVMStructType::getLiteral().

Referenced by unpackOperandVector().

◆ unpackOperandVector()

static SmallVector<Value> unpackOperandVector ( RewriterBase rewriter,
Location  loc,
Value  operand,
NVVM::MMATypes  operandPtxType 
)
static