diff --git a/llvm/lib/Transforms/Instrumentation/GPUSan.cpp b/llvm/lib/Transforms/Instrumentation/GPUSan.cpp
index 1913394707e30..39a575f94b752 100644
--- a/llvm/lib/Transforms/Instrumentation/GPUSan.cpp
+++ b/llvm/lib/Transforms/Instrumentation/GPUSan.cpp
@@ -19,6 +19,7 @@
 #include "llvm/ADT/StringRef.h"
 #include "llvm/Analysis/LoopInfo.h"
 #include "llvm/Analysis/ScalarEvolution.h"
+#include "llvm/Analysis/ScalarEvolutionExpressions.h"
 #include "llvm/Analysis/ValueTracking.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/CallingConv.h"
@@ -45,6 +46,7 @@
 #include "llvm/Support/StringSaver.h"
 #include "llvm/Transforms/Utils/Cloning.h"
 #include "llvm/Transforms/Utils/ModuleUtils.h"
+#include "llvm/Transforms/Utils/ScalarEvolutionExpander.h"
 #include <cstdint>
 #include <optional>
 
@@ -281,6 +283,37 @@ class GPUSanImpl final {
                          {getPtrTy(PO), getPtrTy(PO), Int64Ty, Int32Ty, Int64Ty,
                           Int64Ty, Int64Ty, Int64Ty});
   }
