[DeviceSanitizer] Checking out-of-bounds error on sycl::local_accessor

libdevice/include/asan_libdevice.hpp

@@ -60,7 +60,6 @@ struct DeviceSanitizerReport {
 };
 
 struct LocalArgsInfo {
-  uint32_t ArgIndex = 0;
   uint64_t Size = 0;
   uint64_t SizeWithRedZone = 0;
 };

libdevice/include/spir_global_var.hpp

@@ -15,3 +15,14 @@
 #define __SYCL_LOCAL__ __attribute__((opencl_local))
 #define __SYCL_PRIVATE__ __attribute__((opencl_private))
 #define __SYCL_CONSTANT__ __attribute__((opencl_constant))
+
+#ifndef SPIR_GLOBAL_VAR
+#ifdef __SYCL_DEVICE_ONLY__
+#define SPIR_GLOBAL_VAR __attribute__((sycl_global_var))
+#else
+#define SPIR_GLOBAL_VAR
+#endif
+#endif
+
+extern SPIR_GLOBAL_VAR __SYCL_GLOBAL__ uint64_t *__SYCL_LOCAL__
+    __AsanLaunchInfo;

libdevice/sanitizer_utils.cpp

@@ -23,11 +23,11 @@ using u16 = unsigned short;
 
 DeviceGlobal<uptr> __AsanShadowMemoryGlobalStart;
 DeviceGlobal<uptr> __AsanShadowMemoryGlobalEnd;
-DeviceGlobal<uptr> __AsanShadowMemoryLocalStart;
-DeviceGlobal<uptr> __AsanShadowMemoryLocalEnd;
 DeviceGlobal<DeviceType> __DeviceType;
 DeviceGlobal<uint64_t> __AsanDebug;
-DeviceGlobal<DeviceSanitizerReport> __DeviceSanitizerReportMem;
+
+// Save the pointer to LaunchInfo
+__SYCL_GLOBAL__ uptr *__SYCL_LOCAL__ __AsanLaunchInfo;
 
 #if defined(__SPIR__) || defined(__SPIRV__)
 
@@ -134,6 +134,16 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
   return shadow_ptr;
 }
 
+static __SYCL_CONSTANT__ const char __mem_launch_info[] =
+    "[kernel] launch_info: %p (offset=%p, offset_end=%p, argLocalArgs=%d, "
+    "localArgs=%p)\n";
+
+static __SYCL_CONSTANT__ const char __generic_to[] =
+    "[kernel] %p(4) - %p(%d)\n";
+
+static __SYCL_CONSTANT__ const char __generic_to_fail[] =
+    "[kernel] %p(4) - unknown address space\n";
+
 inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
 
   if (as == ADDRESS_SPACE_GENERIC) {
@@ -172,9 +182,6 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
     }
     return shadow_ptr;
   } else if (as == ADDRESS_SPACE_LOCAL) { // local
-    if (__AsanShadowMemoryLocalStart == 0) {
-      return 0;
-    }
     // The size of SLM is 128KB on PVC
     constexpr unsigned SLM_SIZE = 128 * 1024;
     // work-group linear id
@@ -184,6 +191,19 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
         __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z +
         __spirv_BuiltInWorkgroupId.z;
 
+    auto launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
+    const auto __AsanShadowMemoryLocalStart = launch_info->LocalShadowOffset;
+    const auto __AsanShadowMemoryLocalEnd = launch_info->LocalShadowOffsetEnd;
+
+    if (__AsanShadowMemoryLocalStart == 0) {
+      return 0;
+    }
+
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_launch_info, __AsanShadowMemoryLocalStart,
+                         __AsanShadowMemoryLocalEnd, launch_info->NumLocalArgs,
+                         launch_info->LocalArgs);
+
     uptr shadow_ptr = __AsanShadowMemoryLocalStart +
                       ((wg_lid * SLM_SIZE) >> ASAN_SHADOW_SCALE) +
                       ((addr & (SLM_SIZE - 1)) >> 3);
