19 #include "llvm/IR/IRBuilder.h" 20 #include "llvm/IR/IntrinsicsNVPTX.h" 31 resultType = cast<llvm::StructType>(resultType)->
getElementType(0);
33 case NVVM::ShflKind::bfly:
34 return resultType->isFloatTy()
35 ? llvm::Intrinsic::nvvm_shfl_sync_bfly_f32p
36 : llvm::Intrinsic::nvvm_shfl_sync_bfly_i32p;
37 case NVVM::ShflKind::up:
38 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_up_f32p
39 : llvm::Intrinsic::nvvm_shfl_sync_up_i32p;
40 case NVVM::ShflKind::down:
41 return resultType->isFloatTy()
42 ? llvm::Intrinsic::nvvm_shfl_sync_down_f32p
43 : llvm::Intrinsic::nvvm_shfl_sync_down_i32p;
44 case NVVM::ShflKind::idx:
45 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_idx_f32p
46 : llvm::Intrinsic::nvvm_shfl_sync_idx_i32p;
50 case NVVM::ShflKind::bfly:
51 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_bfly_f32
52 : llvm::Intrinsic::nvvm_shfl_sync_bfly_i32;
53 case NVVM::ShflKind::up:
54 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_up_f32
55 : llvm::Intrinsic::nvvm_shfl_sync_up_i32;
56 case NVVM::ShflKind::down:
57 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_down_f32
58 : llvm::Intrinsic::nvvm_shfl_sync_down_i32;
59 case NVVM::ShflKind::idx:
60 return resultType->isFloatTy() ? llvm::Intrinsic::nvvm_shfl_sync_idx_f32
61 : llvm::Intrinsic::nvvm_shfl_sync_idx_i32;
64 llvm_unreachable(
"unknown shuffle kind");
70 if (layout == NVVM::MMALayout::row) {
73 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16;
75 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16;
77 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16;
79 llvm_unreachable(
"unsupported number of matrix");
85 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16;
87 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16;
89 return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16;
91 llvm_unreachable(
"unsupported number of matrix");
99 class NVVMDialectLLVMIRTranslationInterface
107 convertOperation(
Operation *op, llvm::IRBuilderBase &builder,
110 #include "mlir/Dialect/LLVMIR/NVVMConversions.inc" 119 if (attribute.
getName() == NVVM::NVVMDialect::getKernelFuncAttrName()) {
120 auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
124 llvm::LLVMContext &llvmContext = moduleTranslation.
getLLVMContext();
125 llvm::Function *llvmFunc =
127 llvm::Metadata *llvmMetadata[] = {
128 llvm::ValueAsMetadata::get(llvmFunc),
129 llvm::MDString::get(llvmContext,
"kernel"),
130 llvm::ValueAsMetadata::get(
131 llvm::ConstantInt::get(llvm::Type::getInt32Ty(llvmContext), 1))};
132 llvm::MDNode *llvmMetadataNode =
133 llvm::MDNode::get(llvmContext, llvmMetadata);
135 ->addOperand(llvmMetadataNode);
143 registry.
insert<NVVM::NVVMDialect>();
145 dialect->addInterfaces<NVVMDialectLLVMIRTranslationInterface>();
Include the generated interface declarations.
Operation is a basic unit of execution within MLIR.
llvm::NamedMDNode * getOrInsertNamedModuleMetadata(StringRef name)
Gets the named metadata in the LLVM IR module being constructed, creating it if it does not exist...
static Type getElementType(Type type, ArrayRef< int32_t > indices, function_ref< InFlightDiagnostic(StringRef)> emitErrorFn)
Walks the given type hierarchy with the given indices, potentially down to component granularity...
void appendDialectRegistry(const DialectRegistry ®istry)
Append the contents of the given dialect registry to the registry associated with this context...
llvm::LLVMContext & getLLVMContext() const
Returns the LLVM context in which the IR is being constructed.
NamedAttribute represents a combination of a name and an Attribute value.
llvm::Value * createIntrinsicCall(llvm::IRBuilderBase &builder, llvm::Intrinsic::ID intrinsic, ArrayRef< llvm::Value *> args={}, ArrayRef< llvm::Type *> tys={})
Creates a call to an LLVM IR intrinsic function with the given arguments.
Implementation class for module translation.
LogicalResult success(bool isSuccess=true)
Utility function to generate a LogicalResult.
void addExtension(std::unique_ptr< DialectExtensionBase > extension)
Add the given extension to the registry.
This class represents an efficient way to signal success or failure.
LogicalResult failure(bool isFailure=true)
Utility function to generate a LogicalResult.
Base class for dialect interfaces providing translation to LLVM IR.
StringAttr getName() const
Return the name of the attribute.
static llvm::Intrinsic::ID getLdMatrixIntrinsicId(NVVM::MMALayout layout, int32_t num)
Return the intrinsic ID associated with ldmatrix for the given paramters.
llvm::Function * lookupFunction(StringRef name) const
Finds an LLVM IR function by its name.
The DialectRegistry maps a dialect namespace to a constructor for the matching dialect.
MLIRContext is the top-level object for a collection of MLIR operations.
LLVMTranslationDialectInterface(Dialect *dialect)
void registerNVVMDialectTranslation(DialectRegistry ®istry)
Register the NVVM dialect and the translation from it to the LLVM IR in the given registry;...
static llvm::Intrinsic::ID getShflIntrinsicId(llvm::Type *resultType, NVVM::ShflKind kind, bool withPredicate)