15 #include "llvm/ADT/TypeSwitch.h"
21 llvm::IRBuilderBase &builder,
23 auto kernelBinary = SymbolTable::lookupNearestSymbolFrom<gpu::BinaryOp>(
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,
38 class GPUDialectLLVMIRTranslationInterface
44 convertOperation(
Operation *operation, llvm::IRBuilderBase &builder,
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) {
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.
Operation is the basic unit of execution within MLIR.
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.
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;.