[mlir] fix dereferencing of optional sym_name attribute

[ashay-github]

Summary

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.

[mehdi_amini]

Can you add a test for this please?

[ashay-github]

Added unit tests to check that lowering using either the specified or the
default attribute value.

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

if (hasDescriptorSetAndBinding(globalOp))
       globalVariableMap[calculateGlobalIndex(globalOp)] = globalOp;
   }
   return success();
 }
 
 /// 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:
   //   {spv_module_name}_{function_name}
   auto entryPoints = module.getOps<spirv::EntryPointOp>();
   if (!llvm::hasSingleElement(entryPoints)) {
     return module.emitError(
         "The module must contain exactly one entry point function");

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

-// 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]>>} {
 
   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)>
   //       CHECK: llvm.func @__spv__foo_bar()
 
   //       CHECK: spirv.module @__spv__foo
   //       CHECK:   spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
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]>>} {
   func.func @main() {
     %buffer = memref.alloc() : memref<6xi32>
     %one = arith.constant 1 : index
     gpu.launch_func @foo::@bar blocks in (%one, %one, %one)
         threads in (%one, %one, %one) args(%buffer : memref<6xi32>)
     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
+  }
+}