[DeviceSanitizer] Fix device global type of KernelMetadata (#16357)

OpenCL CPU requires the type of device global wraps with structure type

Author: AllanZyne

llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp

@@ -38,6 +38,29 @@ PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
   auto &DL = M.getDataLayout();
   auto &Ctx = M.getContext();
 
+  // Fix device global type, by wrapping a structure type
+  {
+    assert(KernelMetadata->getValueType()->isArrayTy());
+
+    auto *KernelMetadataOld = KernelMetadata;
+
+    StructType *StructTypeWithArray = StructType::create(Ctx);
+    StructTypeWithArray->setBody(KernelMetadataOld->getValueType());
+
+    KernelMetadata = new GlobalVariable(
+        M, StructTypeWithArray, false, GlobalValue::ExternalLinkage,
+        ConstantStruct::get(StructTypeWithArray,
+                            KernelMetadataOld->getInitializer()),
+        "", nullptr, GlobalValue::NotThreadLocal, 1); // Global AddressSpace
+    KernelMetadata->takeName(KernelMetadataOld);
+    KernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+    KernelMetadata->setDSOLocal(true);
+    KernelMetadata->copyAttributesFrom(KernelMetadataOld);
+    KernelMetadata->copyMetadata(KernelMetadataOld, 0);
+
+    KernelMetadataOld->eraseFromParent();
+  }
+
   // Fix attributes
   KernelMetadata->addAttribute(
       "sycl-device-global-size",

llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll

@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 
 @__asan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
 @__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__asan_kernel to i64), i64 54 }] #2
-; CHECK-IR: @__AsanKernelMetadata {{.*}} !spirv.Decorations
+; CHECK-IR: @__AsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
 

llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll

@@ -18,7 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 
 @__msan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
 @__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }] #0
-; CHECK-IR: @__MsanKernelMetadata {{.*}} !spirv.Decorations
+; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"