[PATCH] D74549: [mlir][spirv] Add pass ConvertGpuLaunchFuncToVulkanCallsPass.

River Riddle via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Feb 13 11:19:45 PST 2020


rriddle added inline comments.


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:83
+                                    mlir::gpu::LaunchFuncOp launchOp,
+                                    SmallVector<Value, 3> &numWorkGroups);
+
----------------
SmallVectorImpl


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:156
+    StringRef name, Location loc, OpBuilder &builder) {
+  std::vector<char> shaderName(name.begin(), name.end());
+  // Append `\0` to follow C style string given that LLVM::createGlobalString()
----------------
SmallString ? That would remove the explicit StringRef conversion below.


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:162
+  std::string entryPointGlobalName =
+      std::string(llvm::formatv("{0}_spv_entry_point_name", name));
+  return LLVM::createGlobalString(
----------------
Seems easier to just write `(name + "_spv_entry_point_name").str()`


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:175
+    if (done) {
+      spirvModule.emitError("should only contain one 'spv.module' op");
+      return failure();
----------------
Return the error directly it auto converts to failure.


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:179
+    done = true;
+    if (failed(spirv::serialize(spirvModule, binary))) {
+      return failure();
----------------
Drop all of the trivial braces.


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:197
+
+    if (!numWorkGroupDimConstant) {
+      return failure();
----------------
Same here.


================
Comment at: mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp:210
+
+// Translates gpu launch op to the sequence of Vulkan runtime calls.
+void GpuLaunchFuncToVulkanCalssPass::translateGpuLaunchCalls(
----------------
///


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D74549





More information about the llvm-commits mailing list