[SYCL][DeviceSanitizer] Add "asanUsed" property into "SYCL/misc properties" and driver support (#11989)

If the kernel is built with the flag "-fsanitize=address", we add a new
property "asanUsed" into "SYCL/misc properties", so that we can enable
the "UR_LAYER_ASAN" layer in the UR loader at piPluginInit.

Author: AllanZyne

clang/lib/Driver/SanitizerArgs.cpp

@@ -1137,6 +1137,24 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
       return;
     GPUSanitize = true;
   }
+  // SPIR sanitizer support is experimental and will pass a fixed set of flags
+  if (TC.getTriple().isSPIR()) {
+    if (Sanitizers.has(SanitizerKind::Address)) {
+      CmdArgs.push_back("-fsanitize=address");
+      CmdArgs.push_back("-fsanitize-address-use-after-return=never");
+      CmdArgs.push_back("-fno-sanitize-address-use-after-scope");
+
+      // -fsanitize-address-outline-instrumentation
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-asan-instrumentation-with-call-threshold=0");
+
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-asan-stack=0");
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-asan-globals=0");
+    }
+    return;
+  }
 
   // Translate available CoverageFeatures to corresponding clang-cc1 flags.
   // Do it even if Sanitizers.empty() since some forms of coverage don't require

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5388,8 +5388,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
 
     // Add -ffine-grained-bitfield-accesses option. This will be added
     // only for SPIR based targets.
-    if (Triple.isSPIR())
-      CmdArgs.push_back("-ffine-grained-bitfield-accesses");
+    if (Triple.isSPIR()) {
+      // It cannot be enabled together with a sanitizer
+      if (!Args.getLastArg(options::OPT_fsanitize_EQ))
+        CmdArgs.push_back("-ffine-grained-bitfield-accesses");
+    }
 
     if (!Args.hasFlag(options::OPT_fsycl_unnamed_lambda,
                       options::OPT_fno_sycl_unnamed_lambda, true))

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1017,11 +1017,18 @@ SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
   getProgramPaths().push_back(getDriver().Dir);
 
   // Diagnose unsupported options only once.