+
+    FunctionCallee getCheckRangeWithBaseFn(PtrOrigin PO, Type* UpperBoundType, Type* LowerBoundType) {
+    return getOrCreateFn(CheckRangeWithBaseFn[PO],
+                         "ompx_check_range_with_base" + getSuffix(PO),
+                         Type::getVoidTy(Ctx),
+                         {
+                             UpperBoundType, /*SCEV max computed address*/
+                             LowerBoundType, /*SCEV min computed address*/
+                             getPtrTy(PO),    /*Start of allocation address*/
+                             Int64Ty, /*Size of allocation, i.e. Length*/
+                             Int32Ty, /*Tag*/
+                             Int64Ty, /*Size of the type that is loaded/stored*/
+                             Int64Ty, /*AccessId, Read/Write*/
+                             Int64Ty, /*SourceId, Allocation source ID*/
+                             Int64Ty  /*PC -- Program Counter*/
+                         });
+  }
+
+  FunctionCallee getCheckRangeFn(PtrOrigin PO, Type* UpperBoundType, Type* LowerBoundType) {
+    return getOrCreateFn(CheckRangeFn[PO], "ompx_check_range" + getSuffix(PO),
+                         Type::getVoidTy(Ctx),
+                         {
+                             UpperBoundType, /*SCEV max computed address*/
+                             LowerBoundType, /*SCEV min computed address*/
+                             Int64Ty, /*Size of the type that is loaded/stored*/
+                             Int64Ty, /*AccessId, Read/Write*/
+                             Int64Ty, /*SourceId, Allocation source ID*/
+                             Int64Ty  /*PC -- Program Counter*/
+                         });
+  }
+
   FunctionCallee getAllocationInfoFn(PtrOrigin PO) {
     assert(PO >= LOCAL && PO <= GLOBAL && "Origin does not need handling.");
     if (auto *F = M.getFunction("ompx_get_allocation_info" + getSuffix(PO)))
@@ -356,6 +389,8 @@ class GPUSanImpl final {
   FunctionCallee LifetimeStartFn;
   FunctionCallee FreeNLocalFn;
   FunctionCallee ThreadIDFn;
+  FunctionCallee CheckRangeWithBaseFn[3];
+  FunctionCallee CheckRangeFn[3];
 
   StringMap<Value *> GlobalStringMap;
   struct AllocationInfoTy {
@@ -1087,10 +1122,176 @@ void GPUSanImpl::instrumentAccess(LoopInfo &LI, Instruction &I, int PtrIdx,
   }
 
   if (Loop *L = LI.getLoopFor(I.getParent())) {
+
     auto &SE = FAM.getResult<ScalarEvolutionAnalysis>(*I.getFunction());
-    const auto &LD = SE.getLoopDisposition(SE.getSCEVAtScope(PtrOp, L), L);
+    SCEVExpander Expander = SCEVExpander(SE, DL, "SCEVExpander");
+    const SCEV *PtrExpr = SE.getSCEV(PtrOp);
+
+    const SCEV *ScStart;
+    const SCEV *ScEnd;
+    const SCEV *Step;
+
+    if (SE.isLoopInvariant(PtrExpr, L)) {
+
+      if (!Expander.isSafeToExpand(PtrExpr))
+        goto handleunhoistable;
+
+      // Assumption: Current loop has one unique predecessor
+      // We can insert at the end of the basic block if it
+      // is not a branch instruction.
+      auto *Entry = L->getLoopPreheader();
+
+      if (!Entry)
+        goto handleunhoistable;
+
+      Instruction *PtrOpInst = dyn_cast<Instruction>(PtrOp);
+
+      if (!PtrOpInst)
+        goto handleunhoistable;
+
+      // Get handle to last instruction.
+      auto LoopEnd = --(Entry->end());
+
+      static int32_t ReadAccessId = -1;
+      static int32_t WriteAccessId = 1;
+      const int32_t &AccessId = IsRead ? ReadAccessId-- : WriteAccessId++;
+
+      auto TySize = DL.getTypeStoreSize(&AccessTy);
+      assert(!TySize.isScalable());
+      Value *Size = ConstantInt::get(Int64Ty, TySize.getFixedValue());
+
+      LoopEnd = --(Entry->end());
+      CallInst *CB;
+      Value *PCVal = getPC(IRB);
+      Instruction *PCInst = dyn_cast<Instruction>(PCVal);
+      if (!PCInst)
+        return;
+
+      Value *AccessIDVal = ConstantInt::get(Int64Ty, AccessId);
+      PCInst->removeFromParent();
+      PCInst->insertBefore(LoopEnd);
+
+      FunctionCallee Callee;
+
+      if (Start) {
+        Callee = getCheckWithBaseFn(PO);
+        CB = createCall(IRB, Callee,
+                        {PtrOp, Start, Length, Tag, Size,
+                         ConstantInt::get(Int64Ty, AccessId), getSourceIndex(I),
+                         PCVal});
+      } else {
+        Callee = getCheckFn(PO);
+        CB = createCall(IRB, Callee,
+                        {PtrOp, Size, ConstantInt::get(Int64Ty, AccessId),
+                         getSourceIndex(I), PCVal});
+      }
+      CB->removeFromParent();
+      CB->insertAfter(PCInst);
+
+      I.setOperand(PtrIdx, IRB.CreatePointerBitCastOrAddrSpaceCast(
+                               CB, PtrOp->getType()));
+
+      return;
+
+    } else {
+      const SCEVAddRecExpr *AddRecExpr = dyn_cast<SCEVAddRecExpr>(PtrExpr);
+      if (AddRecExpr) {
+
+        auto *Entry = L->getLoopPreheader();
+
+        if (!Entry)
+          goto handleunhoistable;
+
+        const SCEV *Ex = SE.getSymbolicMaxBackedgeTakenCount(L);
+
+        ScStart = AddRecExpr->getStart();
+        ScEnd = AddRecExpr->evaluateAtIteration(Ex, SE);
+        Step = AddRecExpr->getStepRecurrence(SE);
+
+        if (const auto *CStep = dyn_cast<SCEVConstant>(Step)) {
+          if (CStep->getValue()->isNegative())
+            std::swap(ScStart, ScEnd);
+        } else {
+          ScStart = SE.getUMinExpr(ScStart, ScEnd);
+          ScEnd = SE.getUMaxExpr(AddRecExpr->getStart(), ScEnd);
+        }
+
+        if (!Expander.isSafeToExpand(ScStart))
+          goto handleunhoistable;
+
+        if (!Expander.isSafeToExpand(ScEnd))
+          goto handleunhoistable;
+
+        // Get handle to last instruction.
+        auto LoopEnd = --(Entry->end());
+        Instruction *LoopEndInst = &*LoopEnd;
+
+        Type *Int64Ty = Type::getInt64Ty(Ctx);
+        Value *LowerBoundCode =
+            Expander.expandCodeFor(ScStart, nullptr, LoopEnd);
+
+        LoopEnd = --(Entry->end());
+
+        Value *UpperBoundCode = Expander.expandCodeFor(ScEnd, nullptr, LoopEnd);
+        static int32_t ReadAccessId = -1;
+        static int32_t WriteAccessId = 1;
+        const int32_t &AccessId = IsRead ? ReadAccessId-- : WriteAccessId++;
+
+        auto TySize = DL.getTypeStoreSize(&AccessTy);
+        assert(!TySize.isScalable());
+        Value *Size = ConstantInt::get(Int64Ty, TySize.getFixedValue());
+
+        LoopEnd = --(Entry->end());
+
+        CallInst *CB;
+        Value *PCVal = getPC(IRB);
+        Instruction *PCInst = dyn_cast<Instruction>(PCVal);
+
+        if (!PCInst)
+          return;
+
+        Value *AccessIDVal = ConstantInt::get(Int64Ty, AccessId);
+        PCInst->removeFromParent();
+        PCInst->insertBefore(LoopEnd);
+
+        FunctionCallee Callee;
+
+        if (Start) {
+          Callee = getCheckRangeWithBaseFn(PO, UpperBoundCode->getType(),
+                                           LowerBoundCode->getType());
+          CB = createCall(IRB, Callee,
+                          {UpperBoundCode, LowerBoundCode, Start, Length, Tag,
+                           Size, AccessIDVal, getSourceIndex(I), PCVal});
+        } else {
+          Callee = getCheckRangeFn(PO, UpperBoundCode->getType(),
+                                   LowerBoundCode->getType());
+          CB = createCall(IRB, Callee,
+                          {UpperBoundCode, LowerBoundCode, Size, AccessIDVal,
+                           getSourceIndex(I), PCVal});
+        }
+        CB->removeFromParent();
+        CB->insertAfter(PCInst);
+
+        // Convert fake pointer to real pointer.
+        Value *PlainPtrOp =
+            IRB.CreatePointerBitCastOrAddrSpaceCast(PtrOp, getPtrTy(PO));
+        auto *CBUnpack =
+            createCall(IRB, getUnpackFn(PO), {PlainPtrOp, getPC(IRB)},
+                       PtrOp->getName() + ".unpack");
+
+        I.setOperand(PtrIdx, IRB.CreatePointerBitCastOrAddrSpaceCast(
+                                 CBUnpack, PtrOp->getType()));
+
+        return;
+
+      } else {
+        goto handleunhoistable;
+      }
+    }
   }
 
+handleunhoistable:
+
   static int32_t ReadAccessId = -1;
   static int32_t WriteAccessId = 1;
   const int32_t &AccessId = IsRead ? ReadAccessId-- : WriteAccessId++;
@@ -1101,6 +1302,7 @@ void GPUSanImpl::instrumentAccess(LoopInfo &LI, Instruction &I, int PtrIdx,
 
   Value *PlainPtrOp =
       IRB.CreatePointerBitCastOrAddrSpaceCast(PtrOp, getPtrTy(PO));
+
   CallInst *CB;
   if (Start) {
     CB = createCall(IRB, getCheckWithBaseFn(PO),
@@ -1114,6 +1316,7 @@ void GPUSanImpl::instrumentAccess(LoopInfo &LI, Instruction &I, int PtrIdx,
                      getSourceIndex(I), getPC(IRB)},
                     I.getName() + ".san");
   }
+
   I.setOperand(PtrIdx,
                IRB.CreatePointerBitCastOrAddrSpaceCast(CB, PtrOp->getType()));
 }
diff --git a/offload/DeviceRTL/src/Sanitizer.cpp b/offload/DeviceRTL/src/Sanitizer.cpp
index 08e401ae592de..5fbb370e89e06 100644
--- a/offload/DeviceRTL/src/Sanitizer.cpp
+++ b/offload/DeviceRTL/src/Sanitizer.cpp
@@ -181,6 +181,27 @@ template <AllocationKind AK> struct AllocationTracker {
     return utils::advancePtr(Start, Offset);
   }
 
+  [[clang::disable_sanitizer_instrumentation]] static void
+  checkWithBaseVoid(_AS_PTR(void, AK) P, _AS_PTR(void, AK) Start,
+                    int64_t Length, uint32_t Tag, int64_t Size,
+                    int64_t AccessId, int64_t SourceId, uint64_t PC) {
+    AllocationPtrTy<AK> AP = AllocationPtrTy<AK>::get(P);
+    if constexpr (AK == AllocationKind::LOCAL)
+      if (Length == 0)
+        Length = getAllocation<AK>(AP, AccessId, PC).Length;
+    if constexpr (AK == AllocationKind::GLOBAL)
+      if (AP.Magic != SanitizerConfig<AllocationKind::GLOBAL>::MAGIC)
+        __sanitizer_trap_info_ptr->garbagePointer<AK>(AP, (void *)P, SourceId,
+                                                      PC);
+    int64_t Offset = AP.Offset;
+    if (OMP_UNLIKELY(
+            Offset > Length - Size ||
+            (SanitizerConfig<AK>::useTags() && Tag != AP.AllocationTag))) {
+      __sanitizer_trap_info_ptr->accessError<AK>(AP, Size, AccessId, SourceId,
+                                                 PC);
+    }
+  }
+
   [[clang::disable_sanitizer_instrumentation]] static _AS_PTR(void, AK)
       check(_AS_PTR(void, AK) P, int64_t Size, int64_t AccessId,
             int64_t SourceId, uint64_t PC) {
@@ -190,6 +211,67 @@ template <AllocationKind AK> struct AllocationTracker {
                          AccessId, SourceId, PC);
   }
 
+  [[clang::disable_sanitizer_instrumentation]] static void
+  checkRangeWithBase(_AS_PTR(void, AK) SCEVMax, _AS_PTR(void, AK) SCEVMin,
+                     _AS_PTR(void, AK) StartAddress, int64_t AllocationLength,
+                     uint32_t Tag, int64_t AccessTypeSize, int64_t AccessId,
+                     int64_t SourceId, uint64_t PC) {
+    // printf("Hello World!\n");
+    AllocationPtrTy<AK> APSCEVMax = AllocationPtrTy<AK>::get(SCEVMax);
+    AllocationPtrTy<AK> APSCEVMin = AllocationPtrTy<AK>::get(SCEVMin);
+    if constexpr (AK == AllocationKind::LOCAL)
+      if (AllocationLength == 0)
+        AllocationLength = getAllocation<AK>(APSCEVMax, AccessId, PC).Length;
+
+    if constexpr (AK == AllocationKind::GLOBAL) {
+      if (APSCEVMax.Magic != SanitizerConfig<AllocationKind::GLOBAL>::MAGIC)
+        __sanitizer_trap_info_ptr->garbagePointer<AK>(
+            APSCEVMax, (void *)SCEVMax, SourceId, PC);
+
+      if (APSCEVMin.Magic != SanitizerConfig<AllocationKind::GLOBAL>::MAGIC)
+        __sanitizer_trap_info_ptr->garbagePointer<AK>(
+            APSCEVMin, (void *)SCEVMin, SourceId, PC);
+    }
+
+    // check upper bound
+    int64_t MaxOffset = APSCEVMax.Offset;
+    if (OMP_UNLIKELY(MaxOffset > AllocationLength - AccessTypeSize ||
+                     (SanitizerConfig<AK>::useTags() &&
+                      Tag != APSCEVMax.AllocationTag))) {
+      __sanitizer_trap_info_ptr->accessError<AK>(APSCEVMax, AccessTypeSize,
+                                                 AccessId, SourceId, PC);
+    }
+
+    // check lower bound
+    auto &AllocationOfMinOffset = getAllocation<AK>(APSCEVMin, AccessId, PC);
+    if (OMP_UNLIKELY(AllocationOfMinOffset.Start != StartAddress ||
+                     (SanitizerConfig<AK>::useTags() &&
+                      Tag != APSCEVMin.AllocationTag))) {
+      __sanitizer_trap_info_ptr->accessError<AK>(APSCEVMin, AccessTypeSize,
+                                                 AccessId, SourceId, PC);
+    }
+  }
+
+  [[clang::disable_sanitizer_instrumentation]] static void
+  checkRange(_AS_PTR(void, AK) SCEVMax, _AS_PTR(void, AK) SCEVMin,
+             int64_t AccessTypeSize, int64_t AccessId, int64_t SourceId,
+             uint64_t PC) {
+    AllocationPtrTy<AK> AP = AllocationPtrTy<AK>::get(SCEVMax);
+    auto &Alloc = getAllocation<AK>(AP, AccessId, PC);
+    return checkRangeWithBase(SCEVMax, SCEVMin, Alloc.Start, Alloc.Length,
+                              Alloc.Tag, AccessTypeSize, AccessId, SourceId,
+                              PC);
+  }
+
+  [[clang::disable_sanitizer_instrumentation]] static void
+  checkVoid(_AS_PTR(void, AK) P, int64_t Size, int64_t AccessId,
+            int64_t SourceId, uint64_t PC) {
+    AllocationPtrTy<AK> AP = AllocationPtrTy<AK>::get(P);
+    auto &Alloc = getAllocation<AK>(AP, AccessId, PC);
+    return checkWithBaseVoid(P, Alloc.Start, Alloc.Length, Alloc.Tag, Size,
+                             AccessId, SourceId, PC);
+  }
+
   [[clang::disable_sanitizer_instrumentation]] static _AS_PTR(void, AK)
       unpack(_AS_PTR(void, AK) P, int64_t SourceId, uint64_t PC) {
     AllocationPtrTy<AK> AP = AllocationPtrTy<AK>::get(P);
@@ -255,6 +337,21 @@ getFakePtrType(void *P, int64_t SourceId, uint64_t PC) {
       AllocationPtrTy<AllocationKind::LOCAL>::get(P), P, SourceId, PC);
 }
 
+static void checkForMagic2(bool IsGlobal, void *P, int64_t SourceId,
+                           int64_t AccessId, uint64_t PC) {
+  if (IsGlobal) {
+    auto AP = AllocationPtrTy<AllocationKind::GLOBAL>::get(P);
+    if (AP.Magic != SanitizerConfig<AllocationKind::GLOBAL>::MAGIC)
+      __sanitizer_trap_info_ptr->garbagePointer2<AllocationKind::GLOBAL>(
+          AP, P, SourceId, AccessId, PC);
+  } else {
+    auto AP = AllocationPtrTy<AllocationKind::LOCAL>::get(P);
+    if (AP.Magic != SanitizerConfig<AllocationKind::LOCAL>::MAGIC)
+      __sanitizer_trap_info_ptr->garbagePointer2<AllocationKind::LOCAL>(
+          AP, P, SourceId, AccessId, PC);
+  }
+}
+
 extern "C" {
 
 #define REAL_PTR_IS_SHARED(PTR) (isSharedMemPtr(PTR))
@@ -452,6 +549,122 @@ ompx_check(void *P, uint64_t Size, uint64_t AccessId, int64_t SourceId,
       P, Start, Length, Tag, Size, AccessId, SourceId, PC);
 }
 
+// Void functions for sanitizing a pointer from base offset and without it
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_void_local(_AS_PTR(void, AllocationKind::LOCAL) P, uint64_t Size,
+                      uint64_t AccessId, int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::LOCAL>::checkVoid(P, Size, AccessId,
+                                                             SourceId, PC);
+}
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_void_global(_AS_PTR(void, AllocationKind::GLOBAL) P, uint64_t Size,
+                       uint64_t AccessId, int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::GLOBAL>::checkVoid(P, Size, AccessId,
+                                                              SourceId, PC);
+}
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_void(void *P, uint64_t Size, uint64_t AccessId, int64_t SourceId,
+                uint64_t PC) {
+  bool IsGlobal = IS_GLOBAL(P);
+  checkForMagic2(IsGlobal, P, SourceId, AccessId, PC);
+  if (IsGlobal)
+    return ompx_check_void_global((_AS_PTR(void, AllocationKind::GLOBAL))P,
+                                  Size, AccessId, SourceId, PC);
+  return ompx_check_void_local((_AS_PTR(void, AllocationKind::LOCAL))P, Size,
+                               AccessId, SourceId, PC);
+}
+
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_with_base_void_local(_AS_PTR(void, AllocationKind::LOCAL) P,
+                                _AS_PTR(void, AllocationKind::LOCAL) Start,
+                                uint64_t Length, uint32_t Tag, uint64_t Size,
+                                uint64_t AccessId, int64_t SourceId,
+                                uint64_t PC) {
+  return AllocationTracker<AllocationKind::LOCAL>::checkWithBaseVoid(
+      P, Start, Length, Tag, Size, AccessId, SourceId, PC);
+}
+
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_with_base_void_global(_AS_PTR(void, AllocationKind::GLOBAL) P,
+                                 _AS_PTR(void, AllocationKind::GLOBAL) Start,
+                                 uint64_t Length, uint32_t Tag, uint64_t Size,
+                                 uint64_t AccessId, int64_t SourceId,
+                                 uint64_t PC) {
+  return AllocationTracker<AllocationKind::GLOBAL>::checkWithBaseVoid(
+      P, Start, Length, Tag, Size, AccessId, SourceId, PC);
+}
+// End of void functions.
+
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_range_with_base_global(_AS_PTR(void, AllocationKind::GLOBAL) SCEVMax,
+                                  _AS_PTR(void, AllocationKind::GLOBAL) SCEVMin,
+                                  _AS_PTR(void, AllocationKind::GLOBAL)
+                                      StartAddress,
+                                  int64_t AllocationLength, uint32_t Tag,
+                                  int64_t AccessTypeSize, int64_t AccessId,
+                                  int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::GLOBAL>::checkRangeWithBase(
+      SCEVMax, SCEVMin, StartAddress, AllocationLength, Tag, AccessTypeSize,
+      AccessId, SourceId, PC);
+}
+
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_range_with_base_local(_AS_PTR(void, AllocationKind::LOCAL) SCEVMax,
+                                 _AS_PTR(void, AllocationKind::LOCAL) SCEVMin,
+                                 _AS_PTR(void, AllocationKind::LOCAL)
+                                     StartAddress,
+                                 int64_t AllocationLength, uint32_t Tag,
+                                 int64_t AccessTypeSize, int64_t AccessId,
+                                 int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::LOCAL>::checkRangeWithBase(
+      SCEVMax, SCEVMin, StartAddress, AllocationLength, Tag, AccessTypeSize,
+      AccessId, SourceId, PC);
+}
+
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_range_local(_AS_PTR(void, AllocationKind::LOCAL) SCEVMax,
+                       _AS_PTR(void, AllocationKind::LOCAL) SCEVMin,
+                       int64_t AccessTypeSize, int64_t AccessId,
+                       int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::LOCAL>::checkRange(
+      SCEVMax, SCEVMin, AccessTypeSize, AccessId, SourceId, PC);
+}
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_range_global(_AS_PTR(void, AllocationKind::GLOBAL) SCEVMax,
+                        _AS_PTR(void, AllocationKind::GLOBAL) SCEVMin,
+                        int64_t AccessTypeSize, int64_t AccessId,
+                        int64_t SourceId, uint64_t PC) {
+  return AllocationTracker<AllocationKind::GLOBAL>::checkRange(
+      SCEVMax, SCEVMin, AccessTypeSize, AccessId, SourceId, PC);
+}
+[[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
+  gnu::used, gnu::retain]] void
+ompx_check_range(void *SCEVMax, void *SCEVMin, int64_t AccessTypeSize,
+                 int64_t AccessId, int64_t SourceId, uint64_t PC) {
+  bool IsGlobalMax = IS_GLOBAL(SCEVMax);
+  bool IsGlobalMin = IS_GLOBAL(SCEVMin);
+  checkForMagic2(IsGlobalMax, SCEVMax, SourceId, AccessId, PC);
+  checkForMagic2(IsGlobalMin, SCEVMin, SourceId, AccessId, PC);
+  if (IsGlobalMax && IsGlobalMin)
+    return ompx_check_range_global(
+        (_AS_PTR(void, AllocationKind::GLOBAL))SCEVMax,
+        (_AS_PTR(void, AllocationKind::GLOBAL))SCEVMin, AccessTypeSize,
+        AccessId, SourceId, PC);
+
+  return ompx_check_range_local((_AS_PTR(void, AllocationKind::LOCAL))SCEVMax,
+                                (_AS_PTR(void, AllocationKind::LOCAL))SCEVMin,
+                                AccessTypeSize, AccessId, SourceId, PC);
+}
+
 [[clang::disable_sanitizer_instrumentation, gnu::flatten, gnu::always_inline,
   gnu::used, gnu::retain]] _AS_PTR(void, AllocationKind::LOCAL)
     ompx_unpack_local(_AS_PTR(void, AllocationKind::LOCAL) P,
diff --git a/offload/include/Shared/Sanitizer.h b/offload/include/Shared/Sanitizer.h
index b1d6fa3552a0d..28d2f3b888b91 100644
--- a/offload/include/Shared/Sanitizer.h
+++ b/offload/include/Shared/Sanitizer.h
@@ -332,6 +332,23 @@ struct SanitizerTrapInfoTy {
     __builtin_trap();
   }
 
+  template <enum AllocationKind AK>
+  [[clang::disable_sanitizer_instrumentation, noreturn, NOINLINE,
+    gnu::cold]] void
+  garbagePointer2(const AllocationPtrTy<AK> AP, void *P, int64_t SourceId,
+                  int64_t Id, uint64_t PC) {
+    ErrorCode = GarbagePointer;
+    AllocationStart = P;
+    AllocationKind = (decltype(AllocationKind))AK;
+    PtrOffset = AP.Offset;
+    PtrSlot = AP.AllocationId;
+    PtrTag = AP.AllocationTag;
+    PtrKind = AP.Kind;
+    AccessId = Id;
+    setCoordinates(SourceId, PC);
+    __builtin_trap();
+  }
+
   template <enum AllocationKind AK>
   [[clang::disable_sanitizer_instrumentation, noreturn, INLINE, gnu::cold]] void
   memoryLeak(const AllocationTy<AK> A, uint64_t Slot) {
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/hoist_check1.c b/offload/test/sanitizer/optimize_sanitizer_tests/hoist_check1.c
new file mode 100644
index 0000000000000..3fc001c03c55c
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/hoist_check1.c
@@ -0,0 +1,38 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+IntTy *foo(IntTy Size) {
+
+  IntTy *a;
+  a = (IntTy *)malloc(sizeof(IntTy) * Size);
+
+  //Ideally the check for this pointer should be hoisted out of the loop.  
+  IntTy *Hoistable = (IntTy *)malloc(sizeof(IntTy));
+  
+#pragma omp target teams map(from : a[0:Size])
+  {
+    for (IntTy I = 0; I < Size; I++) {
+      Hoistable[0] = 1;
+      a[I] = I;
+    }
+  }
+
+  return a;
+}
+
+void printArray(IntTy *a, IntTy Size) {
+
+  for (IntTy I = 0; I < Size; I++) {
+    printf("a: %d ", a[I]);
+  }
+}
+
+int main() {
+
+  int N = 10000000;
+  int *a = foo(N);
+  //printArray(a, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_lower.c b/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_lower.c
new file mode 100644
index 0000000000000..0dd0daeab7f14
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_lower.c
@@ -0,0 +1,35 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+IntTy *foo(int Size) {
+
+  IntTy *a;
+
+  a = (IntTy *)malloc(sizeof(IntTy) * Size);
+
+#pragma omp target teams map(from : a [0:Size])
+  {
+    for (IntTy I = -1; I < Size; I++) {
+      a[I] = I;
+    }
+  }
+
+  return a;
+}
+
+void printArray(int *a, int Size) {
+
+  for (IntTy I = 0; I < Size; I++) {
+    printf("a: %d ", a[I]);
+  }
+}
+
+int main() {
+
+  int N = 1000;
+  int *a = foo(N);
+  printArray(a, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_upper.c b/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_upper.c
new file mode 100644
index 0000000000000..3665b8d7eca22
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/illegal_simple_access_upper.c
@@ -0,0 +1,35 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+IntTy *foo(int Size) {
+
+  IntTy *a;
+
+  a = (IntTy *)malloc(sizeof(IntTy) * Size);
+
+#pragma omp target teams map(from : a [0:Size])
+  {
+    for (IntTy I = 0; I < Size + 1; I++) {
+      a[I] = I;
+    }
+  }
+
+  return a;
+}
+
+void printArray(int *a, int Size) {
+
+  for (IntTy I = 0; I < Size; I++) {
+    printf("a: %d ", a[I]);
+  }
+}
+
+int main() {
+
+  int N = 1000;
+  int *a = foo(N);
+  printArray(a, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable.c b/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable.c
new file mode 100644
index 0000000000000..bcebc2a442c27
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable.c
@@ -0,0 +1,55 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+typedef struct {
+  IntTy *a;
+  IntTy *b;
+  IntTy *c;
+} Product;
+
+Product *foo(IntTy N) {
+
+  IntTy *a;
+  IntTy *b;
+  IntTy *c;
+
+  a = (IntTy *)malloc(sizeof(IntTy) * N);
+  b = (IntTy *)malloc(sizeof(IntTy) * N);
+  c = (IntTy *)malloc(sizeof(IntTy) * N);
+
+#pragma omp target teams map(from : a [0:N], b [0:N], c [0:N])
+  {
+    for (IntTy I = 0; I < N; I++) {
+      a[I] = I;
+      b[I] = I;
+      c[I] = a[I] + b[I];
+    }
+  }
+
+  Product *P = (Product *)malloc(sizeof(Product));
+  P->a = a;
+  P->b = b;
+  P->c = c;
+  return P;
+}
+
+void printProduct(Product *P, IntTy N) {
+
+  IntTy *a = P->a;
+  IntTy *b = P->b;
+  IntTy *c = P->c;
+
+  for (IntTy i = 1; i < N; i++) {
+    printf("a: %d, b:%d, c:%d\n", a[i], b[i], c[i]);
+  }
+}
+
+int main() {
+
+  IntTy N = 1000;
+  Product *P = foo(N);
+  printProduct(P, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable_local_memory.c b/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable_local_memory.c
new file mode 100644
index 0000000000000..cfab7207145d6
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/legal_access_mergable_local_memory.c
@@ -0,0 +1,51 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+typedef struct {
+  IntTy *a;
+  IntTy *b;
+  IntTy *c;
+} Product;
+
+Product *foo(IntTy N) {
+
+  int a[N];
+  int b[N];
+  int c[N];
+
+#pragma omp target teams map(from : a [0:N], b [0:N], c [0:N])
+  {
+    for (IntTy I = 0; I < N; I++) {
+      a[I] = I;
+      b[I] = I;
+      c[I] = a[I] + b[I];
+    }
+  }
+
+  Product *P = (Product *)malloc(sizeof(Product));
+  P->a = a;
+  P->b = b;
+  P->c = c;
+  return P;
+}
+
+void printProduct(Product *P, IntTy N) {
+
+  IntTy *a = P->a;
+  IntTy *b = P->b;
+  IntTy *c = P->c;
+
+  for (IntTy i = 1; i < N; i++) {
+    printf("a: %d, b:%d, c:%d\n", a[i], b[i], c[i]);
+  }
+}
+
+int main() {
+
+  IntTy N = 1000;
+  Product *P = foo(N);
+  printProduct(P, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_monotonic.c b/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_monotonic.c
new file mode 100644
index 0000000000000..16105cc4f7493
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_monotonic.c
@@ -0,0 +1,34 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+IntTy *foo(IntTy Size) {
+
+  IntTy *a;
+  a = (IntTy *)malloc(sizeof(IntTy) * Size);
+
+#pragma omp target teams map(from : a[0:Size])
+  {
+    for (IntTy I = 0; I < Size; I++) {
+      a[I] = I;
+    }
+  }
+
+  return a;
+}
+
+void printArray(IntTy *a, IntTy Size) {
+
+  for (IntTy I = 0; I < Size; I++) {
+    printf("a: %d ", a[I]);
+  }
+}
+
+int main() {
+
+  int N = 1000;
+  int *a = foo(N);
+  printArray(a, N);
+}
diff --git a/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_non_monotonic.c b/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_non_monotonic.c
new file mode 100644
index 0000000000000..dfe5986eea180
--- /dev/null
+++ b/offload/test/sanitizer/optimize_sanitizer_tests/legal_simple_access_non_monotonic.c
@@ -0,0 +1,52 @@
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef int IntTy;
+
+IntTy *foo() {
+
+  IntTy *a;
+  IntTy *b;
+  IntTy *c;
+
+  int N = 1000;
+
+  a = (IntTy *)malloc(sizeof(IntTy) * N * N);
+  b = (IntTy *)malloc(sizeof(IntTy) * N * N);
+  c = (IntTy *)malloc(sizeof(IntTy) * N * N);
+
+  // I ranges from -N, -N + 1, ..., 0, 1, 2, ... N
+  // Square ranges from N^2 - 1, (N-1)^2 ..., 0, 1, 4, 9, 16, ... N^2 - 1
+#pragma omp target teams map(from : a [0:N * N])
+  {
+    int Square;
+    for (IntTy I = -N; I <= N; I++) {
+      Square = I * I;
+      // legal boundary access.
+      if (I == N || I == -N) {
+        Square = Square - 1;
+      }
+      a[Square] = I;
+    }
+  }
+
+  return a;
+}
+
+void printArray(int *a, int N) {
+
+  for (IntTy I = -N; I <= N; I++) {
+    int Index = I * I;
+    if (I == N || I == -N) {
+      Index = Index - 1;
+    }
+    printf("a: %d ", a[Index]);
+  }
+}
+
+int main() {
+  int N = 1000;
+  int *a = foo();
+  printArray(a, N);
+}