[mlir] fix dereferencing of optional sym_name attribute

`sym_name` is an optional attribute of `ModuleOp`, so it is unsafe to
fetch the underlying value without checking whether it is non-empty.
Such unsafe dereferencing causes the lower-host-to-llvm-calls_fail.mlir
test to segfault.  Although this bug existed for four months, it wasn't
triggered, since previous tests executed a code path that used a default
value instead of one fetched from the module attribute.

This patch makes the code use a default value if the optional attribute
does not have a value.

Reviewed By: stella.stamenova

Differential Revision: https://reviews.llvm.org/D141941

diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
index 2017d55..08512e9 100644 (file)
--- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
@@ -126,7 +126,7 @@ static LogicalResult getKernelGlobalVariables(
 /// Encodes the SPIR-V module's symbolic name into the name of the entry point
 /// function.
 static LogicalResult encodeKernelName(spirv::ModuleOp module) {
-  StringRef spvModuleName = *module.getSymName();
+  StringRef spvModuleName = module.getSymName().value_or(kSPIRVModule);
   // We already know that the module contains exactly one entry point function
   // based on `getKernelGlobalVariables()` call. Update this function's name
   // to:

diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
index f46b23c..550eca3 100644 (file)
--- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
+// RUN: mlir-opt --lower-host-to-llvm %s -split-input-file | FileCheck %s
 
 module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
 
@@ -45,3 +45,35 @@ module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#s
     return
   }
 }
+
+// -----
+
+// Check using a specified sym_name attribute.
+module {
+  spirv.module Logical GLSL450 attributes { sym_name = "spirv.sym" } {
+    // CHECK: spirv.func @spirv.sym_bar
+    // CHECK: spirv.EntryPoint "GLCompute" @spirv.sym_bar
+    // CHECK: spirv.ExecutionMode @spirv.sym_bar "LocalSize", 1, 1, 1
+    spirv.func @bar() "None" {
+      spirv.Return
+    }
+    spirv.EntryPoint "GLCompute" @bar
+    spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
+  }
+}
+
+// -----
+
+// Check using the default sym_name attribute.
+module {
+  spirv.module Logical GLSL450 {
+    // CHECK: spirv.func @__spv___bar
+    // CHECK: spirv.EntryPoint "GLCompute" @__spv___bar
+    // CHECK: spirv.ExecutionMode @__spv___bar "LocalSize", 1, 1, 1
+    spirv.func @bar() "None" {
+      spirv.Return
+    }
+    spirv.EntryPoint "GLCompute" @bar
+    spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
+  }
+}