Rebaseusers/pcc/spr/codegen-introduce-machinefunctiongetpreferredalignment

Created using spr 1.3.6-beta.1

45 files changed, 1380 insertions, 87 deletions

diff --git a/offload/ci/.gitignore b/offload/ci/.gitignore
new file mode 100644
index 0000000..f2198f0
--- /dev/null
+++ b/offload/ci/.gitignore
@@ -0,0 +1 @@
+/*.workdir
diff --git a/offload/ci/openmp-offload-amdgpu-clang-flang.py b/offload/ci/openmp-offload-amdgpu-clang-flang.py
new file mode 100755
index 0000000..f55ae02
--- /dev/null
+++ b/offload/ci/openmp-offload-amdgpu-clang-flang.py
@@ -0,0 +1,71 @@
+#! /usr/bin/env python3
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+import os
+import sys
+
+# Adapt to location in source tree
+llvmsrcroot = os.path.normpath(f"{__file__}/../../..")
+
+sys.path.insert(0, os.path.join(llvmsrcroot, ".ci/buildbot"))
+import worker
+
+llvmbuilddir = "llvm.build"
+llvminstalldir = "llvm.inst"
+
+with worker.run(
+    __file__,
+    llvmsrcroot,
+    clobberpaths=[llvmbuilddir, llvminstalldir],
+    workerjobs=64,
+) as w:
+    with w.step("configure-openmp", halt_on_fail=True):
+        w.run_command(
+            [
+                "cmake",
+                f"-S{w.in_llvmsrc('llvm')}",
+                f"-B{llvmbuilddir}",
+                "-GNinja",
+                "-DCMAKE_BUILD_TYPE=Release",
+                "-DLLVM_ENABLE_ASSERTIONS=ON",
+                f"-DLLVM_LIT_ARGS=-vv --show-unsupported --show-xfail -j {w.jobs} --time-tests --timeout 100",
+                f"-DCMAKE_INSTALL_PREFIX={w.in_workdir(llvminstalldir)}",
+                "-DCLANG_DEFAULT_LINKER=lld",
+                "-DLLVM_TARGETS_TO_BUILD=X86;AMDGPU",
+                "-DLLVM_ENABLE_ASSERTIONS=ON",
+                "-DCMAKE_C_COMPILER_LAUNCHER=ccache",
+                "-DCMAKE_CXX_COMPILER_LAUNCHER=ccache",
+                "-DFLANG_RUNTIME_F128_MATH_LIB=libquadmath",
+                "-DLLVM_ENABLE_PER_TARGET_RUNTIME_DIR=ON",
+                "-DCMAKE_CXX_STANDARD=17",
+                "-DBUILD_SHARED_LIBS=ON",
+                "-DLIBOMPTARGET_PLUGINS_TO_BUILD=amdgpu;host",
+                "-DRUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES=compiler-rt;openmp",
+                "-DLLVM_RUNTIME_TARGETS=default;amdgcn-amd-amdhsa",
+                "-DCOMPILER_RT_BUILD_ORC=OFF",
+                "-DCOMPILER_RT_BUILD_XRAY=OFF",
+                "-DCOMPILER_RT_BUILD_MEMPROF=OFF",
+                "-DCOMPILER_RT_BUILD_LIBFUZZER=OFF",
+                "-DCOMPILER_RT_BUILD_SANITIZERS=ON",
+                "-DLLVM_ENABLE_PROJECTS=clang;lld;mlir;flang;llvm",
+                "-DLLVM_ENABLE_RUNTIMES=flang-rt;offload;compiler-rt;openmp",
+            ]
+        )
+
+    with w.step("compile-openmp", halt_on_fail=True):
+        w.run_ninja(builddir=llvmbuilddir, ccache_stats=True)
+
+    with w.step("test-openmp"):
+        w.run_ninja(
+            ["check-openmp"], add_env={"HSA_ENABLE_SDMA": "0"}, builddir=llvmbuilddir
+        )
+
+    with w.step("Add check check-offload"):
+        w.run_ninja(
+            ["check-offload"], add_env={"HSA_ENABLE_SDMA": "0"}, builddir=llvmbuilddir
+        )
+
+    with w.step("LLVM: Install", halt_on_fail=True):
+        w.run_ninja(["install"], builddir=llvmbuilddir)
diff --git a/offload/ci/openmp-offload-amdgpu-runtime.py b/offload/ci/openmp-offload-amdgpu-runtime.py
new file mode 100755
index 0000000..e93f54f
--- /dev/null
+++ b/offload/ci/openmp-offload-amdgpu-runtime.py
@@ -0,0 +1,60 @@
+#! /usr/bin/env python3
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+import os
+import sys
+
+# Adapt to location in source tree
+llvmsrcroot = os.path.normpath(f"{__file__}/../../..")
+
+sys.path.insert(0, os.path.join(llvmsrcroot, ".ci/buildbot"))
+import worker
+
+llvmbuilddir = "llvm.build"
+llvminstalldir = "llvm.inst"
+
+with worker.run(
+    __file__,
+    llvmsrcroot,
+    clobberpaths=[llvmbuilddir, llvminstalldir],
+    workerjobs=32,
+) as w:
+    with w.step("configure-openmp", halt_on_fail=True):
+        w.run_command(
+            [
+                "cmake",
+                f"-S{w.in_llvmsrc('llvm')}",
+                f"-B{llvmbuilddir}",
+                "-GNinja",
+                "-DCMAKE_BUILD_TYPE=Release",
+                "-DLLVM_ENABLE_ASSERTIONS=ON",
+                f"-DLLVM_LIT_ARGS=-vv --show-unsupported --show-xfail -j {w.jobs} --time-tests --timeout 100",
+                f"-DCMAKE_INSTALL_PREFIX={w.in_workdir(llvminstalldir)}",
+                "-DCLANG_DEFAULT_LINKER=lld",
+                "-DLLVM_TARGETS_TO_BUILD=X86;AMDGPU",
+                "-DCMAKE_C_COMPILER_LAUNCHER=ccache",
+                "-DCMAKE_CXX_COMPILER_LAUNCHER=ccache",
+                "-DRUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES=compiler-rt;openmp",
+                "-DLLVM_RUNTIME_TARGETS=default;amdgcn-amd-amdhsa",
+                "-DLLVM_ENABLE_PROJECTS=clang;lld;llvm",
+                "-DLLVM_ENABLE_RUNTIMES=offload;compiler-rt;openmp",
+            ]
+        )
+
+    with w.step("compile-openmp", halt_on_fail=True):
+        w.run_ninja(builddir=llvmbuilddir, ccache_stats=True)
+
+    with w.step("test-openmp"):
+        w.run_ninja(
+            ["check-openmp"], add_env={"HSA_ENABLE_SDMA": "0"}, builddir=llvmbuilddir
+        )
+
+    with w.step("Add check check-offload"):
+        w.run_ninja(
+            ["check-offload"], add_env={"HSA_ENABLE_SDMA": "0"}, builddir=llvmbuilddir
+        )
+
+    with w.step("LLVM: Install", halt_on_fail=True):
+        w.run_ninja(["install"], builddir=llvmbuilddir)
diff --git a/offload/cmake/caches/FlangOffload.cmake b/offload/cmake/caches/FlangOffload.cmake
new file mode 100644
index 0000000..ed75df6
--- /dev/null
+++ b/offload/cmake/caches/FlangOffload.cmake
@@ -0,0 +1,13 @@
+set(LLVM_ENABLE_PROJECTS "clang;flang;mlir;clang-tools-extra;lld" CACHE STRING "")
+set(LLVM_ENABLE_RUNTIMES "compiler-rt;flang-rt;libunwind;libcxx;libcxxabi;openmp;offload" CACHE STRING "")
+set(LLVM_ENABLE_PER_TARGET_RUNTIME_DIR ON CACHE BOOL "")
+
+set(LLVM_RUNTIME_TARGETS default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda CACHE STRING "") 
+set(RUNTIMES_nvptx64-nvidia-cuda_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/NVPTX.cmake" CACHE STRING "")
+set(RUNTIMES_amdgcn-amd-amdhsa_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/AMDGPU.cmake" CACHE STRING "")
+set(RUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES "compiler-rt;libc;openmp;libcxx;libcxxabi;flang-rt" CACHE STRING "")
+set(RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES "compiler-rt;libc;openmp;libcxx;libcxxabi;flang-rt" CACHE STRING "")
+set(RUNTIMES_nvptx64-nvidia-cuda_FLANG_RT_LIBC_PROVIDER llvm CACHE STRING "")
+set(RUNTIMES_nvptx64-nvidia-cuda_FLANG_RT_LIBCXX_PROVIDER llvm CACHE STRING "")
+set(RUNTIMES_amdgcn-amd-amdhsa_FLANG_RT_LIBC_PROVIDER llvm CACHE STRING "")
+set(RUNTIMES_amdgcn-amd-amdhsa_FLANG_RT_LIBCXX_PROVIDER llvm CACHE STRING "")
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index f1c3386..0f98a44 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -45,7 +45,7 @@
 
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/Support/Format.h"
-#include "llvm/Support/circular_raw_ostream.h"
+#include "llvm/Support/raw_ostream.h"
 
 /// 32-Bit field data attributes controlling information presented to the user.
 enum OpenMPInfoType : uint32_t {
@@ -142,7 +142,7 @@ inline uint32_t getInfoLevel() { return getInfoLevelInternal().load(); }
 #define INFO(_flags, _id, ...)                                                 \
   do {                                                                         \
     if (::llvm::offload::debug::isDebugEnabled()) {                            \
-      DP(__VA_ARGS__);                                                         \
+      INFO_DEBUG_INT(_flags, _id, __VA_ARGS__);                                \
     } else if (getInfoLevel() & _flags) {                                      \
       INFO_MESSAGE(_id, __VA_ARGS__);                                          \
     }                                                                          \
@@ -174,6 +174,14 @@ private:
   bool ShouldEmitNewLineOnDestruction;
   bool NeedEndNewLine = false;
 
+  /// Buffer to reduce interference between different threads
+  /// writing at the same time to the underlying stream.
+  static constexpr size_t BufferSize = 256;
+  llvm::SmallString<BufferSize> Buffer;
+
+  // Stream to write into Buffer. Its flushed to Os upon destruction.
+  llvm::raw_svector_ostream BufferStrm;
+
   /// If the stream is muted, writes to it are ignored
   bool Muted = false;
 
@@ -203,13 +211,13 @@ private:
       NeedEndNewLine = true;
     }
   }
-  void emitPrefix() { Os.write(Prefix.c_str(), Prefix.size()); }
+  void emitPrefix() { BufferStrm.write(Prefix.c_str(), Prefix.size()); }
   void writeWithPrefix(StringRef Str) {
     if (ShouldPrefixNextString) {
       emitPrefix();
       ShouldPrefixNextString = false;
     }
-    Os.write(Str.data(), Str.size());
+    BufferStrm.write(Str.data(), Str.size());
   }
 
 public:
@@ -218,26 +226,29 @@ public:
                         bool ShouldEmitNewLineOnDestruction = true)
       : Prefix(std::move(Prefix)), Os(Os), BaseLevel(BaseLevel),
         ShouldPrefixNextString(ShouldPrefixNextString),
-        ShouldEmitNewLineOnDestruction(ShouldEmitNewLineOnDestruction) {
+        ShouldEmitNewLineOnDestruction(ShouldEmitNewLineOnDestruction),
+        BufferStrm(Buffer) {
     SetUnbuffered();
   }
   ~odbg_ostream() final {
     if (ShouldEmitNewLineOnDestruction && NeedEndNewLine)
-      Os << '\n';
+      BufferStrm << '\n';
+    Os << BufferStrm.str();
   }
   odbg_ostream(const odbg_ostream &) = delete;
   odbg_ostream &operator=(const odbg_ostream &) = delete;
-  odbg_ostream(odbg_ostream &&other) : Os(other.Os) {
+  odbg_ostream(odbg_ostream &&other) : Os(other.Os), BufferStrm(Buffer) {
     Prefix = std::move(other.Prefix);
     BaseLevel = other.BaseLevel;
     ShouldPrefixNextString = other.ShouldPrefixNextString;
     ShouldEmitNewLineOnDestruction = other.ShouldEmitNewLineOnDestruction;
     NeedEndNewLine = other.NeedEndNewLine;
     Muted = other.Muted;
+    BufferStrm << other.BufferStrm.str();
   }
 
   /// Forward the current_pos method to the underlying stream.
-  uint64_t current_pos() const final { return Os.tell(); }
+  uint64_t current_pos() const final { return BufferStrm.tell(); }
 
   /// Some of the `<<` operators expect an lvalue, so we trick the type
   /// system.
@@ -247,17 +258,8 @@ public:
   void shouldMute(const OnlyLevel Filter) { Muted = BaseLevel != Filter; }
 };
 
-/// dbgs - Return a circular-buffered debug stream.
-[[maybe_unused]] static llvm::raw_ostream &dbgs() {
-  // Do one-time initialization in a thread-safe way.
-  static struct dbgstream {
-    llvm::circular_raw_ostream strm;
-
-    dbgstream() : strm(llvm::errs(), "*** Debug Log Output ***\n", 0) {}
-  } thestrm;
-
-  return thestrm.strm;
-}
+/// dbgs - Return the debug stream for offload debugging (just llvm::errs()).
+[[maybe_unused]] static llvm::raw_ostream &dbgs() { return llvm::errs(); }
 
 #ifdef OMPTARGET_DEBUG
 
@@ -306,8 +308,14 @@ struct DebugSettings {
       return;
 
     Settings.Enabled = true;
-    if (EnvRef.equals_insensitive("all"))
-      return;
+
+    if (EnvRef.starts_with_insensitive("all")) {
+      auto Spec = parseDebugFilter(EnvRef);
+      if (Spec.Type.equals_insensitive("all")) {
+        Settings.DefaultLevel = Spec.Level;
+        return;
+      }
+    }
 
     if (!EnvRef.getAsInteger(10, Settings.DefaultLevel))
       return;
@@ -618,7 +626,38 @@ static inline odbg_ostream reportErrorStream() {
 #define FORMAT_TO_STR(Format, ...)                                             \
   ::llvm::omp::target::debug::formatToStr(Format __VA_OPT__(, ) __VA_ARGS__)
 
-#define DP(...) ODBG() << FORMAT_TO_STR(__VA_ARGS__);
+template <uint32_t InfoId> static constexpr const char *InfoIdToODT() {
+  constexpr auto getId = []() {
+    switch (InfoId) {
+    case OMP_INFOTYPE_KERNEL_ARGS:
+      return "KernelArgs";
+    case OMP_INFOTYPE_MAPPING_EXISTS:
+      return "MappingExists";
+    case OMP_INFOTYPE_DUMP_TABLE:
+      return "DumpTable";
+    case OMP_INFOTYPE_MAPPING_CHANGED:
+      return "MappingChanged";
+    case OMP_INFOTYPE_PLUGIN_KERNEL:
+      return "PluginKernel";
+    case OMP_INFOTYPE_DATA_TRANSFER:
+      return "DataTransfer";
+    case OMP_INFOTYPE_EMPTY_MAPPING:
+      return "EmptyMapping";
+    case OMP_INFOTYPE_ALL:
+      return "Default";
+    }
+    return static_cast<const char *>(nullptr);
+  };
+
+  constexpr const char *result = getId();
+  static_assert(result != nullptr, "Unknown InfoId being used");
+  return result;
+}
+
+// Transform the INFO id to the corresponding debug type and print the message
+#define INFO_DEBUG_INT(_flags, _id, ...)                                       \
+  ODBG(::llvm::omp::target::debug::InfoIdToODT<_flags>())                      \
+      << FORMAT_TO_STR(__VA_ARGS__);
 
 // Define default format for pointers
 static inline raw_ostream &operator<<(raw_ostream &Os, void *Ptr) {
@@ -627,9 +666,11 @@ static inline raw_ostream &operator<<(raw_ostream &Os, void *Ptr) {
 }
 
 #else
-#define DP(...)                                                                \
+
+#define INFO_DEBUG_INT(_flags, _id, ...)                                       \
   {                                                                            \
   }
+
 #endif // OMPTARGET_DEBUG
 
 // New REPORT macro in the same style as ODBG
diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h
index fd458fa0..5dddd00 100644
--- a/offload/include/omptarget.h
+++ b/offload/include/omptarget.h
@@ -80,6 +80,10 @@ enum tgt_map_type {
   // Attach pointer and pointee, after processing all other maps.
   // Applicable to map-entering directives. Does not change ref-count.
   OMP_TGT_MAPTYPE_ATTACH = 0x4000,
+  // When a lookup fails, fall back to using null as the translated pointer,
+  // instead of preserving the original pointer's value. Currently only
+  // useful in conjunction with RETURN_PARAM.
+  OMP_TGT_MAPTYPE_FB_NULLIFY = 0x8000,
   // descriptor for non-contiguous target-update
   OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000,
   // member of struct, member given by [16 MSBs] - 1
diff --git a/offload/liboffload/API/Queue.td b/offload/liboffload/API/Queue.td
index ededa9c..4008432 100644
--- a/offload/liboffload/API/Queue.td
+++ b/offload/liboffload/API/Queue.td
@@ -125,3 +125,15 @@ def olLaunchHostFunction : Function {
   ];
   let returns = [];
 }
+
+def olQueryQueue : Function {
+    let desc = "Query for queue work completion in a non-blocking manner.";
+    let details = [
+      "The function checks if a queue work has completed enqueued work without blocking the calling thread."
+    ];
+    let params = [
+        Param<"ol_queue_handle_t", "Queue", "handle of the queue", PARAM_IN>,
+        Param<"bool *", "IsQueueWorkCompleted", " A flag indicating if the queue work has completed", PARAM_OUT_OPTIONAL>
+    ];
+    let returns = [];
+}
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 4a4f665..6c37f5b 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1005,10 +1005,28 @@ Error olMemcpy_impl(ol_queue_handle_t Queue, void *DstPtr,
     if (auto Res =
             DstDevice->Device->dataSubmit(DstPtr, SrcPtr, Size, QueueImpl))
       return Res;
-  } else {
+  } else if (SrcDevice->Platform.Plugin == DstDevice->Platform.Plugin &&
+             SrcDevice->Platform.Plugin->isDataExchangable(
+                 SrcDevice->Device->getDeviceId(),
+                 DstDevice->Device->getDeviceId())) {
     if (auto Res = SrcDevice->Device->dataExchange(SrcPtr, *DstDevice->Device,
                                                    DstPtr, Size, QueueImpl))
       return Res;
+  } else {
+    if (Queue)
+      if (auto Res = olSyncQueue_impl(Queue))
+        return Res;
+
+    void *Buffer = malloc(Size);
+    if (!Buffer)
+      return createOffloadError(ErrorCode::OUT_OF_RESOURCES,
+                                "Couldn't allocate a buffer for transfer");
+    Error Res = SrcDevice->Device->dataRetrieve(Buffer, SrcPtr, Size, nullptr);
+    if (!Res)
+      Res = DstDevice->Device->dataSubmit(DstPtr, Buffer, Size, nullptr);
+
+    free(Buffer);
+    return Res;
   }
 
   return Error::success();
@@ -1231,5 +1249,17 @@ Error olMemUnregister_impl(ol_device_handle_t Device, void *Ptr) {
   return Device->Device->dataUnlock(Ptr);
 }
 
+Error olQueryQueue_impl(ol_queue_handle_t Queue, bool *IsQueueWorkCompleted) {
+  if (Queue->AsyncInfo->Queue) {
+    if (auto Err = Queue->Device->Device->queryAsync(Queue->AsyncInfo, false,
+                                                     IsQueueWorkCompleted))
+      return Err;
+  } else if (IsQueueWorkCompleted) {
+    // No underlying queue means there's no work to complete.
+    *IsQueueWorkCompleted = true;
+  }
+  return Error::success();
+}
+
 } // namespace offload
 } // namespace llvm
diff --git a/offload/libomptarget/OpenMP/InteropAPI.cpp b/offload/libomptarget/OpenMP/InteropAPI.cpp
index a694af1..2eaf376 100644
--- a/offload/libomptarget/OpenMP/InteropAPI.cpp
+++ b/offload/libomptarget/OpenMP/InteropAPI.cpp
@@ -219,7 +219,7 @@ omp_interop_val_t *__tgt_interop_get(ident_t *LocRef, int32_t InteropType,
 
   auto DeviceOrErr = PM->getDevice(DeviceNum);
   if (!DeviceOrErr) {
-    [[maybe_unused]] std::string ErrStr = toString(DeviceOrErr.takeError());
+    std::string ErrStr = toString(DeviceOrErr.takeError());
     ODBG(ODT_Interface) << "Couldn't find device " << DeviceNum
                         << " while constructing interop object: " << ErrStr;
     return omp_interop_none;
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index c71728c..b8edd79 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -83,10 +83,14 @@ int MappingInfoTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin,
                /*UseHoldRefCount=*/false, /*Name=*/nullptr,
                /*IsRefCountINF=*/true))
            .first->HDTT;