-  // All sanitizer options are not currently supported.
-  for (auto A :
-       Args.filtered(options::OPT_fsanitize_EQ, options::OPT_fcf_protection_EQ))
+  // All sanitizer options are not currently supported, except AddressSanitizer
+  for (auto *A : Args.filtered(options::OPT_fsanitize_EQ,
+                               options::OPT_fcf_protection_EQ)) {
+    if (A->getOption().getID() == options::OPT_fsanitize_EQ &&
+        A->getValues().size() == 1) {
+      std::string SanitizeVal = A->getValue();
+      if (SanitizeVal == "address")
+        continue;
+    }
     D.getDiags().Report(clang::diag::warn_drv_unsupported_option_for_target)
         << A->getAsString(Args) << getTriple().str();
+  }
 }
 
 void SYCLToolChain::addClangTargetOptions(
@@ -1049,6 +1056,15 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
     auto Opt(A->getOption().getID());
     switch (Opt) {
     case options::OPT_fsanitize_EQ:
+      if (A->getValues().size() == 1) {
+        std::string SanitizeVal = A->getValue();
+        if (SanitizeVal == "address") {
+          if (IsNewDAL)
+            DAL->append(A);
+          continue;
+        }
+      }
+      [[fallthrough]];
     case options::OPT_fcf_protection_EQ:
       if (!IsNewDAL)
         DAL->eraseArg(Opt);
@@ -1428,3 +1444,7 @@ void SYCLToolChain::AddClangCXXStdlibIncludeArgs(const ArgList &Args,
                                                  ArgStringList &CC1Args) const {
   HostTC.AddClangCXXStdlibIncludeArgs(Args, CC1Args);
 }
+
+SanitizerMask SYCLToolChain::getSupportedSanitizers() const {
+  return SanitizerKind::Address;
+}

clang/lib/Driver/ToolChains/SYCL.h

@@ -215,6 +215,8 @@ class LLVM_LIBRARY_VISIBILITY SYCLToolChain : public ToolChain {
       const llvm::opt::ArgList &Args,
       llvm::opt::ArgStringList &CC1Args) const override;
 
+  SanitizerMask getSupportedSanitizers() const override;
+
   const ToolChain &HostTC;
   const bool IsSYCLNativeCPU;
 

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

@@ -0,0 +1,15 @@
+// RUN: %clangxx -fsycl -fsanitize=address -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-ASAN %s
+// SYCL-ASAN: clang{{.*}} "-fsycl-is-device"
+// SYCL-ASAN-SAME: -fsanitize=address
+// 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-stack=0"
+// SYCL-ASAN-SAME: "-mllvm" "-asan-globals=0"
+
+// RUN: %clangxx -fsycl -fsanitize=address -mllvm -asan-stack=1 -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-ASAN-FILTER %s
+// SYCL-ASAN-FILTER: clang{{.*}} "-fsycl-is-device"
+// SYCL-ASAN-FILTER-SAME: -fsanitize=address
+// SYCL-ASAN-FILTER-SAME: "-mllvm" "-asan-stack=0"

clang/test/Driver/sycl-unsupported.cpp

@@ -1,14 +1,4 @@
 /// Diagnose unsupported options specific to SYCL compilations
-// RUN: %clangxx -fsycl -fsanitize=address -fsycl-targets=spir64 -### %s 2>&1 \
-// RUN:  | FileCheck %s -DARCH=spir64 -DOPT=-fsanitize=address
-// RUN: %clang_cl -fsycl -fsanitize=address -fsycl-targets=spir64 -### %s 2>&1 \
-// RUN:  | FileCheck %s -DARCH=spir64 -DOPT=-fsanitize=address
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -fsanitize=address -### %s 2>&1 \
-// RUN:  | FileCheck %s -DARCH=spir64_gen -DOPT=-fsanitize=address
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga -fsanitize=address -### %s 2>&1 \
-// RUN:  | FileCheck %s -DARCH=spir64_fpga -DOPT=-fsanitize=address
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 -fsanitize=address -### %s 2>&1 \
-// RUN:  | FileCheck %s -DARCH=spir64_x86_64 -DOPT=-fsanitize=address
 
 // RUN: %clangxx -fsycl -fcf-protection -fsycl-targets=spir64 -### %s 2>&1 \
 // RUN:  | FileCheck %s -DARCH=spir64 -DOPT=-fcf-protection

llvm/test/tools/sycl-post-link/sycl-sanitize/asan.ll

@@ -0,0 +1,66 @@
+; This test checks that the post-link tool properly generates "asanUsed=1"
+; in [SYCL/misc properties]
+
+; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.prop
+; CHECK: [SYCL/misc properties]
+; CHECK: asanUsed=1
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%struct.AssertHappened.10 = type { i32, [257 x i8], [257 x i8], [129 x i8], i32, i64, i64, i64, i64, i64, i64 }
+%"class.sycl::_V1::range.12" = type { %"class.sycl::_V1::detail::array.11" }
+%"class.sycl::_V1::detail::array.11" = type { [2 x i64] }
+%"class.sycl::_V1::detail::RoundedRangeIDGenerator.15" = type <{ %"class.sycl::_V1::range.12", %"class.sycl::_V1::range.12", %"class.sycl::_V1::range.12", %"class.sycl::_V1::range.12", i8, [7 x i8] }>
+%"class.sycl::_V1::detail::RoundedRangeKernel.17" = type <{ %"class.sycl::_V1::range.12", %class.anon.16, [7 x i8] }>
+%class.anon.16 = type { i8 }
+%struct.__devicelib_div_t_32.13 = type { i32, i32 }
+%struct.__devicelib_div_t_64.14 = type { i64, i64 }
+
+@llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @asan.module_ctor, ptr null }]
+@__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@SPIR_AssertHappenedMem = linkonce_odr dso_local addrspace(1) global %struct.AssertHappened.10 zeroinitializer, align 8
+@__spirv_BuiltInWorkgroupId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__spirv_BuiltInGlobalLinearId = external dso_local local_unnamed_addr addrspace(1) constant i64, align 8
+@__spirv_BuiltInWorkgroupSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+
+; Function Attrs: nounwind uwtable
+define internal void @asan.module_ctor() #0 {
+  call void @__asan_init()
+  call void @__asan_version_mismatch_check_v8()
+  ret void
+}
+
+declare void @__asan_init()
+
+declare void @__asan_version_mismatch_check_v8()
+
+; Function Attrs: mustprogress norecurse nounwind sanitize_address uwtable
+define weak_odr dso_local spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperIZZ4mainENKUlRNS0_7handlerEE_clES4_E9TheKernelEE(ptr noundef byval(%"class.sycl::_V1::range.12") align 8 %_arg_UserRange) local_unnamed_addr #1 {
+entry:
+  ret void
+}
+
+attributes #0 = { nounwind uwtable "frame-pointer"="all" "sycl-optlevel"="2" }
+attributes #1 = { mustprogress norecurse nounwind sanitize_address uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
+
+!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0}
+!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
+!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
+!llvm.module.flags = !{!3, !4, !5}
+!sycl.specialization-constants = !{}
+!sycl.specialization-constants-default-values = !{}
+
+!0 = !{i32 1, i32 2}
+!1 = !{i32 4, i32 100000}
+!2 = !{!"clang version 18.0.0 (https://github.com/AllanZyne/llvm.git 97c052ed8efa30f750dacf8d89e8e64743ec03f7)"}
+!3 = !{i32 1, !"wchar_size", i32 4}
+!4 = !{i32 7, !"uwtable", i32 2}
+!5 = !{i32 7, !"frame-pointer", i32 2}
+!6 = !{i32 8544724}
+!7 = !{i32 -1}
+!8 = !{}
+!9 = !{i1 false, i1 true}

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -312,6 +312,11 @@ std::vector<StringRef> getKernelNamesUsingAssert(const Module &M) {
   return SPIRKernelNames;
 }
 
