15#include "llvm/ADT/TypeSwitch.h"
21 llvm::IRBuilderBase &builder,
24 launchOp, launchOp.getKernelModuleName());
26 launchOp.emitError(
"Couldn't find the binary holding the kernel: ")
27 << launchOp.getKernelModuleName();
30 auto offloadingHandler =
31 dyn_cast<gpu::OffloadingLLVMTranslationAttrInterface>(
32 kernelBinary.getOffloadingHandlerAttr());
33 assert(offloadingHandler &&
"Invalid offloading handler.");
34 return offloadingHandler.launchKernel(launchOp, kernelBinary, builder,
38class GPUDialectLLVMIRTranslationInterface
44 convertOperation(Operation *operation, llvm::IRBuilderBase &builder,
45 LLVM::ModuleTranslation &moduleTranslation)
const override {
46 return llvm::TypeSwitch<Operation *, LogicalResult>(operation)
47 .Case([&](gpu::GPUModuleOp) {
return success(); })
48 .Case([&](gpu::BinaryOp op) {
49 auto offloadingHandler =
50 dyn_cast<gpu::OffloadingLLVMTranslationAttrInterface>(
51 op.getOffloadingHandlerAttr());
52 assert(offloadingHandler &&
"Invalid offloading handler.");
53 return offloadingHandler.embedBinary(op, builder, moduleTranslation);
55 .Case([&](gpu::LaunchFuncOp op) {
58 .Default([&](Operation *op) {
67 registry.
insert<gpu::GPUDialect>();
69 dialect->addInterfaces<GPUDialectLLVMIRTranslationInterface>();
static void launchKernel(sycl::queue *queue, sycl::kernel *kernel, size_t gridX, size_t gridY, size_t gridZ, size_t blockX, size_t blockY, size_t blockZ, size_t sharedMemBytes, void **params, size_t paramsCount)
The DialectRegistry maps a dialect namespace to a constructor for the matching dialect.
bool addExtension(TypeID extensionID, std::unique_ptr< DialectExtensionBase > extension)
Add the given extension to the registry.
Base class for dialect interfaces providing translation to LLVM IR.
LLVMTranslationDialectInterface(Dialect *dialect)
Implementation class for module translation.
MLIRContext is the top-level object for a collection of MLIR operations.
void appendDialectRegistry(const DialectRegistry ®istry)
Append the contents of the given dialect registry to the registry associated with this context.
InFlightDiagnostic emitError(const Twine &message={})
Emit an error about fatal conditions with this operation, reporting up to any diagnostic handlers tha...
OperationName getName()
The name of an operation is the key identifier for it.
static Operation * lookupNearestSymbolFrom(Operation *from, StringAttr symbol)
Returns the operation registered with the given symbol name within the closest parent operation of,...
Include the generated interface declarations.
void registerGPUDialectTranslation(DialectRegistry ®istry)
Register the GPU dialect and the translation from it to the LLVM IR in the given registry;.