[DevSanitizer] Fix several misc issues in device memory sanitizer (#18026)

* Disable inserting unreachable instruction since we have another
  mechanism to exit current kernel execution.
* Cleanup clean shadow after it's allocated.
* Remove redzone maximum limitation.

Author: zhaomaosu

llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp

@@ -2193,7 +2193,8 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
       Value *Cmp = convertToBool(ConvertedShadow, IRB, "_mscmp");
       Instruction *CheckTerm = SplitBlockAndInsertIfThen(
           Cmp, &*IRB.GetInsertPoint(),
-          /* Unreachable */ !MS.Recover, MS.ColdCallWeights);
+          /* Unreachable */ SpirOrSpirv ? false : !MS.Recover,
+          MS.ColdCallWeights);
 
       IRB.SetInsertPoint(CheckTerm);
       insertWarningFn(IRB, Origin);

sycl/test-e2e/AddressSanitizer/common/options-redzone.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -DUNSAFE -O0 -g -o %t1.out
-// RUN: env UR_LAYER_ASAN_OPTIONS=redzone:64 %{run} not %t1.out 2>&1 | FileCheck %s
+// RUN: env UR_LAYER_ASAN_OPTIONS=redzone:4000 %{run} not %t1.out 2>&1 | FileCheck %s
 // RUN: %{build} %device_asan_flags -DSAFE -O0 -g -o %t2.out
 
 // clang-format off
@@ -21,10 +21,11 @@ int main() {
      h.single_task<class Test>([=]() { ++array[0]; });
 #endif
    }).wait();
+  // CHECK: <SANITIZER>[WARNING]: Increasing the redzone size may cause excessive memory overhead
   // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
   // CHECK: {{READ of size 1 at kernel <.*Test> LID\(0, 0, 0\) GID\(0, 0, 0\)}}
-  // CHECK: {{  #0 .* .*options-redzone.cpp:}}[[@LINE-7]]
-  // CHECK-MIN: The valid range of "redzone" is [16, 2048]. Setting to the minimum value 16.
+  // CHECK: {{  #0 .* .*options-redzone.cpp:}}[[@LINE-8]]
+  // CHECK-MIN: The valid range of "redzone" is [16, 18446744073709551615]. Setting to the minimum value 16.
 
   sycl::free(array, q);
   return 0;

unified-runtime/source/loader/layers/sanitizer/msan/msan_interceptor.cpp

@@ -479,6 +479,9 @@ ur_result_t MsanInterceptor::prepareLaunch(
   UR_CALL(getContext()->urDdiTable.USM.pfnDeviceAlloc(
       ContextInfo->Handle, DeviceInfo->Handle, nullptr, nullptr,
       ContextInfo->MaxAllocatedSize, (void **)&LaunchInfo.Data->CleanShadow));
+  UR_CALL(EnqueueUSMBlockingSet(Queue, (void *)LaunchInfo.Data->CleanShadow, 0,
+                                ContextInfo->MaxAllocatedSize, 0, nullptr,
+                                nullptr));
 
   if (LaunchInfo.LocalWorkSize.empty()) {
     LaunchInfo.LocalWorkSize.resize(LaunchInfo.WorkDim);

unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_common.hpp

@@ -34,6 +34,21 @@ inline constexpr bool IsPowerOfTwo(uptr x) {
   return (x & (x - 1)) == 0 && x != 0;
 }
 
+inline constexpr uptr RoundUpToPowerOfTwo(uptr x) {
+  if (x == 0)
+    return 1;
+  x--;
+  x |= x >> 1;
+  x |= x >> 2;
+  x |= x >> 4;
+  x |= x >> 8;
+  x |= x >> 16;
+  x |= x >> 32;
+  x++;
+  assert(IsPowerOfTwo(x));
+  return x;
+}
+
 inline constexpr uptr RoundUpTo(uptr Size, uptr boundary) {
   assert(IsPowerOfTwo(boundary));
   return (Size + boundary - 1) & ~(boundary - 1);
@@ -48,16 +63,12 @@ inline constexpr bool IsAligned(uptr a, uptr alignment) {
   return (a & (alignment - 1)) == 0;
 }
 
-// Valid redzone sizes are 16, 32, 64, ... 2048, so we encode them in 3 bits.
+// Valid redzone sizes are 16, 32, 64 ..., so we encode them in 3 bits.
 // We use adaptive redzones: for larger allocation larger redzones are used.
-inline constexpr uptr RZLog2Size(uptr rz_log) {
-  assert(rz_log < 8);
-  return 16 << rz_log;
-}
+inline constexpr uptr RZLog2Size(uptr rz_log) { return 16 << rz_log; }
 
 inline constexpr uptr RZSize2Log(uptr rz_size) {
   assert(rz_size >= 16);
-  assert(rz_size <= 2048);
   assert(IsPowerOfTwo(rz_size));
   uptr res = log2(rz_size) - 4;
   assert(rz_size == RZLog2Size(res));

unified-runtime/source/loader/layers/sanitizer/sanitizer_common/sanitizer_options.cpp

@@ -12,6 +12,7 @@
  */
 
 #include "sanitizer_options.hpp"
+#include "sanitizer_common.hpp"
 #include "sanitizer_options_impl.hpp"
 
 #include <cstring>
@@ -47,7 +48,13 @@ void SanitizerOptions::Init(const std::string &EnvName,
   Parser.ParseBool("halt_on_error", HaltOnError);
 
   Parser.ParseUint64("quarantine_size_mb", MaxQuarantineSizeMB, 0, UINT32_MAX);
-  Parser.ParseUint64("redzone", MinRZSize, 16, 2048);
+  Parser.ParseUint64("redzone", MinRZSize, 16);
+  MinRZSize =
+      IsPowerOfTwo(MinRZSize) ? MinRZSize : RoundUpToPowerOfTwo(MinRZSize);
+  if (MinRZSize > 16) {
+    Logger.warning(
+        "Increasing the redzone size may cause excessive memory overhead");
+  }
 }
 
 } // namespace ur_sanitizer_layer