+bool isModuleUsingAsan(const Module &M) {
+  auto *AsanInitFunction = M.getFunction("__asan_init");
+  return AsanInitFunction;
+}
+
 // Gets reqd_work_group_size information for function Func.
 std::vector<uint32_t> getKernelReqdWorkGroupSizeMetadata(const Function &Func) {
   MDNode *ReqdWorkGroupSizeMD = Func.getMetadata("reqd_work_group_size");
@@ -537,6 +542,11 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
       PropSet[PropSetRegTy::SYCL_ASSERT_USED].insert({FName, true});
   }
 
+  {
+    if (isModuleUsingAsan(M))
+      PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"asanUsed", true});
+  }
+
   if (GlobProps.EmitDeviceGlobalPropSet) {
     // Extract device global maps per module
     auto DevGlobalPropertyMap = collectDeviceGlobalProperties(M);

sycl/include/sycl/detail/pi.h

@@ -2693,6 +2693,13 @@ __SYCL_EXPORT pi_result piextSignalExternalSemaphore(
     pi_uint32 num_events_in_wait_list, const pi_event *event_wait_list,
     pi_event *event);
 
+typedef enum {
+  _PI_SANITIZE_TYPE_NONE = 0x0,
+  _PI_SANITIZE_TYPE_ADDRESS = 0x1,
+  _PI_SANITIZE_TYPE_MEMORY = 0x2,
+  _PI_SANITIZE_TYPE_THREAD = 0x3
+} _pi_sanitize_type;
+
 struct _pi_plugin {
   // PI version supported by host passed to the plugin. The Plugin
   // checks and writes the appropriate Function Pointers in
@@ -2709,6 +2716,8 @@ struct _pi_plugin {
 #define _PI_API(api) decltype(::api) *api;
 #include <sycl/detail/pi.def>
   } PiFunctionTable;
+
+  _pi_sanitize_type SanitizeType;
 };
 
 #ifdef __cplusplus

sycl/plugins/unified_runtime/pi_unified_runtime.cpp

@@ -1302,7 +1302,15 @@ __SYCL_EXPORT pi_result piPluginInit(pi_plugin *PluginInit) {
   strncpy(PluginInit->PluginVersion, SupportedVersion, PluginVersionSize);
 
   // Initialize UR and discover adapters
-  HANDLE_ERRORS(urLoaderInit(0, nullptr));
+  ur_loader_config_handle_t LoaderConfig{};
+  HANDLE_ERRORS(urLoaderConfigCreate(&LoaderConfig));
+
+  if (PluginInit->SanitizeType == _PI_SANITIZE_TYPE_ADDRESS) {
+    HANDLE_ERRORS(urLoaderConfigEnableLayer(LoaderConfig, "UR_LAYER_ASAN"));
+  }
+
+  HANDLE_ERRORS(urLoaderInit(0, LoaderConfig));
+
   uint32_t NumAdapters;
   HANDLE_ERRORS(urAdapterGet(0, nullptr, &NumAdapters));
   if (NumAdapters > 0) {

sycl/source/detail/pi.cpp

@@ -455,10 +455,15 @@ static void initializePlugins(std::vector<PluginPtr> &Plugins) {
   std::vector<std::tuple<std::string, backend, void *>> LoadedPlugins =
       loadPlugins(std::move(PluginNames));
 
+  bool IsAsanUsed = ProgramManager::getInstance().kernelUsesAsan();
+
   for (auto &[Name, Backend, Library] : LoadedPlugins) {
-    std::shared_ptr<PiPlugin> PluginInformation = std::make_shared<PiPlugin>(
-        PiPlugin{_PI_H_VERSION_STRING, _PI_H_VERSION_STRING,
-                 /*Targets=*/nullptr, /*FunctionPointers=*/{}});
+    std::shared_ptr<PiPlugin> PluginInformation =
+        std::make_shared<PiPlugin>(PiPlugin{
+            _PI_H_VERSION_STRING, _PI_H_VERSION_STRING,
+            /*Targets=*/nullptr, /*FunctionPointers=*/{},
+            /*IsAsanUsed*/
+            IsAsanUsed ? _PI_SANITIZE_TYPE_ADDRESS : _PI_SANITIZE_TYPE_NONE});
 
     if (!Library) {
       if (trace(PI_TRACE_ALL)) {

sycl/source/detail/program_manager/program_manager.cpp

@@ -1019,7 +1019,7 @@ loadDeviceLibFallback(const ContextImplPtr Context, DeviceLibExt Extension,
   return LibProg;
 }
 
-ProgramManager::ProgramManager() {
+ProgramManager::ProgramManager() : m_AsanFoundInImage(false) {
   const char *SpvFile = std::getenv(UseSpvEnv);
   // If a SPIR-V file is specified with an environment variable,
   // register the corresponding image
@@ -1437,6 +1437,13 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) {
 
     cacheKernelUsesAssertInfo(*Img);
 
+    // check if kernel uses asan
+    {
+      pi_device_binary_property Prop = Img->getProperty("asanUsed");
+      m_AsanFoundInImage |=
+          Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0);
+    }
+
     // Sort kernel ids for faster search
     std::sort(m_BinImg2KernelIDs[Img.get()]->begin(),
               m_BinImg2KernelIDs[Img.get()]->end(), LessByHash<kernel_id>{});

sycl/source/detail/program_manager/program_manager.hpp

@@ -291,6 +291,8 @@ class ProgramManager {
 
   bool kernelUsesAssert(const std::string &KernelName) const;
 
+  bool kernelUsesAsan() const { return m_AsanFoundInImage; }
+
   std::set<RTDeviceBinaryImage *>
   getRawDeviceImages(const std::vector<kernel_id> &KernelIDs);
 
@@ -399,6 +401,9 @@ class ProgramManager {
 
   std::set<std::string> m_KernelUsesAssert;
 
+  // True iff there is a device image compiled with AddressSanitizer
+  bool m_AsanFoundInImage;
+
   // Maps between device_global identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<DeviceGlobalMapEntry>>
       m_DeviceGlobals;