-  ODBG(ODT_Mapping) << "Creating new map entry: HstBase=" << NewEntry.HstPtrBase
-                    << ", HstBegin=" << NewEntry.HstPtrBegin
-                    << ", HstEnd=" << NewEntry.HstPtrEnd
-                    << ", TgtBegin=" << NewEntry.TgtPtrBegin
+  ODBG(ODT_Mapping) << "Creating new map entry: HstBase="
+                    << reinterpret_cast<void *>(NewEntry.HstPtrBase)
+                    << ", HstBegin="
+                    << reinterpret_cast<void *>(NewEntry.HstPtrBegin)
+                    << ", HstEnd="
+                    << reinterpret_cast<void *>(NewEntry.HstPtrEnd)
+                    << ", TgtBegin="
+                    << reinterpret_cast<void *>(NewEntry.TgtPtrBegin)
                     << ", DynRefCount=" << NewEntry.dynRefCountToStr()
                     << ", HoldRefCount=" << NewEntry.holdRefCountToStr();
   (void)NewEntry;
@@ -502,9 +506,11 @@ int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
                                          int64_t Size) {
   assert(Entry && "Trying to deallocate a null entry.");
 
-  ODBG(ODT_Mapping) << "Deleting tgt data " << Entry->TgtPtrBegin << " of size "
-                    << Size << " by freeing allocation "
-                    << "starting at " << Entry->TgtAllocBegin;
+  ODBG(ODT_Mapping) << "Deleting tgt data "
+                    << reinterpret_cast<void *>(Entry->TgtPtrBegin)
+                    << " of size " << Size << " by freeing allocation "
+                    << "starting at "
+                    << reinterpret_cast<void *>(Entry->TgtAllocBegin);
 
   void *Event = Entry->getEvent();
   if (Event && Device.destroyEvent(Event) != OFFLOAD_SUCCESS) {
diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp
index 924a986..41b653a 100644
--- a/offload/libomptarget/PluginManager.cpp
+++ b/offload/libomptarget/PluginManager.cpp
@@ -59,7 +59,7 @@ void PluginManager::deinit() {
       continue;
 
     if (auto Err = Plugin->deinit()) {
-      [[maybe_unused]] std::string InfoMsg = toString(std::move(Err));
+      std::string InfoMsg = toString(std::move(Err));
       ODBG(ODT_Deinit) << "Failed to deinit plugin: " << InfoMsg;
     }
     Plugin.release();
@@ -73,7 +73,7 @@ bool PluginManager::initializePlugin(GenericPluginTy &Plugin) {
     return true;
 
   if (auto Err = Plugin.init()) {
-    [[maybe_unused]] std::string InfoMsg = toString(std::move(Err));
+    std::string InfoMsg = toString(std::move(Err));
     ODBG(ODT_Init) << "Failed to init plugin: " << InfoMsg;
     return false;
   }
@@ -106,7 +106,7 @@ bool PluginManager::initializeDevice(GenericPluginTy &Plugin,
 
   auto Device = std::make_unique<DeviceTy>(&Plugin, UserId, DeviceId);
   if (auto Err = Device->init()) {
-    [[maybe_unused]] std::string InfoMsg = toString(std::move(Err));
+    std::string InfoMsg = toString(std::move(Err));
     ODBG(ODT_Init) << "Failed to init device " << DeviceId << ": " << InfoMsg;
     return false;
   }
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 676fda5..bd99ede 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -600,8 +600,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
-    bool UpdateRef =
-        !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0);
+    bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF);
 
     MappingInfoTy::HDTTMapAccessorTy HDTTMap =
         Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();
@@ -707,14 +706,20 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
         // to references to a local device pointer that refers to this device
         // address.
         //
-        // TODO: Add a new map-type bit to support OpenMP 6.1's `fb_nullify`
-        // and set the result to `nullptr - Delta`. Note that `fb_nullify` is
-        // already the default for `need_device_ptr`, but clang/flang do not
-        // support its codegen yet.
-        TgtPtrBase = reinterpret_cast<void *>(
-            reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
-        ODBG(ODT_Mapping) << "Returning host pointer " << TgtPtrBase
-                          << " as fallback (lookup failed)";
+        // OpenMP 6.1's `fb_nullify` fallback behavior: when the FB_NULLIFY bit
+        // is set by the compiler, e.g. for `use/need_device_ptr(fb_nullify)`),
+        // return `nullptr - Delta` when lookup fails.
+        if (ArgTypes[I] & OMP_TGT_MAPTYPE_FB_NULLIFY) {
+          TgtPtrBase = reinterpret_cast<void *>(
+              reinterpret_cast<uintptr_t>(nullptr) - Delta);
+          ODBG(ODT_Mapping) << "Returning offsetted null pointer " << TgtPtrBase
+                            << " as fallback (lookup failed)";
+        } else {
+          TgtPtrBase = reinterpret_cast<void *>(
+              reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
+          ODBG(ODT_Mapping) << "Returning host pointer " << TgtPtrBase
+                            << " as fallback (lookup failed)";
+        }
       }
       ArgsBase[I] = TgtPtrBase;
     }
@@ -1104,9 +1109,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     void *HstPtrBegin = Args[I];
     int64_t DataSize = ArgSizes[I];
     bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
-    bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
-                      (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) &&
-                     !(FromMapper && I == 0);
+    bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
+                     (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
     bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
@@ -1254,12 +1258,12 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
                   << "Restoring target descriptor " << ShadowPtr.TgtPtrAddr
                   << " to its original content (" << ShadowPtr.PtrSize
                   << " bytes), containing pointee address "
-                  << ShadowPtr.TgtPtrContent.data();
+                  << static_cast<const void *>(ShadowPtr.TgtPtrContent.data());
             } else {
               ODBG(ODT_Mapping)
                   << "Restoring target pointer " << ShadowPtr.TgtPtrAddr
                   << " to its original value "
-                  << ShadowPtr.TgtPtrContent.data();
+                  << static_cast<const void *>(ShadowPtr.TgtPtrContent.data());
             }
             Ret = Device.submitData(ShadowPtr.TgtPtrAddr,
                                     ShadowPtr.TgtPtrContent.data(),
@@ -1299,12 +1303,14 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
                     << "Restoring host descriptor " << ShadowPtr.HstPtrAddr
                     << " to its original content (" << ShadowPtr.PtrSize
                     << " bytes), containing pointee address "
-                    << ShadowPtr.HstPtrContent.data();
+                    << static_cast<const void *>(
+                           ShadowPtr.HstPtrContent.data());
               } else {
                 ODBG(ODT_Mapping)
                     << "Restoring host pointer " << ShadowPtr.HstPtrAddr
                     << " to its original value "
-                    << ShadowPtr.HstPtrContent.data();
+                    << static_cast<const void *>(
+                           ShadowPtr.HstPtrContent.data());
               }
               std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(),
                           ShadowPtr.PtrSize);
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index 287bb14..379c8ec 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -2431,7 +2431,10 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
   }
 
   /// Query for the completion of the pending operations on the async info.
-  Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override {
+  Error queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                       bool *IsQueueWorkCompleted) override {
+    if (IsQueueWorkCompleted)
+      *IsQueueWorkCompleted = false;
     AMDGPUStreamTy *Stream =
         reinterpret_cast<AMDGPUStreamTy *>(AsyncInfo.Queue);
     assert(Stream && "Invalid stream");
@@ -2444,11 +2447,16 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     if (!(*CompletedOrErr))
       return Plugin::success();
 
+    if (IsQueueWorkCompleted)
+      *IsQueueWorkCompleted = true;
     // Once the stream is completed, return it to stream pool and reset
     // AsyncInfo. This is to make sure the synchronization only works for its
     // own tasks.
-    AsyncInfo.Queue = nullptr;
-    return AMDGPUStreamManager.returnResource(Stream);
+    if (ReleaseQueue) {
+      AsyncInfo.Queue = nullptr;
+      return AMDGPUStreamManager.returnResource(Stream);
+    }
+    return Plugin::success();
   }
 
   /// Pin the host buffer and return the device pointer that should be used for
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index caf86a9..19db44c 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -854,8 +854,10 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
 
   /// Query for the completion of the pending operations on the __tgt_async_info
   /// structure in a non-blocking manner.
-  Error queryAsync(__tgt_async_info *AsyncInfo);
-  virtual Error queryAsyncImpl(__tgt_async_info &AsyncInfo) = 0;
+  Error queryAsync(__tgt_async_info *AsyncInfo, bool ReleaseQueue = true,
+                   bool *IsQueueWorkCompleted = nullptr);
+  virtual Error queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                               bool *IsQueueWorkCompleted) = 0;
 
   /// Check whether the architecture supports VA management
   virtual bool supportVAManagement() const { return false; }
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 4ec8366..807df0f 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -849,7 +849,8 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
 }
 Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
                                                       StringRef InputTgtImage) {
-  ODBG(OLDT_Init) << "Load data from image " << InputTgtImage.bytes_begin();
+  ODBG(OLDT_Init) << "Load data from image "
+                  << static_cast<const void *>(InputTgtImage.bytes_begin());
 
   std::unique_ptr<MemoryBuffer> Buffer;
   if (identify_magic(InputTgtImage) == file_magic::bitcode) {
@@ -1198,12 +1199,14 @@ Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo,
   return Plugin::success();
 }
 
-Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo) {
+Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo,
+                                  bool ReleaseQueue,
+                                  bool *IsQueueWorkCompleted) {
   if (!AsyncInfo || !AsyncInfo->Queue)
     return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                          "invalid async info queue");
 
