[flang-commits] [flang] [MLIR] Enabling Intel GPU Integration. (PR #65539)

[Fabian Mora via flang-commits]

================
@@ -0,0 +1,70 @@
+//===- SerializeToSPIRV.cpp - Convert GPU kernel to SPIRV blob -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This pass iterates all the SPIR-V modules in the top module and serializes
+/// each SPIR-V module to SPIR-V binary and then attachs the binary blob as a
+/// string attribute to the corresponding gpu module.
+///
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
+#include "mlir/Target/SPIRV/Serialization.h"
+
+namespace mlir {
+#define GEN_PASS_DEF_GPUSERIALIZETOSPIRVPASS
+#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"
+} // namespace mlir
+
+using namespace mlir;
+
+struct GpuSerializeToSPIRVPass : public mlir::impl::GpuSerializeToSPIRVPassBase<GpuSerializeToSPIRVPass> {
+public:
+  void runOnOperation() override {
+    auto mod = getOperation();
+    llvm::SmallVector<uint32_t, 0> spvBinary;
+    for (mlir::gpu::GPUModuleOp gpuMod : mod.getOps<gpu::GPUModuleOp>()) {
+      auto name = gpuMod.getName();
+      // check that the spv module has the same name with gpu module except the
+      // prefix "__spv__"
+      auto isSameMod = [&](spirv::ModuleOp spvMod) -> bool {
+        auto spvModName = spvMod.getName();
+        return spvModName->consume_front("__spv__") && spvModName == name;
+      };
+      auto spvMods = mod.getOps<spirv::ModuleOp>();
+      auto it = llvm::find_if(spvMods, isSameMod);
+      if (it == spvMods.end()) {
+        gpuMod.emitError() << "Unable to find corresponding SPIR-V module";
+        signalPassFailure();
+        return;
+      }
+      auto spvMod = *it;
+
+      spvBinary.clear();
+      // serialize the spv module to spv binary
+      if (mlir::failed(spirv::serialize(spvMod, spvBinary))) {
+        spvMod.emitError() << "Failed to serialize SPIR-V module";
+        signalPassFailure();
+        return;
+      }
+
+      // attach the spv binary to the gpu module
+      auto spvData =
+          llvm::StringRef(reinterpret_cast<const char *>(spvBinary.data()),
+                          spvBinary.size() * sizeof(uint32_t));
+      auto spvAttr = mlir::StringAttr::get(&getContext(), spvData);
+      gpuMod->setAttr(gpu::getDefaultGpuBinaryAnnotation(), spvAttr);
+      spvMod->erase();
+    }
+  }
+};
----------------
fabianmcg wrote:

@silee2 here are the steps:
1. Implement a target attribute, see for example: [NVVMTargetAttr](https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td#L1679-L1741). The idea of this attribute is to hold properties intrinsic to the target, like triple, chip, flags, etc.
2.  Add a pass to attach the target to a module, see: [GpuNVVMAttachTarget](https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td#L85-L128) and [Dialect/GPU/Transforms/NVVMAttachTarget.cpp](https://github.com/llvm/llvm-project/blob/main/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp). The idea of this pass is to attach the SPIRV target to GPU modules, so it must know how to create them.
3. We're currently implementing `TargetAttrs` as external models to keep libraries separated, see [NVVM/Target.cpp](https://github.com/llvm/llvm-project/blob/main/mlir/lib/Target/LLVM/NVVM/Target.cpp#L44-L50), so `GpuSerializeToSPIRVPass::run` would be there.
4. Modify `getModuleLoadFn` & `createKernelLaunch` appropriately in [SelectObjectAttr.cpp#L125-L15](https://github.com/llvm/llvm-project/blob/main/mlir/lib/Target/LLVMIR/Dialect/GPU/SelectObjectAttr.cpp#L125-L152) instead of adding the changes in `GPUToLLVMCommon`.
5. Then the compilation workflow should look something similar to this: [GPU: Compilation Overview](https://mlir.llvm.org/docs/Dialects/GPU/#compilation-overview)

I'll take care of adding a pointer to the top module symbol table so it can be used be the `SPIRVTarget`.

If you have any questions just ping me in discord or discourse `@fabianmc`.

https://github.com/llvm/llvm-project/pull/65539