LLVMIRToNVVMTranslation.cpp File Reference

#include "mlir/Target/LLVMIR/Dialect/NVVM/LLVMIRToNVVMTranslation.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Target/LLVMIR/ModuleImport.h"
#include "llvm/IR/ConstantRange.h"
#include "mlir/Dialect/LLVMIR/NVVMConvertibleLLVMIRIntrinsics.inc"
#include "mlir/Dialect/LLVMIR/NVVMFromLLVMIRConversions.inc"

Go to the source code of this file.

Functions
static bool	isBarrierSyncIntrinsic (llvm::Intrinsic::ID id)
	Returns true if id is one of the bar.sync intrinsic IDs imported into nvvm.barrier.
static LogicalResult	convertBarrierSyncIntrinsic (OpBuilder &odsBuilder, llvm::CallInst *inst, LLVM::ModuleImport &moduleImport, ArrayRef< llvm::Value * > llvmOperands, ArrayRef< llvm::OperandBundleUse > llvmOpBundles)
	Imports one of the four bar.sync LLVM intrinsic variants into a single nvvm.barrier op, deriving the aligned attribute and the optional numberOfThreads operand from the specific intrinsic ID.
static bool	isConvertibleIntrinsic (llvm::Intrinsic::ID id)
	Returns true if the LLVM intrinsic id has a corresponding MLIR NVVM op.
static ArrayRef< unsigned >	getSupportedIntrinsicsImpl ()
	Returns the LLVM intrinsic IDs that have a corresponding MLIR NVVM op.
static LogicalResult	convertIntrinsicImpl (OpBuilder &odsBuilder, llvm::CallInst *inst, LLVM::ModuleImport &moduleImport)
	Converts the LLVM intrinsic to an MLIR NVVM dialect operation if a conversion exits.

Variables
static constexpr llvm::Intrinsic::ID	kBarrierSyncIntrinsics []
	LLVM intrinsic IDs of the four bar.sync variants imported into nvvm.barrier.

Function Documentation

convertBarrierSyncIntrinsic()

LogicalResult convertBarrierSyncIntrinsic ( OpBuilder & odsBuilder, llvm::CallInst * inst, LLVM::ModuleImport & moduleImport, ArrayRef< llvm::Value * > llvmOperands, ArrayRef< llvm::OperandBundleUse > llvmOpBundles )

static

Imports one of the four bar.sync LLVM intrinsic variants into a single nvvm.barrier op, deriving the aligned attribute and the optional numberOfThreads operand from the specific intrinsic ID.

Definition at line 40 of file LLVMIRToNVVMTranslation.cpp.

References mlir::LLVM::ModuleImport::convertIntrinsicArguments(), mlir::Builder::getBoolAttr(), mlir::LLVM::ModuleImport::mapNoResultOp(), success(), and mlir::LLVM::ModuleImport::translateLoc().

Referenced by convertIntrinsicImpl().

convertIntrinsicImpl()

LogicalResult convertIntrinsicImpl ( OpBuilder & odsBuilder, llvm::CallInst * inst, LLVM::ModuleImport & moduleImport )

static

Converts the LLVM intrinsic to an MLIR NVVM dialect operation if a conversion exits.

Returns failure otherwise.

Definition at line 88 of file LLVMIRToNVVMTranslation.cpp.

References convertBarrierSyncIntrinsic(), isBarrierSyncIntrinsic(), and isConvertibleIntrinsic().

getSupportedIntrinsicsImpl()

ArrayRef< unsigned > getSupportedIntrinsicsImpl ( )

static

Returns the LLVM intrinsic IDs that have a corresponding MLIR NVVM op.

Definition at line 74 of file LLVMIRToNVVMTranslation.cpp.

References kBarrierSyncIntrinsics.

isBarrierSyncIntrinsic()

bool isBarrierSyncIntrinsic ( llvm::Intrinsic::ID id )

static

Returns true if id is one of the bar.sync intrinsic IDs imported into nvvm.barrier.

Definition at line 32 of file LLVMIRToNVVMTranslation.cpp.

References kBarrierSyncIntrinsics.

Referenced by convertIntrinsicImpl(), and isConvertibleIntrinsic().

isConvertibleIntrinsic()

bool isConvertibleIntrinsic ( llvm::Intrinsic::ID id )

static

Returns true if the LLVM intrinsic id has a corresponding MLIR NVVM op.

Definition at line 66 of file LLVMIRToNVVMTranslation.cpp.

References isBarrierSyncIntrinsic().

Referenced by convertIntrinsicImpl().

Variable Documentation

kBarrierSyncIntrinsics

llvm::Intrinsic::ID kBarrierSyncIntrinsics[]

staticconstexpr

Initial value:

= {
    llvm::Intrinsic::nvvm_barrier_cta_sync_all,
    llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all,
    llvm::Intrinsic::nvvm_barrier_cta_sync_count,
    llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count,
}

LLVM intrinsic IDs of the four bar.sync variants imported into nvvm.barrier.

Definition at line 23 of file LLVMIRToNVVMTranslation.cpp.

Referenced by getSupportedIntrinsicsImpl(), and isBarrierSyncIntrinsic().