-  return queryAsyncImpl(*AsyncInfo);
+  return queryAsyncImpl(*AsyncInfo, ReleaseQueue, IsQueueWorkCompleted);
 }
 
 Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) {
@@ -1656,9 +1659,10 @@ int32_t GenericPluginTy::is_initialized() const { return Initialized; }
 
 int32_t GenericPluginTy::isPluginCompatible(StringRef Image) {
   auto HandleError = [&](Error Err) -> bool {
-    [[maybe_unused]] std::string ErrStr = toString(std::move(Err));
-    ODBG(OLDT_Init) << "Failure to check validity of image " << Image.data()
-                    << ": " << ErrStr;
+    std::string ErrStr = toString(std::move(Err));
+    ODBG(OLDT_Init) << "Failure to check validity of image "
+                    << static_cast<const void *>(Image.data()) << ": "
+                    << ErrStr;
     return false;
   };
   switch (identify_magic(Image)) {
@@ -1685,8 +1689,9 @@ int32_t GenericPluginTy::isPluginCompatible(StringRef Image) {
 
 int32_t GenericPluginTy::isDeviceCompatible(int32_t DeviceId, StringRef Image) {
   auto HandleError = [&](Error Err) -> bool {
-    [[maybe_unused]] std::string ErrStr = toString(std::move(Err));
-    ODBG(OLDT_Init) << "Failure to check validity of image " << Image << ": "
+    std::string ErrStr = toString(std::move(Err));
+    ODBG(OLDT_Init) << "Failure to check validity of image "
+                    << static_cast<const void *>(Image.data()) << ": "
                     << ErrStr;
     return false;
   };
@@ -2069,7 +2074,7 @@ int32_t GenericPluginTy::use_auto_zero_copy(int32_t DeviceId) {
 int32_t GenericPluginTy::is_accessible_ptr(int32_t DeviceId, const void *Ptr,
                                            size_t Size) {
   auto HandleError = [&](Error Err) -> bool {
-    [[maybe_unused]] std::string ErrStr = toString(std::move(Err));
+    std::string ErrStr = toString(std::move(Err));
     ODBG(OLDT_Device) << "Failure while checking accessibility of pointer "
                       << Ptr << " for device " << DeviceId << ": " << ErrStr;
     return false;
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 621c90e..d5ab0b3 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -793,7 +793,10 @@ struct CUDADeviceTy : public GenericDeviceTy {
   }
 
   /// Query for the completion of the pending operations on the async info.
-  Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override {
+  Error queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                       bool *IsQueueWorkCompleted) override {
+    if (IsQueueWorkCompleted)
+      *IsQueueWorkCompleted = false;
     CUstream Stream = reinterpret_cast<CUstream>(AsyncInfo.Queue);
     CUresult Res = cuStreamQuery(Stream);
 
@@ -801,12 +804,16 @@ struct CUDADeviceTy : public GenericDeviceTy {
     if (Res == CUDA_ERROR_NOT_READY)
       return Plugin::success();
 
+    if (IsQueueWorkCompleted)
+      *IsQueueWorkCompleted = true;
     // Once the stream is synchronized and the operations completed (or an error
     // occurs), return it to stream pool and reset AsyncInfo. This is to make
     // sure the synchronization only works for its own tasks.
-    AsyncInfo.Queue = nullptr;
-    if (auto Err = CUDAStreamManager.returnResource(Stream))
-      return Err;
+    if (ReleaseQueue) {
+      AsyncInfo.Queue = nullptr;
+      if (auto Err = CUDAStreamManager.returnResource(Stream))
+        return Err;
+    }
 
     return Plugin::check(Res, "error in cuStreamQuery: %s");
   }
diff --git a/offload/plugins-nextgen/host/src/rtl.cpp b/offload/plugins-nextgen/host/src/rtl.cpp
index 81fbb67..6033796 100644
--- a/offload/plugins-nextgen/host/src/rtl.cpp
+++ b/offload/plugins-nextgen/host/src/rtl.cpp
@@ -336,7 +336,10 @@ struct GenELF64DeviceTy : public GenericDeviceTy {
 
   /// All functions are already synchronous. No need to do anything on this
   /// query function.
-  Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override {
+  Error queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                       bool *IsQueueWorkCompleted) override {
+    if (IsQueueWorkCompleted)
+      *IsQueueWorkCompleted = true;
     return Plugin::success();
   }
 
diff --git a/offload/plugins-nextgen/level_zero/include/L0Device.h b/offload/plugins-nextgen/level_zero/include/L0Device.h
index d14e710..001a41b 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Device.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Device.h
@@ -576,7 +576,8 @@ public:
                      AsyncInfoWrapperTy &AsyncInfoWrapper) override;
   Error synchronizeImpl(__tgt_async_info &AsyncInfo,
                         bool ReleaseQueue) override;
-  Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override;
+  Error queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                       bool *IsQueueWorkCompleted) override;
   Error dataSubmitImpl(void *TgtPtr, const void *HstPtr, int64_t Size,
                        AsyncInfoWrapperTy &AsyncInfoWrapper) override;
   Error dataRetrieveImpl(void *HstPtr, const void *TgtPtr, int64_t Size,
diff --git a/offload/plugins-nextgen/level_zero/src/L0Device.cpp b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
index 2cae1e4..4db3c4e 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Device.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Device.cpp
@@ -192,8 +192,7 @@ Error L0DeviceTy::initImpl(GenericPluginTy &Plugin) {
   CALL_ZE_RET_ERROR(zeDeviceGetCacheProperties, zeDevice, &Count,
                     &CacheProperties);
 
-  DeviceName =
-      std::string(DeviceProperties.name, sizeof(DeviceProperties.name));
+  DeviceName = std::string(DeviceProperties.name);
 
   ODBG(OLDT_Device) << "Found a GPU device, Name = " << DeviceProperties.name;
 
@@ -356,10 +355,15 @@ L0DeviceTy::hasPendingWorkImpl(AsyncInfoWrapperTy &AsyncInfoWrapper) {
   return true;
 }
 
-Error L0DeviceTy::queryAsyncImpl(__tgt_async_info &AsyncInfo) {
+Error L0DeviceTy::queryAsyncImpl(__tgt_async_info &AsyncInfo, bool ReleaseQueue,
+                                 bool *IsQueueWorkCompleted) {
+  if (IsQueueWorkCompleted)
+    *IsQueueWorkCompleted = true;
   const bool IsAsync = AsyncInfo.Queue && asyncEnabled();
   if (!IsAsync)
     return Plugin::success();
+  if (IsQueueWorkCompleted)
+    *IsQueueWorkCompleted = false;
 
   auto &Plugin = getPlugin();
   auto *AsyncQueue = static_cast<AsyncQueueTy *>(AsyncInfo.Queue);
@@ -367,6 +371,9 @@ Error L0DeviceTy::queryAsyncImpl(__tgt_async_info &AsyncInfo) {
   if (!AsyncQueue->WaitEvents.empty())
     return Plugin::success();
 
+  if (IsQueueWorkCompleted)
+    *IsQueueWorkCompleted = true;
+
   // Commit delayed USM2M copies.
   for (auto &USM2M : AsyncQueue->USM2MList) {
     std::copy_n(static_cast<const char *>(std::get<0>(USM2M)),
@@ -377,9 +384,11 @@ Error L0DeviceTy::queryAsyncImpl(__tgt_async_info &AsyncInfo) {
     std::copy_n(static_cast<char *>(std::get<0>(H2M)), std::get<2>(H2M),
                 static_cast<char *>(std::get<1>(H2M)));
   }
-  Plugin.releaseAsyncQueue(AsyncQueue);
-  getStagingBuffer().reset();
-  AsyncInfo.Queue = nullptr;
+  if (ReleaseQueue) {
+    Plugin.releaseAsyncQueue(AsyncQueue);
+    getStagingBuffer().reset();
+    AsyncInfo.Queue = nullptr;
+  }
 
   return Plugin::success();
 }
diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp
new file mode 100644
index 0000000..562e283
--- /dev/null
+++ b/offload/test/mapping/declare_mapper_target_checks.cpp
@@ -0,0 +1,145 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// XFAIL: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+#define TRUE 1
+#define FALSE 0
+
+struct TY1 {
+  int i1, i2, i3;
+  static constexpr auto name = "TY1";
+};
+struct TY2 {
+  int i1, i2, i3;
+  static constexpr auto name = "TY2";
+};
+
+// TY1 is not mapped, TY2 is
+#pragma omp declare mapper(TY2 t) map(to : t.i1) map(from : t.i3)
+
+struct TY3 {
+  TY2 n;
+  static constexpr auto name = "TY3";
+};
+struct TY4 {
+  int a;
+  TY2 n;
+  int b;
+  static constexpr auto name = "TY4";
+};
+
+template <typename T> int testType() {
+  T t1[2], t2[3], t3[4];
+  for (int i = 0; i < 2; i++)
+    t1[i].i1 = t3[i].i1 = 1;
+
+#pragma omp target map(tofrom : t1, t2, t3)
+  for (int i = 0; i < 2; i++) {
+    t1[i].i3 = t3[i].i3 = t1[i].i1;
+    t1[i].i1 = t3[i].i1 = 7;
+  }
+
+  for (int i = 0; i < 2; i++) {
+    if (t1[i].i3 != 1) {
+      printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i,
+             t1[i].i3, i, t1[i].i1);
+      return 1;
+    }
+    if (t3[i].i3 != 1) {
+      printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i,
+             t3[i].i3, i, t3[i].i1);
+      return 1;
+    }
+  }
+
+  int pt0 = omp_target_is_present(&t1[0], omp_get_default_device());
+  int pt1 = omp_target_is_present(&t2[1], omp_get_default_device());
+  int pt2 = omp_target_is_present(&t3[2], omp_get_default_device());
+
+  printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name,
+         pt0, pt1, pt2);
+  return pt0 + pt1 + pt2;
+}
+
+template <typename T> int testTypeNestedPtr(T t1[2], T t2[3], T t3[4]) {
+  for (int i = 0; i < 2; i++)
+    t1[i].n.i1 = t3[i].n.i1 = 1;
+
+#pragma omp target map(tofrom : t1[0 : 2], t2[0 : 3], t3[0 : 4])
+  for (int i = 0; i < 2; i++) {
+    t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1;
+    t1[i].n.i1 = t3[i].n.i1 = 7;
+  }
+
+  for (int i = 0; i < 2; i++) {
+    if (t1[i].n.i3 != t1[i].n.i1) {
+      printf("failed %s-ptr. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i,
+             t1[i].n.i3, i, t1[i].n.i1);
+      return 1;
+    }
+    if (t3[i].n.i3 != t3[i].n.i1) {
+      printf("failed %s-ptr. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i,
+             t3[i].n.i3, i, t3[i].n.i1);
+      return 1;
+    }
+  }
+
+  int pt0 = omp_target_is_present(&t1[0], omp_get_default_device());
+  int pt1 = omp_target_is_present(&t2[1], omp_get_default_device());
+  int pt2 = omp_target_is_present(&t3[2], omp_get_default_device());
+
+  printf("present check for %s-ptr: t1 %i, t2 %i, t3 %i, expected 3x 0\n",
+         T::name, pt0, pt1, pt2);
+  return pt0 + pt1 + pt2;
+}
+
+template <typename T> int testTypeNested() {
+  T t1[2], t2[3], t3[4];
+  testTypeNestedPtr(t1, t2, t3);
+  for (int i = 0; i < 2; i++)
+    t1[i].n.i1 = t3[i].n.i1 = 1;
+
+#pragma omp target map(tofrom : t1, t2, t3)
+  for (int i = 0; i < 2; i++) {
+    t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1;
+    t1[i].n.i1 = t3[i].n.i1 = 7;
+  }
+
+  for (int i = 0; i < 2; i++) {
+    if (t1[i].n.i3 != t1[i].n.i1) {
+      printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i,
+             t1[i].n.i3, i, t1[i].n.i1);
+      return 1;
+    }
+    if (t3[i].n.i3 != t3[i].n.i1) {
+      printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i,
+             t3[i].n.i3, i, t3[i].n.i1);
+      return 1;
+    }
+  }
+
+  int pt0 = omp_target_is_present(&t1[0], omp_get_default_device());
+  int pt1 = omp_target_is_present(&t2[1], omp_get_default_device());
+  int pt2 = omp_target_is_present(&t3[2], omp_get_default_device());
+
+  printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name,
+         pt0, pt1, pt2);
+  return pt0 + pt1 + pt2;
+}
+
+int main(int argc, char **argv) {
+  int r = 0;
+  r += testType<TY1>();
+  // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0
+  r += testType<TY2>();
+  // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0
+  r += testTypeNested<TY3>();
+  // CHECK: present check for TY3-ptr: t1 0, t2 0, t3 0, expected 3x 0
+  // CHECK: present check for TY3: t1 0, t2 0, t3 0, expected 3x 0
+  r += testTypeNested<TY4>();
+  // CHECK: present check for TY4-ptr: t1 0, t2 0, t3 0, expected 3x 0
+  // CHECK: present check for TY4: t1 0, t2 0, t3 0, expected 3x 0
+  return r;
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp
new file mode 100644
index 0000000..5c232d5
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp
@@ -0,0 +1,32 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value by
+// default.
+//
+// This is necessary because we must assume that the
+// pointee is device-accessible, even if it was not
+// previously mapped.
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(a)
+    printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp
new file mode 100644
index 0000000..fca0eee
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp
@@ -0,0 +1,28 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:          0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_nullify : a)
+    printf("%p\n", a); // OFFLOAD-NEXT:   (nil)
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp
new file mode 100644
index 0000000..51944c5
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp
@@ -0,0 +1,28 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+#include <stdio.h>
+
+int x = 0;
+
+struct ST {
+  int *a = &x;
+
+  void f1() {
+    printf("%p\n", a); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : a)
+    printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f1();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp
new file mode 100644
index 0000000..59a8fac
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp
@@ -0,0 +1,33 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value by
+// default.
+//
+// This is necessary because we must assume that the
+// pointee is device-accessible, even if it was not
+// previously mapped.
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(b)
+    printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp
new file mode 100644
index 0000000..65c7173
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp
@@ -0,0 +1,29 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:          0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_nullify : b)
+    printf("%p\n", b); // OFFLOAD-NEXT:   (nil)
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp
new file mode 100644
index 0000000..beeb752
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp
@@ -0,0 +1,29 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+#include <stdio.h>
+
+int x = 0;
+int *y = &x;
+
+struct ST {
+  int *&b = y;
+
+  void f2() {
+    printf("%p\n", b); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : b)
+    printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+  }
+};
+
+int main() {
+  ST s;
+  s.f2();
+}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp
index 33a3634..5be209a 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp
@@ -1,4 +1,8 @@
-// RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
 
 // Test that when a use_device_ptr lookup fails, the
 // privatized pointer retains its original value by
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp
new file mode 100644
index 0000000..984744cd
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp
@@ -0,0 +1,21 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+
+void f1() {
+  printf("%p\n", xp); // CHECK:          0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_nullify : xp)
+  printf("%p\n", xp); // OFFLOAD-NEXT:   (nil)
+                      // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp
new file mode 100644
index 0000000..197704f
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp
@@ -0,0 +1,21 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+
+void f1() {
+  printf("%p\n", xp); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : xp)
+  printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f1(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp
new file mode 100644
index 0000000..1060ed9
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp
@@ -0,0 +1,26 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value by
+// default.
+//
+// This is necessary because we must assume that the
+// pointee is device-accessible, even if it was not
+// previously mapped.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(xpr)
+  printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp
new file mode 100644
index 0000000..7fa76dd6
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer is set to null because of fb_nullify.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:          0x[[#%x,ADDR:]]
+  // FIXME: We won't get "nil" until we start privatizing xpr.
+#pragma omp target data use_device_ptr(fb_nullify : xpr)
+  printf("%p\n", xpr); // EXPECTED-OFFLOAD-NEXT:   (nil)
+                       // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+                       // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp
new file mode 100644
index 0000000..e7f8bd4
--- /dev/null
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp
@@ -0,0 +1,22 @@
+// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// Test that when a use_device_ptr lookup fails, the
+// privatized pointer retains its original value
+// because of fb_preserve.
+
+#include <stdio.h>
+int x;
+int *xp = &x;
+int *&xpr = xp;
+
+void f2() {
+  printf("%p\n", xpr); // CHECK:      0x[[#%x,ADDR:]]
+#pragma omp target data use_device_ptr(fb_preserve : xpr)
+  printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]]
+}
+
+int main() { f2(); }
diff --git a/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90 b/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90
new file mode 100644
index 0000000..6b87e81
--- /dev/null
+++ b/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90
@@ -0,0 +1,53 @@
+! Regression test for default mappers on nested derived types with allocatable
+! members when mapping a parent object and running an optimized target region.
+
+! REQUIRES: flang, amdgpu
+
+! RUN: %libomptarget-compile-fortran-generic -O3
+! RUN: %libomptarget-run-generic | %fcheck-generic
+
+program test_default_mapper_enter_data_teams_collapse
+  implicit none
+
+  type inner_type
+    real, allocatable :: data(:)
+  end type inner_type
+
+  type outer_type
+    type(inner_type) :: inner
+    character(len=19) :: desc = ' '
+  end type outer_type
+
+  type(outer_type) :: obj
+  integer, parameter :: n = 10
+  integer :: i, j
+  real :: expected, actual
+
+  allocate(obj%inner%data(n))
+  obj%inner%data = 0.0
+
+  !$omp target enter data map(to: obj)
+
+  !$omp target teams distribute parallel do collapse(2)
+  do i = 1, n
+    do j = 1, n
+      obj%inner%data(i) = real(i)
+    end do
+  end do
+  !$omp end target teams distribute parallel do
+
+  !$omp target exit data map(from: obj)
+
+  expected = real(n * (n + 1)) / 2.0
+  actual = sum(obj%inner%data)
+
+  if (abs(actual - expected) < 1.0e-6) then
+    print *, "PASS"
+  else
+    print *, "FAIL", actual, expected
+  end if
+
+  deallocate(obj%inner%data)
+end program test_default_mapper_enter_data_teams_collapse
+
+! CHECK: PASS
diff --git a/offload/test/offloading/fortran/recursive-default-mapper.f90 b/offload/test/offloading/fortran/recursive-default-mapper.f90
new file mode 100644
index 0000000..47b706d
--- /dev/null
+++ b/offload/test/offloading/fortran/recursive-default-mapper.f90
@@ -0,0 +1,40 @@
+! Offloading test for recursive default mapper emission
+! REQUIRES: flang, amdgpu
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+module recursive_mapper_mod
+  implicit none
+
+  type :: inner
+    integer :: value
+    type(inner), pointer :: next
+  end type inner
+
+  type :: outer
+    integer, allocatable :: arr(:)
+    type(inner), pointer :: head
+  end type outer
+
+contains
+
+end module recursive_mapper_mod
+
+program main
+  use recursive_mapper_mod
+  implicit none
+
+  type(outer) :: o
+
+  allocate(o%arr(2))
+  o%arr = [1, 2]
+
+  !$omp target map(tofrom: o)
+    o%arr(1) = o%arr(1) + 1
+    o%arr(2) = o%arr(2) + 1
+  !$omp end target
+
+  print *, o%arr(1), o%arr(2)
+end program main
+
+! CHECK: 2 3
diff --git a/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90 b/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90
new file mode 100644
index 0000000..d2d8f7a
--- /dev/null
+++ b/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90
@@ -0,0 +1,49 @@
+! Offload test that ensures defaultmap(tofrom: scalar) does not suppress
+! implicit default mapper generation for allocatable derived types.
+! REQUIRES: flang, amdgpu
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program defaultmap_implicit_mapper
+  implicit none
+
+  type :: payload_t
+    integer, allocatable :: arr(:)
+  end type payload_t
+
+  type(payload_t), allocatable :: obj
+  integer, parameter :: n = 8
+  integer :: i
+  integer :: scalar
+  logical :: ok
+
+  allocate(obj)
+  allocate(obj%arr(n))
+  obj%arr = 1
+  scalar = 2
+
+  !$omp target defaultmap(tofrom: scalar)
+    do i = 1, n
+      obj%arr(i) = obj%arr(i) + scalar
+    end do
+    scalar = 7
+  !$omp end target
+
+  ok = .true.
+  do i = 1, n
+    if (obj%arr(i) /= 3) ok = .false.
+  end do
+  if (scalar /= 7) ok = .false.
+
+  if (ok) then
+    print *, "Test passed!"
+  else
+    print *, "Test failed!"
+    print *, obj%arr
+    print *, scalar
+  end if
+
+  deallocate(obj%arr)
+  deallocate(obj)
+end program defaultmap_implicit_mapper
+
+! CHECK: Test passed!
diff --git a/offload/test/offloading/fortran/target-parameter-array.f90 b/offload/test/offloading/fortran/target-parameter-array.f90
new file mode 100644
index 0000000..b85fb06
--- /dev/null
+++ b/offload/test/offloading/fortran/target-parameter-array.f90
@@ -0,0 +1,131 @@
+! Offload test for parameter (constant) arrays and character scalars accessed
+! with dynamic indices/substrings in OpenMP target regions.
+
+! REQUIRES: flang, amdgpu
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+
+program test_parameter_mapping
+  implicit none
+  integer, parameter :: dp = selected_real_kind(15, 307)
+  logical :: all_tests_pass
+
+  all_tests_pass = .true.
+
+  ! Test 1: Parameter array with dynamic index
+  call test_param_array_dynamic_index(all_tests_pass)
+
+  ! Test 2: Integer parameter array
+  call test_int_param_array(all_tests_pass)
+
+  ! Test 3: Character scalar with dynamic substring
+  call test_char_substring(all_tests_pass)
+
+  ! Test 4: Verify scalar parameters work (inlined)
+  call test_scalar_param(all_tests_pass)
+
+  if (all_tests_pass) then
+    print *, "PASS"
+  else
+    print *, "FAIL"
+  endif
+
+contains
+
+! Test 1: Parameter array with dynamic index in target region
+subroutine test_param_array_dynamic_index(test_pass)
+  logical, intent(inout) :: test_pass
+  real(dp), parameter :: const_array(3) = [1.0_dp, 2.0_dp, 3.0_dp]
+  integer :: idx
+  real(dp) :: result
+  real(dp), parameter :: expected = 2.0_dp
+  real(dp), parameter :: tolerance = 1.0e-10_dp
+
+  idx = 2
+  result = 0.0_dp
+
+  !$omp target map(tofrom:result) map(to:idx)
+    ! Access parameter array with dynamic index
+    result = const_array(idx)
+  !$omp end target
+
+  if (abs(result - expected) > tolerance) then
+    print *, "Test 1 FAILED: expected", expected, "got", result
+    test_pass = .false.
+  endif
+end subroutine test_param_array_dynamic_index
+
+! Test 2: Integer parameter array with different indices
+subroutine test_int_param_array(test_pass)
+  logical, intent(inout) :: test_pass
+  integer, parameter :: int_array(4) = [10, 20, 30, 40]
+  integer :: idx1, idx2
+  integer :: result1, result2
+
+  idx1 = 1
+  idx2 = 4
+  result1 = 0
+  result2 = 0
+
+  !$omp target map(tofrom:result1, result2) map(to:idx1, idx2)
+    ! Access parameter array with different dynamic indices
+    result1 = int_array(idx1)
+    result2 = int_array(idx2)
+  !$omp end target
+
+  if (result1 /= 10 .or. result2 /= 40) then
+    print *, "Test 2 FAILED: expected 10, 40 got", result1, result2
+    test_pass = .false.
+  endif
+end subroutine test_int_param_array
+
+! Test 3: Character scalar parameter with dynamic substring access
+subroutine test_char_substring(test_pass)
+  logical, intent(inout) :: test_pass
+  character(len=20), parameter :: char_scalar = "constant_string_data"
+  integer :: start_idx, end_idx
+  character(len=8) :: result
+  character(len=8), parameter :: expected = "string_d"
+
+  start_idx = 10
+  end_idx = 17
+  result = ""
+
+  !$omp target map(tofrom:result) map(to:start_idx, end_idx)
+    ! Dynamic substring access - character scalar must be mapped
+    result = char_scalar(start_idx:end_idx)
+  !$omp end target
+
+  if (result /= expected) then
+    print *, "Test 3 FAILED: expected '", expected, "' got '", result, "'"
+    test_pass = .false.
+  endif
+end subroutine test_char_substring
+
+! Test 4: Scalar parameter (can be inlined, no mapping needed)
+subroutine test_scalar_param(test_pass)
+  logical, intent(inout) :: test_pass
+  integer, parameter :: scalar_const = 42
+  real(dp), parameter :: real_const = 3.14159_dp
+  integer :: int_result
+  real(dp) :: real_result
+  real(dp), parameter :: tolerance = 1.0e-5_dp
+
+  int_result = 0
+  real_result = 0.0_dp
+
+  !$omp target map(tofrom:int_result, real_result)
+    ! Scalar parameters should be inlined (no mapping needed)
+    int_result = scalar_const
+    real_result = real_const
+  !$omp end target
+
+  if (int_result /= 42 .or. abs(real_result - real_const) > tolerance) then
+    print *, "Test 4 FAILED: expected 42, 3.14159 got", int_result, real_result
+    test_pass = .false.
+  endif
+end subroutine test_scalar_param
+
+end program test_parameter_mapping
+
+! CHECK: PASS
diff --git a/offload/test/offloading/strided_multiple_update.c b/offload/test/offloading/strided_multiple_update_from.c
index 4f2df81..4f2df81 100644
--- a/offload/test/offloading/strided_multiple_update.c
+++ b/offload/test/offloading/strided_multiple_update_from.c
diff --git a/offload/test/offloading/strided_multiple_update_to.c b/offload/test/offloading/strided_multiple_update_to.c
new file mode 100644
index 0000000..bb16d7a
--- /dev/null
+++ b/offload/test/offloading/strided_multiple_update_to.c
@@ -0,0 +1,124 @@
+// This test checks that #pragma omp target update to(data1[0:3:4],
+// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
+// from the host to the device.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+// XFAIL: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int len = 12;
+  double data1[len], data2[len];
+
+  // Initialize host arrays
+  for (int i = 0; i < len; i++) {
+    data1[i] = i;
+    data2[i] = i * 10;
+  }
+
+  printf("original host array values:\n");
+  printf("data1:\n");
+  for (int i = 0; i < len; i++)
+    printf("%.1f\n", data1[i]);
+  printf("data2:\n");
+  for (int i = 0; i < len; i++)
+    printf("%.1f\n", data2[i]);
+
+  // CHECK: original host array values:
+  // CHECK-NEXT: data1:
+  // CHECK-NEXT: 0.0
+  // CHECK-NEXT: 1.0
+  // CHECK-NEXT: 2.0
+  // CHECK-NEXT: 3.0
+  // CHECK-NEXT: 4.0
+  // CHECK-NEXT: 5.0
+  // CHECK-NEXT: 6.0
+  // CHECK-NEXT: 7.0
+  // CHECK-NEXT: 8.0
+  // CHECK-NEXT: 9.0
+  // CHECK-NEXT: 10.0
+  // CHECK-NEXT: 11.0
+  // CHECK-NEXT: data2:
+  // CHECK-NEXT: 0.0
+  // CHECK-NEXT: 10.0
+  // CHECK-NEXT: 20.0
+  // CHECK-NEXT: 30.0
+  // CHECK-NEXT: 40.0
+  // CHECK-NEXT: 50.0
+  // CHECK-NEXT: 60.0
+  // CHECK-NEXT: 70.0
+  // CHECK-NEXT: 80.0
+  // CHECK-NEXT: 90.0
+  // CHECK-NEXT: 100.0
+  // CHECK-NEXT: 110.0
+
+#pragma omp target data map(tofrom : data1[0 : len], data2[0 : len])
+  {
+    // Initialize device arrays to 20
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++) {
+        data1[i] = 20.0;
+        data2[i] = 20.0;
+      }
+    }
+
+    // Modify host arrays for strided elements
+    data1[0] = 10.0;
+    data1[4] = 10.0;
+    data1[8] = 10.0;
+    data2[0] = 10.0;
+    data2[5] = 10.0;
+
+    // data1[0:3:4]  // indices 0,4,8
+    // data2[0:2:5]  // indices 0,5
+#pragma omp target update to(data1[0 : 3 : 4], data2[0 : 2 : 5])
+
+    // Verify on device by adding 5
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++)
+        data1[i] += 5.0;
+      for (int i = 0; i < len; i++)
+        data2[i] += 5.0;
+    }
+  }
+
+  printf("device array values after update to:\n");
+  printf("data1:\n");
+  for (int i = 0; i < len; i++)
+    printf("%.1f\n", data1[i]);
+  printf("data2:\n");
+  for (int i = 0; i < len; i++)
+    printf("%.1f\n", data2[i]);
+
+  // CHECK: device array values after update to:
+  // CHECK-NEXT: data1:
+  // CHECK-NEXT: 15.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 15.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 15.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: data2:
+  // CHECK-NEXT: 15.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 15.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+  // CHECK-NEXT: 25.0
+}
diff --git a/offload/test/offloading/strided_partial_update.c b/offload/test/offloading/strided_partial_update_from.c
index 4a2977a..4a2977a 100644
--- a/offload/test/offloading/strided_partial_update.c
+++ b/offload/test/offloading/strided_partial_update_from.c
diff --git a/offload/test/offloading/strided_partial_update_to.c b/offload/test/offloading/strided_partial_update_to.c
new file mode 100644
index 0000000..f9c960f
--- /dev/null
+++ b/offload/test/offloading/strided_partial_update_to.c
@@ -0,0 +1,74 @@
+// This test checks that #pragma omp target update to(data[0:4:3]) correctly
+// updates every third element (stride 3) from the host to the device, partially
+// across the array
+
+// RUN: %libomptarget-compile-run-and-check-generic
+// XFAIL: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int len = 11;
+  double data[len];
+
+  // Initialize on host
+  for (int i = 0; i < len; i++)
+    data[i] = i;
+
+  // Initial values
+  printf("original host array values:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+  printf("\n");
+
+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 2.000000
+  // CHECK: 3.000000
+  // CHECK: 4.000000
+  // CHECK: 5.000000
+  // CHECK: 6.000000
+  // CHECK: 7.000000
+  // CHECK: 8.000000
+  // CHECK: 9.000000
+  // CHECK: 10.000000
+
+#pragma omp target data map(tofrom : data[0 : len])
+  {
+    // Initialize device array to 20
+#pragma omp target
+    for (int i = 0; i < len; i++)
+      data[i] = 20.0;
+
+    // Modify host data for strided elements
+    data[0] = 10.0;
+    data[3] = 10.0;
+    data[6] = 10.0;
+    data[9] = 10.0;
+
+#pragma omp target update to(data[0 : 4 : 3]) // indices 0,3,6,9
+
+    // Verify on device by adding 5
+#pragma omp target
+    for (int i = 0; i < len; i++)
+      data[i] += 5.0;
+  }
+
+  printf("device array values after update to:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+  printf("\n");
+
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+}
diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update_from.c
index 9910bed..9910bed 100644
--- a/offload/test/offloading/strided_update.c
+++ b/offload/test/offloading/strided_update_from.c
diff --git a/offload/test/offloading/strided_update_to.c b/offload/test/offloading/strided_update_to.c
new file mode 100644
index 0000000..eca20ab
--- /dev/null
+++ b/offload/test/offloading/strided_update_to.c
@@ -0,0 +1,74 @@
+// This test checks that "update to" clause in OpenMP is supported when the
+// elements are updated in a non-contiguous manner. This test checks that
+// #pragma omp target update to(data[0:4:2]) correctly updates only every
+// other element (stride 2) from the host to the device
+
+// RUN: %libomptarget-compile-run-and-check-generic
+// XFAIL: intelgpu
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int len = 8;
+  double data[len];
+
+  // Initialize on host
+  for (int i = 0; i < len; i++) {
+    data[i] = i;
+  }
+
+  // Initial values
+  printf("original host array values:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+  printf("\n");
+
+#pragma omp target data map(tofrom : len, data[0 : len])
+  {
+    // Initialize device to 20
+#pragma omp target
+    for (int i = 0; i < len; i++) {
+      data[i] = 20.0;
+    }
+
+    // Modify host for strided elements
+    data[0] = 10.0;
+    data[2] = 10.0;
+    data[4] = 10.0;
+    data[6] = 10.0;
+
+#pragma omp target update to(data[0 : 4 : 2])
+
+    // Verify on device by adding 5
+#pragma omp target
+    for (int i = 0; i < len; i++) {
+      data[i] += 5.0;
+    }
+  }
+
+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 2.000000
+  // CHECK: 3.000000
+  // CHECK: 4.000000
+  // CHECK: 5.000000
+  // CHECK: 6.000000
+  // CHECK: 7.000000
+
+  printf("device array values after update to:\n");
+  for (int i = 0; i < len; i++)
+    printf("%f\n", data[i]);
+  printf("\n");
+
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+  // CHECK: 15.000000
+  // CHECK: 25.000000
+
+  return 0;
+}
diff --git a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
index 47a299f..718e9e7 100644
--- a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
+++ b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
@@ -15,6 +15,7 @@
 
 #include "llvm/Frontend/Offloading/Utility.h"
 #include "llvm/Support/CommandLine.h"
+#include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/JSON.h"
 #include "llvm/Support/MemoryBuffer.h"
 
@@ -60,11 +61,11 @@ int main(int argc, char **argv) {
       MemoryBuffer::getFile(InputFilename, /*isText=*/true,
                             /*RequiresNullTerminator=*/true);
   if (!KernelInfoMB)
-    report_fatal_error("Error reading the kernel info json file");
+    reportFatalUsageError("Error reading the kernel info json file");
   Expected<json::Value> JsonKernelInfo =
       json::parse(KernelInfoMB.get()->getBuffer());
   if (auto Err = JsonKernelInfo.takeError())
-    report_fatal_error("Cannot parse the kernel info json file");
+    reportFatalUsageError("Cannot parse the kernel info json file");
 
   auto NumTeamsJson =
       JsonKernelInfo->getAsObject()->getInteger("NumTeamsClause");
@@ -104,7 +105,7 @@ int main(int argc, char **argv) {
       MemoryBuffer::getFile(KernelEntryName + ".image", /*isText=*/false,
                             /*RequiresNullTerminator=*/false);
   if (!ImageMB)
-    report_fatal_error("Error reading the kernel image.");
+    reportFatalUsageError("Error reading the kernel image.");
 
   __tgt_device_image DeviceImage;
   DeviceImage.ImageStart = const_cast<char *>(ImageMB.get()->getBufferStart());
@@ -145,7 +146,7 @@ int main(int argc, char **argv) {
                             /*RequiresNullTerminator=*/false);
 
   if (!DeviceMemoryMB)
-    report_fatal_error("Error reading the kernel input device memory.");
+    reportFatalUsageError("Error reading the kernel input device memory.");
 
   // On AMD for currently unknown reasons we cannot copy memory mapped data to
   // device. This is a work-around.
@@ -178,14 +179,15 @@ int main(int argc, char **argv) {
                               /*isText=*/false,
                               /*RequiresNullTerminator=*/false);
     if (!OriginalOutputMB)
-      report_fatal_error("Error reading the kernel original output file, make "
-                         "sure LIBOMPTARGET_SAVE_OUTPUT is set when recording");
+      reportFatalUsageError(
+          "Error reading the kernel original output file, make sure "
+          "LIBOMPTARGET_SAVE_OUTPUT is set when recording");
     ErrorOr<std::unique_ptr<MemoryBuffer>> ReplayOutputMB =
         MemoryBuffer::getFile(KernelEntryName + ".replay.output",
                               /*isText=*/false,
                               /*RequiresNullTerminator=*/false);
     if (!ReplayOutputMB)
-      report_fatal_error("Error reading the kernel replay output file");
+      reportFatalUsageError("Error reading the kernel replay output file");
 
     StringRef OriginalOutput = OriginalOutputMB.get()->getBuffer();
     StringRef ReplayOutput = ReplayOutputMB.get()->getBuffer();
diff --git a/offload/unittests/OffloadAPI/CMakeLists.txt b/offload/unittests/OffloadAPI/CMakeLists.txt
index 546d7da..fce38c5 100644
--- a/offload/unittests/OffloadAPI/CMakeLists.txt
+++ b/offload/unittests/OffloadAPI/CMakeLists.txt
@@ -48,7 +48,8 @@ add_offload_unittest("queue"
     queue/olGetQueueInfo.cpp
     queue/olGetQueueInfoSize.cpp
     queue/olWaitEvents.cpp
-    queue/olLaunchHostFunction.cpp)
+    queue/olLaunchHostFunction.cpp
+    queue/olQueryQueue.cpp)
 
 add_offload_unittest("symbol"
     symbol/olGetSymbol.cpp
diff --git a/offload/unittests/OffloadAPI/queue/olQueryQueue.cpp b/offload/unittests/OffloadAPI/queue/olQueryQueue.cpp
new file mode 100644
index 0000000..4efd939
--- /dev/null
+++ b/offload/unittests/OffloadAPI/queue/olQueryQueue.cpp
@@ -0,0 +1,24 @@
+//===------- Offload API tests - olQueryQueue ----------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===--------------------------------------------------------------------===//
+
+#include "../common/Fixtures.hpp"
+#include <OffloadAPI.h>
+#include <gtest/gtest.h>
+
+using olQueryQueueTest = OffloadQueueTest;
+OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olQueryQueueTest);
+
+TEST_P(olQueryQueueTest, SuccessEmptyAsyncQueue) {
+  ASSERT_SUCCESS(olQueryQueue(Queue, nullptr));
+}
+
+TEST_P(olQueryQueueTest, SuccessEmptyAsyncQueueCheckResult) {
+  bool IsQueueWorkCompleted;
+  ASSERT_SUCCESS(olQueryQueue(Queue, &IsQueueWorkCompleted));
+  ASSERT_TRUE(IsQueueWorkCompleted);
+}