[PATCH] D72696: [mlir][spirv] mlir-vulkan-runner

Wed Feb 12 18:20:07 PST 2020

antiagainst requested changes to this revision.
antiagainst added a comment.
This revision now requires changes to proceed.

Super nice! This is so great! Thanks a ton, Denis! I don't have major concerns given that a large portion of this patch has already went through reviews multiple times previously. Just a few nits. We can land and then iterate on it. :)

================
Comment at: mlir/include/mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h:9
+//
+// The file declares a pass to convert GPU to Vulkan runtime calls.
+//
----------------
Nit: "This file declares passes to convert GPU dialect ops to Vulkan ..." so later if we add more entry points we don't need to update this.

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:10
+// This file implements a pass to convert gpu.launch_func op into a sequence of
+// Vulkan runtime calls. The Vulkan runtime API is huge, so currently we don't
+// instrument a host part with calls to Vulkan API, instead we istrument a host
----------------
The Vulkan runtime API surface is huge so currently we don't expose separate external functions in IR for each of them, instead we expose a few external functions to wrapper libraries which ...

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:72
+                                   std::vector<char> &binaryShader);
+  Value createEntryPointNameConstant(StringRef name, Location loc,
+                                     OpBuilder &builder);
----------------
Insert an empty line and document this?

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:75
+
+  /// Creates a LLVM constant for each number of local workgroup and
+  /// populates the given `numWorkGroups`.
----------------
for each dimension of 

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:83
+  void declareVulkanFunctions(Location loc);
+  void translateGpuLaunchCalls(mlir::gpu::LaunchFuncOp launchOp);
+
----------------
Insert an empty line and document this?

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:86
+public:
+  void runOnModule() override {
+    llvmDialect = getContext().getRegisteredDialect<LLVM::LLVMDialect>();
----------------
Could we move the impl out side of the declaration like other large methods?

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:165
+  for (auto spirvModule : module.getOps<spirv::ModuleOp>()) {
+    if (done) {
+      return failure();
----------------
Print an errror saying "should only contain one 'spv.module' op"?

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:208
+  // Declare runtime functions.
+  declareVulkanFunctions(loc);
+
----------------
What about moving this to the end so that if we have errors we don't need to do this at all?

================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:267
+    pass("launch-func-to-vulkan",
+         "Convert all launch_func ops to Vulkan runtime calls");
----------------
gpu.launch_func

================
Comment at: mlir/test/Conversion/GPUToVulkan/simple.mlir:1
+// RUN: mlir-opt %s -launch-func-to-vulkan | FileCheck %s
+
----------------
Can we have a more descriptive name? simple.mlir does not actually say anything. What about `invoke-vulkan.mlir`?

================
Comment at: mlir/test/Conversion/GPUToVulkan/simple.mlir:14
+    spv.globalVariable @kernel_1_arg_1 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
+    func @kernel_1() attributes {workgroup_attributions = 0 : i64} {
+      %0 = spv._address_of @kernel_1_arg_1 : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
----------------
Super nit: just call it `kernel` would be fine. The `_1` does not provide additional value.

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:57
+    break;
+  default:
+    return failure();
----------------
Print an error saying "unsupported storage class"?

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:72
+    break;
+  default:
+    return failure();
----------------
Same here

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:84
+        memorySize += resourceDataBindingPair.second.size;
+      } else {
+        return failure();
----------------
Print an error?

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:133
+  // Wait for device.
+  RETURN_ON_VULKAN_ERROR(vkDeviceWaitIdle(device), "vkDeviceWaitIdle");
+  vkDestroyDevice(device, nullptr);
----------------
Move to the beginning to be safe?

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:645
+  commandBufferBeginInfo.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
+  commandBufferBeginInfo.pInheritanceInfo = 0;
+
----------------
Prefer to use `nullptr` for pointer fields.

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.cpp:669
+LogicalResult VulkanRuntime::submitCommandBuffersToQueue() {
+  VkSubmitInfo submitInfo;
+  submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
----------------
`= {}` 

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.h:7
+//
+// =============================================================================
+//
----------------
This line is a bit different. ;)

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.h:68
+/// VulkanHostMemoryBuffer mapped into a descriptor set and a binding.
+using ResourceData =
+    llvm::DenseMap<DescriptorSetIndex,
----------------
What about ResourceBufferBindingMap?

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.h:73
+/// StorageClass mapped into a descriptor set and a binding.
+using ResourceStorageClassData =
+    llvm::DenseMap<DescriptorSetIndex,
----------------
ResourceStorageClassBindingMap?

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.h:89
+/// Vulkan runtime.
+/// The purpose of this class is to run SPIR-V computation shader on Vulkan
+/// device.
----------------
compute shader

================
Comment at: mlir/tools/mlir-vulkan-runner/VulkanRuntime.h:92
+/// Before the run, user must provide and set resource data with descriptors,
+/// spir-v shader, number of work groups and entry point. After the creation of
+/// VulkanRuntime, special methods must be called in the following
----------------
SPIR-V

================
Comment at: mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp:28
+
+static LogicalResult runMLIRPasses(ModuleOp m) {
+  PassManager pm(m.getContext());
----------------
Let's name it as `module`. `m` and other single char names should really have a very small scope.

================
Comment at: mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp:10
+// Implements C runtime wrappers around the VulkanRuntime.
+// Also adds VulkanRuntimeManager class to manage VulkanRuntime.
+//
----------------
We can delete this line given it's implementation details.

================
Comment at: mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp:22
+/// wrappers. It's designed to handle single SPIR-V compute shader.
+class VulkanRuntimeManager {
+  public:
----------------
Wrap this class in anonymous namespace to hide?

================
Comment at: mlir/tools/mlir-vulkan-runner/vulkan-runtime-wrappers.cpp:28
+
+    static VulkanRuntimeManager *instance() {
+      static VulkanRuntimeManager *runtimeManager = new VulkanRuntimeManager;
----------------
Hmm, instead of having this singleton, what about exposing two more functions, `initVulkan` and `deinitVulkan` to be more explicit? `initVulkan` can return a pointer that gets passed into other functions like `setResourceData`, etc. This file will be compiled into libraries so having this kind of `static` stuff is not friendly. I'm fine to use `llvm::ManagedStatic` for now to avoid huge changes again given this patch is already fair large and has been waited for a long time. We can put a TODO here and add `initVulkan` and `deinitVulkan` later.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D72696/new/

https://reviews.llvm.org/D72696