[SYCL][DeviceSanitizer] Use -asan-constructor-kind=none to disable ctor/dtor (#13259)

The removed code fails to prevent ctor/dtor generation.

Author: wenju-he

clang/lib/Driver/SanitizerArgs.cpp

@@ -1151,6 +1151,10 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
       CmdArgs.push_back("-mllvm");
       CmdArgs.push_back("-asan-instrumentation-with-call-threshold=0");
 
+      // asan initialization is done in unified runtime rather than in ctor.
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-asan-constructor-kind=none");
+
       CmdArgs.push_back("-mllvm");
       CmdArgs.push_back("-asan-stack=0");
       CmdArgs.push_back("-mllvm");

clang/test/Driver/sycl-device-sanitizer.cpp

@@ -5,6 +5,7 @@
 // SYCL-ASAN-SAME: -fsanitize-address-use-after-return=never
 // SYCL-ASAN-SAME: -fno-sanitize-address-use-after-scope
 // SYCL-ASAN-SAME: "-mllvm" "-asan-instrumentation-with-call-threshold=0"
+// SYCL-ASAN-SAME: "-mllvm" "-asan-constructor-kind=none"
 // SYCL-ASAN-SAME: "-mllvm" "-asan-stack=0"
 // SYCL-ASAN-SAME: "-mllvm" "-asan-globals=0"
 

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -2911,12 +2911,6 @@ bool ModuleAddressSanitizer::instrumentModule(Module &M) {
     }
   }
 
-  // SPIR kernel needn't AsanCtorFunction & AsanDtorFunction
-  if (TargetTriple.isSPIR()) {
-    AsanCtorFunction = nullptr;
-    AsanDtorFunction = nullptr;
-  }
-
   const uint64_t Priority = GetCtorAndDtorPriority(TargetTriple);
 
   // Put the constructor and destructor in comdat if both

sycl/test/check_device_code/device_asan_no_ctor.cpp

@@ -0,0 +1,23 @@
+// REQUIRES: linux
+// RUN: %clangxx -fsycl-device-only -fno-sycl-early-optimizations -Xarch_device -fsanitize=address -emit-llvm %s -S -o %t.ll
+// RUN: FileCheck %s --input-file %t.ll
+
+// Check asan ctor is not generated since device asan initialization is done in
+// unified runtime rather than in ctor.
+
+// CHECK-NOT: asan.module_ctor
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+  sycl::queue queue;
+  {
+    queue.submit(
+        [&](sycl::handler &cgh) { cgh.single_task<class t>([=]() {}); });
+    queue.wait();
+  }
+
+  return 0;
+}