@@ -268,22 +288,18 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) {
 bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
   const int Expected = ASAN_REPORT_NONE;
   int Desired = ASAN_REPORT_START;
-  if (atomicCompareAndSet(&__DeviceSanitizerReportMem.get().Flag, Desired,
-                          Expected) == Expected) {
-    __DeviceSanitizerReportMem.get().ErrorType = error_type;
+  auto &SanitizerReport =
+      ((__SYCL_GLOBAL__ LaunchInfo *)__AsanLaunchInfo)->SanitizerReport;
+  if (atomicCompareAndSet(&SanitizerReport.Flag, Desired, Expected) ==
+      Expected) {
+    SanitizerReport.ErrorType = error_type;
     // Show we've done copying
-    atomicStore(&__DeviceSanitizerReportMem.get().Flag, ASAN_REPORT_FINISH);
+    atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH);
     return true;
   }
   return false;
 }
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define __DEVICE_SANITIZER_REPORT_ACCESSOR __DeviceSanitizerReportMem.get()
-#else // __SYCL_DEVICE_ONLY__
-#define __DEVICE_SANITIZER_REPORT_ACCESSOR
-#endif // __SYCL_DEVICE_ONLY__
-
 bool __asan_internal_report_save(
     uptr ptr, uint32_t as, const char __SYCL_CONSTANT__ *file, uint32_t line,
     const char __SYCL_CONSTANT__ *func, bool is_write, uint32_t access_size,
@@ -292,8 +308,12 @@ bool __asan_internal_report_save(
 
   const int Expected = ASAN_REPORT_NONE;
   int Desired = ASAN_REPORT_START;
-  if (atomicCompareAndSet(&__DEVICE_SANITIZER_REPORT_ACCESSOR.Flag, Desired,
-                          Expected) == Expected) {
+
+  auto &SanitizerReport =
+      ((__SYCL_GLOBAL__ LaunchInfo *)__AsanLaunchInfo)->SanitizerReport;
+
+  if (atomicCompareAndSet(&SanitizerReport.Flag, Desired, Expected) ==
+      Expected) {
 
     int FileLength = 0;
     int FuncLength = 0;
@@ -305,39 +325,40 @@ bool __asan_internal_report_save(
       for (auto *C = func; *C != '\0'; ++C, ++FuncLength)
         ;
 
-    int MaxFileIdx = sizeof(__DEVICE_SANITIZER_REPORT_ACCESSOR.File) - 1;
-    int MaxFuncIdx = sizeof(__DEVICE_SANITIZER_REPORT_ACCESSOR.Func) - 1;
+    int MaxFileIdx = sizeof(SanitizerReport.File) - 1;
+    int MaxFuncIdx = sizeof(SanitizerReport.Func) - 1;
 
     if (FileLength < MaxFileIdx)
       MaxFileIdx = FileLength;
     if (FuncLength < MaxFuncIdx)
       MaxFuncIdx = FuncLength;
 
     for (int Idx = 0; Idx < MaxFileIdx; ++Idx)
-      __DEVICE_SANITIZER_REPORT_ACCESSOR.File[Idx] = file[Idx];
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.File[MaxFileIdx] = '\0';
+      SanitizerReport.File[Idx] = file[Idx];
+    SanitizerReport.File[MaxFileIdx] = '\0';
 
     for (int Idx = 0; Idx < MaxFuncIdx; ++Idx)
-      __DEVICE_SANITIZER_REPORT_ACCESSOR.Func[Idx] = func[Idx];
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.Func[MaxFuncIdx] = '\0';
-
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.Line = line;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.GID0 = __spirv_GlobalInvocationId_x();
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.GID1 = __spirv_GlobalInvocationId_y();
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.GID2 = __spirv_GlobalInvocationId_z();
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.LID0 = __spirv_LocalInvocationId_x();
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.LID1 = __spirv_LocalInvocationId_y();
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.LID2 = __spirv_LocalInvocationId_z();
-
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.Address = ptr;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.IsWrite = is_write;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.AccessSize = access_size;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.ErrorType = error_type;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.MemoryType = memory_type;
-    __DEVICE_SANITIZER_REPORT_ACCESSOR.IsRecover = is_recover;
+      SanitizerReport.Func[Idx] = func[Idx];
+    SanitizerReport.Func[MaxFuncIdx] = '\0';
+
+    SanitizerReport.Line = line;
+    SanitizerReport.GID0 = __spirv_GlobalInvocationId_x();
+    SanitizerReport.GID1 = __spirv_GlobalInvocationId_y();
+    SanitizerReport.GID2 = __spirv_GlobalInvocationId_z();
+    SanitizerReport.LID0 = __spirv_LocalInvocationId_x();
+    SanitizerReport.LID1 = __spirv_LocalInvocationId_y();
+    SanitizerReport.LID2 = __spirv_LocalInvocationId_z();
+
+    SanitizerReport.Address = ptr;
+    SanitizerReport.IsWrite = is_write;
+    SanitizerReport.AccessSize = access_size;
+    SanitizerReport.ErrorType = error_type;
+    SanitizerReport.MemoryType = memory_type;
+    SanitizerReport.IsRecover = is_recover;
 
     // Show we've done copying
-    atomicStore(&__DEVICE_SANITIZER_REPORT_ACCESSOR.Flag, ASAN_REPORT_FINISH);
+    atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH);
+    return true;
   }
   return false;
 }
@@ -545,7 +566,7 @@ ASAN_REPORT_ERROR(store, true, 4)
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size(                           \
       uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
       uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
-    u##size *shadow_address = (u##size *)MemToShadow(addr, as);                \
+    auto *shadow_address = (__SYCL_GLOBAL__ u##size *)MemToShadow(addr, as);   \
     if (shadow_address && *shadow_address) {                                   \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
                                  func);                                        \
@@ -554,7 +575,7 @@ ASAN_REPORT_ERROR(store, true, 4)
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort(                 \
       uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file,              \
       uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
-    u##size *shadow_address = (u##size *)MemToShadow(addr, as);                \
+    auto *shadow_address = (__SYCL_GLOBAL__ u##size *)MemToShadow(addr, as);   \
     if (shadow_address && *shadow_address) {                                   \
       __asan_report_access_error(addr, as, size, is_write, addr, file, line,   \
                                  func, true);                                  \
@@ -595,7 +616,7 @@ static __SYCL_CONSTANT__ const char __mem_set_shadow_local[] =
     "[kernel] set_shadow_local(beg=%p, end=%p, val:%02X)\n";
 
 DEVICE_EXTERN_C_NOINLINE void
-__asan_set_shadow_local_memory(uptr ptr, size_t size,
+__asan_set_shadow_static_local(uptr ptr, size_t size,
                                size_t size_with_redzone) {
   // Since ptr is aligned to ASAN_SHADOW_GRANULARITY,
   // if size != aligned_size, then the buffer tail of ptr is not aligned
@@ -638,4 +659,47 @@ __asan_set_shadow_local_memory(uptr ptr, size_t size,
   }
 }
 
+static __SYCL_CONSTANT__ const char __mem_local_arg[] =
+    "[kernel] local_arg(index=%d, size=%d, size_rz=%d)\n";
+
+static __SYCL_CONSTANT__ const char __mem_set_shadow_dynamic_local_begin[] =
+    "[kernel] BEGIN __asan_set_shadow_dynamic_local\n";
+static __SYCL_CONSTANT__ const char __mem_set_shadow_dynamic_local_end[] =
+    "[kernel] END   __asan_set_shadow_dynamic_local\n";
+static __SYCL_CONSTANT__ const char __mem_report_arg_count_incorrect[] =
+    "[kernel] ERROR: The number of local args is incorrect, expect %d, actual "
+    "%d\n";
+
+DEVICE_EXTERN_C_NOINLINE void
+__asan_set_shadow_dynamic_local(uptr ptr, uint32_t num_args) {
+  if (__AsanDebug)
+    __spirv_ocl_printf(__mem_set_shadow_dynamic_local_begin);
+
+  auto *launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
+  if (num_args != launch_info->NumLocalArgs) {
+    __spirv_ocl_printf(__mem_report_arg_count_incorrect, num_args,
+                       launch_info->NumLocalArgs);
+    return;
+  }
+
+  uptr *args = (uptr *)ptr;
+  if (__AsanDebug)
+    __spirv_ocl_printf(__mem_launch_info, launch_info->LocalShadowOffset,
+                       launch_info->LocalShadowOffsetEnd,
+                       launch_info->NumLocalArgs, launch_info->LocalArgs);
+
+  for (uint32_t i = 0; i < num_args; ++i) {
+    auto *local_arg = &launch_info->LocalArgs[i];
+    if (__AsanDebug)
+      __spirv_ocl_printf(__mem_local_arg, i, local_arg->Size,
+                         local_arg->SizeWithRedZone);
+
+    __asan_set_shadow_static_local(args[i], local_arg->Size,
+                                   local_arg->SizeWithRedZone);
+  }
+
+  if (__AsanDebug)
+    __spirv_ocl_printf(__mem_set_shadow_dynamic_local_end);
+}
+
 #endif // __SPIR__ || __SPIRV__

llvm/lib/Transforms/IPO/DeadArgumentElimination.cpp

@@ -540,6 +540,13 @@ void DeadArgumentEliminationPass::surveyFunction(const Function &F) {
     return;
   }
 
+  // Don't touch sanitized functions. The "__asan_launch" argument needs to be
+  // present at all times, even if it's not used.
+  if (F.hasFnAttribute(Attribute::SanitizeAddress)) {
+    markLive(F);
+    return;
+  }
+
   unsigned RetCount = numRetVals(&F);
 
   // Assume all return values are dead