[Openmp-commits] [llvm] [openmp] [offload] Use HSA SVM for AMDGPU shared memory (PR #215801)
Robert Imschweiler via Openmp-commits
openmp-commits at lists.llvm.org
Wed Oct 21 00:22:09 PDT 2026
https://github.com/ro-i updated https://github.com/llvm/llvm-project/pull/215801
>From d3347a2d11103cf6d772d5e540ce5fb963a197e6 Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Fri, 9 Oct 2026 07:48:18 -0500
Subject: [PATCH 1/2] [offload] Use HSA SVM for AMDGPU shared memory
---
.../amdgpu/dynamic_hsa/hsa.cpp | 2 +
.../plugins-nextgen/amdgpu/dynamic_hsa/hsa.h | 1 +
.../amdgpu/dynamic_hsa/hsa_ext_amd.h | 28 ++-
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 205 +++++++++++++++---
.../api/amdgpu_managed_memory_accessible.c | 31 +++
.../test/api/omp_device_managed_memory_ops.c | 68 ++++++
openmp/docs/design/Runtimes.md | 12 +
7 files changed, 317 insertions(+), 30 deletions(-)
create mode 100644 offload/test/api/amdgpu_managed_memory_accessible.c
create mode 100644 offload/test/api/omp_device_managed_memory_ops.c
diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp
index 6f8d7710475059..6e5edd2d8fcd0c 100644
--- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp
+++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp
@@ -67,6 +67,7 @@ DLWRAP(hsa_amd_agents_allow_access, 4)
DLWRAP(hsa_amd_memory_lock, 5)
DLWRAP(hsa_amd_memory_unlock, 1)
DLWRAP(hsa_amd_memory_fill, 3)
+DLWRAP(hsa_amd_svm_attributes_set, 4)
DLWRAP(hsa_amd_register_system_event_handler, 2)
DLWRAP(hsa_amd_signal_create, 5)
DLWRAP(hsa_amd_signal_async_handler, 5)
@@ -77,6 +78,7 @@ DLWRAP(hsa_code_object_reader_create_from_memory, 3)
DLWRAP(hsa_code_object_reader_destroy, 1)
DLWRAP(hsa_executable_load_agent_code_object, 5)
DLWRAP(hsa_amd_vmem_address_reserve, 4)
+DLWRAP(hsa_amd_vmem_address_reserve_align, 5)
DLWRAP(hsa_amd_vmem_address_free, 2)
DLWRAP(hsa_amd_vmem_handle_create, 5)
DLWRAP(hsa_amd_vmem_handle_release, 1)
diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
index c303902ca3da7f..108a8ca9635079 100644
--- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
+++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
@@ -123,6 +123,7 @@ typedef enum {
HSA_SYSTEM_INFO_VERSION_MINOR = 1,
HSA_SYSTEM_INFO_TIMESTAMP = 2,
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY = 3,
+ HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED = 0x201,
HSA_AMD_SYSTEM_INFO_EXT_VERSION_MAJOR = 0x207,
HSA_AMD_SYSTEM_INFO_EXT_VERSION_MINOR = 0x208,
} hsa_system_info_t;
diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h
index fe2d241b9754d1..3bd3e543882b4f 100644
--- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h
+++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h
@@ -137,6 +137,20 @@ hsa_status_t hsa_amd_memory_unlock(void* host_ptr);
hsa_status_t hsa_amd_memory_fill(void *ptr, uint32_t value, size_t count);
+typedef enum hsa_amd_svm_attribute_s {
+ HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE = 0x200,
+} hsa_amd_svm_attribute_t;
+
+typedef struct hsa_amd_svm_attribute_pair_s {
+ uint64_t attribute;
+ uint64_t value;
+} hsa_amd_svm_attribute_pair_t;
+
+hsa_status_t
+hsa_amd_svm_attributes_set(void *ptr, size_t size,
+ hsa_amd_svm_attribute_pair_t *attribute_list,
+ size_t attribute_count);
+
typedef enum hsa_amd_event_type_s {
HSA_AMD_GPU_MEMORY_FAULT_EVENT = 0,
} hsa_amd_event_type_t;
@@ -175,7 +189,8 @@ typedef enum {
typedef enum {
HSA_EXT_POINTER_TYPE_UNKNOWN = 0,
HSA_EXT_POINTER_TYPE_HSA = 1,
- HSA_EXT_POINTER_TYPE_LOCKED = 2
+ HSA_EXT_POINTER_TYPE_LOCKED = 2,
+ HSA_EXT_POINTER_TYPE_RESERVED_ADDR = 5
} hsa_amd_pointer_type_t;
typedef struct hsa_amd_pointer_info_s {
@@ -186,6 +201,8 @@ typedef struct hsa_amd_pointer_info_s {
size_t sizeInBytes;
void *userData;
hsa_agent_t agentOwner;
+ uint32_t global_flags;
+ bool registered;
} hsa_amd_pointer_info_t;
typedef enum {
@@ -223,6 +240,15 @@ hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t *queue,
hsa_status_t hsa_amd_vmem_address_reserve(void **va, size_t size,
uint64_t address, uint64_t flags);
+typedef enum hsa_amd_vmem_address_reserve_flag_s {
+ HSA_AMD_VMEM_ADDRESS_NO_REGISTER = (1UL << 0),
+} hsa_amd_vmem_address_reserve_flag_t;
+
+hsa_status_t hsa_amd_vmem_address_reserve_align(void **va, size_t size,
+ uint64_t address,
+ uint64_t alignment,
+ uint64_t flags);
+
hsa_status_t hsa_amd_vmem_address_free(void *va, size_t size);
hsa_status_t
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index ea9cb6fc5bcf6f..ff057d9501ab3b 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -10,6 +10,7 @@
//
//===----------------------------------------------------------------------===//
+#include <algorithm>
#include <atomic>
#include <cassert>
#include <cstddef>
@@ -46,6 +47,7 @@
#include "llvm/Support/FileOutputBuffer.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/Support/MemoryBuffer.h"
+#include "llvm/Support/Process.h"
#include "llvm/Support/Program.h"
#include "llvm/Support/Signals.h"
#include "llvm/Support/VersionTuple.h"
@@ -237,6 +239,41 @@ static std::optional<llvm::VersionTuple> getAMDExtensionVersion() {
return std::nullopt;
return llvm::VersionTuple(Major, Minor);
}
+
+/// Get the information of the allocation containing \p Ptr.
+static Expected<hsa_amd_pointer_info_t> getPointerInfo(const void *Ptr) {
+ hsa_amd_pointer_info_t Info{};
+ Info.size = sizeof(hsa_amd_pointer_info_t);
+ hsa_status_t Status = hsa_amd_pointer_info(Ptr, &Info, /*Allocator=*/nullptr,
+ /*num_agents_accessible=*/nullptr,
+ /*accessible=*/nullptr);
+ if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))
+ return std::move(Err);
+ return Info;
+}
+
+/// Whether \p Info describes shared virtual memory (SVM), i.e., addresses that
+/// are reserved without being registered in the driver. Runtimes before version
+/// 1.13 of the AMD extension do not report the registration.
+static bool isSVM(const hsa_amd_pointer_info_t &Info) {
+ return Info.type == HSA_EXT_POINTER_TYPE_RESERVED_ADDR &&
+ Info.size >= offsetof(hsa_amd_pointer_info_t, registered) +
+ sizeof(Info.registered) &&
+ !Info.registered;
+}
+
+/// Fill \p Count copies of the 4 byte \p Pattern at \p Ptr. hsa_amd_memory_fill
+/// fills SVM on the host, where runtimes before ROCr commit 9ab39639dd ("fix
+/// CPU side hsa_amd_memory_fill", May 2026) repeat only the lowest byte of the
+/// pattern.
+static hsa_status_t fillMemory(void *Ptr, uint32_t Pattern, size_t Count,
+ bool IsSVM) {
+ if (IsSVM && (Pattern & 0xff) * 0x01010101u != Pattern) {
+ std::fill_n(static_cast<uint32_t *>(Ptr), Count, Pattern);
+ return HSA_STATUS_SUCCESS;
+ }
+ return hsa_amd_memory_fill(Ptr, Pattern, Count);
+}
} // namespace hsa_utils
/// Utility class representing generic resource references to AMDGPU resources.
@@ -2717,6 +2754,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
if (TgtPtr == nullptr)
return Plugin::success();
+ // See allocate() for the registration of host / shared memory as pinned
+ // memory.
+ if (Kind == TARGET_ALLOC_HOST || Kind == TARGET_ALLOC_SHARED)
+ if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr))
+ return Err;
+
AMDGPUMemoryPoolTy *MemoryPool = nullptr;
switch (Kind) {
case TARGET_ALLOC_DEFAULT:
@@ -2726,21 +2769,27 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
case TARGET_ALLOC_HOST:
MemoryPool = &HostDevice.getFineGrainedMemoryPool();
break;
- case TARGET_ALLOC_SHARED:
+ case TARGET_ALLOC_SHARED: {
+ // Shared allocations are SVM allocations unless SVM is unsupported.
+ auto InfoOrErr = hsa_utils::getPointerInfo(TgtPtr);
+ if (!InfoOrErr)
+ return InfoOrErr.takeError();
+
+ if (hsa_utils::isSVM(*InfoOrErr)) {
+ hsa_status_t Status =
+ hsa_amd_vmem_address_free(TgtPtr, InfoOrErr->sizeInBytes);
+ return Plugin::check(Status, "error in hsa_amd_vmem_address_free: %s");
+ }
+
MemoryPool = &HostDevice.getFineGrainedMemoryPool();
break;
}
+ }
if (!MemoryPool)
return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
"no memory pool for the specified allocation kind");
- // See allocate() for the registration of host / shared memory as pinned
- // memory.
- if (Kind == TARGET_ALLOC_HOST || Kind == TARGET_ALLOC_SHARED)
- if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr))
- return Err;
-
if (auto Err = MemoryPool->deallocate(TgtPtr))
return Err;
@@ -2829,9 +2878,13 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
return std::move(Err);
// The buffer may be locked or allocated through HSA allocators. Assume that
- // the buffer is host pinned if the runtime reports a HSA type.
- if (Info.type != HSA_EXT_POINTER_TYPE_LOCKED &&
- Info.type != HSA_EXT_POINTER_TYPE_HSA)
+ // the buffer is host pinned if the runtime reports a HSA type. SVM is not
+ // pinned, but all the agents access it through its host address, so it
+ // must not be locked either.
+ if (hsa_utils::isSVM(Info))
+ Info.agentBaseAddress = Info.hostBaseAddress;
+ else if (Info.type != HSA_EXT_POINTER_TYPE_LOCKED &&
+ Info.type != HSA_EXT_POINTER_TYPE_HSA)
return false;
assert(Info.hostBaseAddress && "Invalid host pinned address");
@@ -3042,6 +3095,11 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
// Fast case, where we can use the 4 byte hsa_amd_memory_fill
if (Size % 4 == 0 &&
(PatternSize == 4 || PatternSize == 2 || PatternSize == 1)) {
+ auto InfoOrErr = hsa_utils::getPointerInfo(TgtPtr);
+ if (!InfoOrErr)
+ return InfoOrErr.takeError();
+ const bool IsSVM = hsa_utils::isSVM(*InfoOrErr);
+
uint32_t Pattern;
if (PatternSize == 1) {
auto *Byte = reinterpret_cast<const uint8_t *>(PatternPtr);
@@ -3065,14 +3123,15 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
void *Dst;
uint32_t Pattern;
int64_t Size;
+ bool IsSVM;
};
- auto *Args = new MemFillArgsTy{TgtPtr, Pattern, Size / 4};
+ auto *Args = new MemFillArgsTy{TgtPtr, Pattern, Size / 4, IsSVM};
auto Fill = [](void *Data) {
MemFillArgsTy *Args = reinterpret_cast<MemFillArgsTy *>(Data);
assert(Args && "Invalid arguments");
- auto Status =
- hsa_amd_memory_fill(Args->Dst, Args->Pattern, Args->Size);
+ auto Status = hsa_utils::fillMemory(Args->Dst, Args->Pattern,
+ Args->Size, Args->IsSVM);
delete Args;
auto Err =
Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n");
@@ -3087,7 +3146,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
return Stream->pushHostCallback(Fill, Args);
}
// If there is no pending work, do the fill synchronously
- auto Status = hsa_amd_memory_fill(TgtPtr, Pattern, Size / 4);
+ auto Status = hsa_utils::fillMemory(TgtPtr, Pattern, Size / 4, IsSVM);
return Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n");
}
@@ -3508,6 +3567,11 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))
return std::move(Err);
+ // SVM is accessible by all the agents, see initSVM().
+ if (hsa_utils::isSVM(Info))
+ return Size <=
+ Info.sizeInBytes - utils::getPtrDiff(Ptr, Info.hostBaseAddress);
+
// Checks if the pointer is known by HSA and accessible by the device
for (uint32_t i = 0; i < Count; i++) {
if (Agents[i].handle == getAgent().handle)
@@ -4006,12 +4070,14 @@ struct AMDGPUPluginContextTy final : public PluginContextTy {
Expected<PluginAllocInfoTy> getAllocInfo(const void *Ptr) override;
private:
- // Track each allocation's Kind so we can tell HOST from SHARED — HSA
- // backs both with the same host fine-grained pool.
- // TODO: drop once TARGET_ALLOC_SHARED has its own backing pool.
+ // Track each allocation's kind, owning device and requested size, which HSA
+ // cannot tell: it backs HOST and pool-backed SHARED allocations with the same
+ // host fine-grained pool, reports no owner for SVM, and rounds SVM sizes up
+ // to whole pages.
struct AllocInfo {
TargetAllocTy Kind;
GenericDeviceTy *Device;
+ size_t Size;
};
llvm::DenseMap<const void *, AllocInfo> Allocations;
std::mutex AllocationsMutex;
@@ -4103,6 +4169,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
"no AMDGPU host agents with kernel args and "
"fine-grained memory pools");
+ initSVM(HostAgents);
+
return NumDevices;
}
@@ -4201,6 +4269,62 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
return KernelAgents;
}
+ /// Enable SVM if the user, the driver and the runtime support it.
+ void initSVM(ArrayRef<hsa_agent_t> HostAgents) {
+ bool Supported = false;
+ if (!OMPX_UseSVM ||
+ hsa_system_get_info(HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED, &Supported) !=
+ HSA_STATUS_SUCCESS ||
+ !Supported)
+ return;
+
+ // SVM is recognized through the registration of its addresses, which the
+ // runtime reports since version 1.13 of the AMD extension.
+ auto Version = hsa_utils::getAMDExtensionVersion();
+ if (!Version || *Version < llvm::VersionTuple(1, 13))
+ return;
+
+ // Accesses may incur a page fault and the migration of the memory to the
+ // accessing agent.
+ for (hsa_agent_t Agent :
+ llvm::concat<const hsa_agent_t>(HostAgents, KernelAgents))
+ SVMAttrs.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, Agent.handle});
+ }
+
+ /// Allocate shared virtual memory (SVM), which is accessible by all the
+ /// agents and which the driver migrates between them on demand. Returns null
+ /// if SVM is unsupported.
+ Expected<void *> allocateSVM(size_t Size, size_t Alignment) {
+ // The reserved addresses are page aligned. Some runtimes report the wrong
+ // base address for larger alignments.
+ const size_t PageSize = llvm::sys::Process::getPageSizeEstimate();
+ if (SVMAttrs.empty() || Alignment > PageSize)
+ return nullptr;
+
+ // Reserve whole pages that are not registered in the driver, so that the
+ // SVM attributes govern the memory backing them.
+ const size_t ReservedSize = llvm::alignTo(Size, PageSize);
+ void *Ptr = nullptr;
+ hsa_status_t Status = hsa_amd_vmem_address_reserve_align(
+ &Ptr, ReservedSize, /*address=*/0, /*alignment=*/0,
+ HSA_AMD_VMEM_ADDRESS_NO_REGISTER);
+ if (auto Err = Plugin::check(
+ Status, "error in hsa_amd_vmem_address_reserve_align: %s"))
+ return std::move(Err);
+
+ Status = hsa_amd_svm_attributes_set(Ptr, ReservedSize, SVMAttrs.data(),
+ SVMAttrs.size());
+ if (auto Err =
+ Plugin::check(Status, "error in hsa_amd_svm_attributes_set: %s")) {
+ Status = hsa_amd_vmem_address_free(Ptr, ReservedSize);
+ return joinErrors(
+ std::move(Err),
+ Plugin::check(Status, "error in hsa_amd_vmem_address_free: %s"));
+ }
+
+ return Ptr;
+ }
+
/// Create an HSA signal for the RPC doorbell and return the fields needed
/// for the GPU to fire interrupts that wake the server thread.
Error initRPCDoorbell(uint64_t *&Value, uint64_t *&Mailbox,
@@ -4341,6 +4465,14 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
/// The host devices, one per HSA host agent.
llvm::SmallVector<AMDHostDeviceTy *> HostDevices;
+
+ /// Whether the user allows serving shared allocations through SVM.
+ BoolEnvar OMPX_UseSVM =
+ BoolEnvar("LIBOMPTARGET_AMDGPU_SHARED_ALLOC_SVM", true);
+
+ /// The attributes making SVM accessible by all the agents. Empty if SVM is
+ /// unsupported.
+ llvm::SmallVector<hsa_amd_svm_attribute_pair_t> SVMAttrs;
};
Expected<void *> AMDGPUPluginContextTy::allocate(
@@ -4351,7 +4483,7 @@ Expected<void *> AMDGPUPluginContextTy::allocate(
if (!PtrOrErr || !*PtrOrErr)
return PtrOrErr;
std::lock_guard<std::mutex> Lock(AllocationsMutex);
- Allocations[*PtrOrErr] = {Kind, &Device};
+ Allocations[*PtrOrErr] = {Kind, &Device, static_cast<size_t>(Size)};
return PtrOrErr;
}
@@ -4372,27 +4504,25 @@ Error AMDGPUPluginContextTy::deallocate(GenericDeviceTy &Device, void *Ptr,
Expected<PluginAllocInfoTy>
AMDGPUPluginContextTy::getAllocInfo(const void *Ptr) {
// HSA gives the base of the region containing Ptr, so interior pointers
- // resolve to the same tracker entry.
- hsa_amd_pointer_info_t HsaInfo{};
- HsaInfo.size = sizeof(hsa_amd_pointer_info_t);
- hsa_status_t Status = hsa_amd_pointer_info(
- const_cast<void *>(Ptr), &HsaInfo, /*Allocator=*/nullptr,
- /*num_agents_accessible=*/nullptr, /*accessible=*/nullptr);
- if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))
- return std::move(Err);
+ // resolve to the same tracker entry. SVM only reports its host base address,
+ // which all the agents share.
+ auto HsaInfoOrErr = hsa_utils::getPointerInfo(Ptr);
+ if (!HsaInfoOrErr)
+ return HsaInfoOrErr.takeError();
+ void *Base = hsa_utils::isSVM(*HsaInfoOrErr) ? HsaInfoOrErr->hostBaseAddress
+ : HsaInfoOrErr->agentBaseAddress;
AllocInfo Info;
{
std::lock_guard<std::mutex> Lock(AllocationsMutex);
- auto It = Allocations.find(HsaInfo.agentBaseAddress);
+ auto It = Allocations.find(Base);
if (It == Allocations.end())
return Plugin::error(ErrorCode::NOT_FOUND,
"pointer is not a known allocation in this context");
Info = It->second;
}
- return PluginAllocInfoTy{Info.Device, Info.Kind, HsaInfo.agentBaseAddress,
- HsaInfo.sizeInBytes};
+ return PluginAllocInfoTy{Info.Device, Info.Kind, Base, Info.Size};
}
Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
@@ -4601,6 +4731,23 @@ Expected<void *> AMDGPUDeviceTy::allocate(size_t Size, void *,
if (Size == 0)
return nullptr;
+ // Shared (managed) allocations use SVM, which the driver migrates between
+ // the host and the devices on demand. They fall back to the fine-grained host
+ // memory pool below if SVM is unsupported or cannot serve the alignment.
+ if (Kind == TARGET_ALLOC_SHARED) {
+ auto AllocOrErr =
+ static_cast<AMDGPUPluginTy &>(Plugin).allocateSVM(Size, Alignment);
+ if (!AllocOrErr)
+ return AllocOrErr.takeError();
+
+ // Register as pinned memory like the memory pool allocations below.
+ if (void *Alloc = *AllocOrErr) {
+ if (auto Err = PinnedAllocs.registerHostBuffer(Alloc, Alloc, Size))
+ return std::move(Err);
+ return Alloc;
+ }
+ }
+
// Find the correct memory pool.
AMDGPUMemoryPoolTy *MemoryPool = nullptr;
switch (Kind) {
diff --git a/offload/test/api/amdgpu_managed_memory_accessible.c b/offload/test/api/amdgpu_managed_memory_accessible.c
new file mode 100644
index 00000000000000..9cf01aed56d83c
--- /dev/null
+++ b/offload/test/api/amdgpu_managed_memory_accessible.c
@@ -0,0 +1,31 @@
+// RUN: %libomptarget-compile-run-and-check-amdgcn-amd-amdhsa
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+// Device managed memory is accessible by the device it was allocated for.
+
+#include <omp.h>
+#include <stdio.h>
+
+void *llvm_omp_target_alloc_shared(size_t, int);
+void llvm_omp_target_free_shared(void *, int);
+
+int main() {
+ const size_t Size = 1024;
+ const int Device = omp_get_default_device();
+
+ char *Shared = llvm_omp_target_alloc_shared(Size, Device);
+ if (!Shared) {
+ printf("FAIL: allocation\n");
+ return 1;
+ }
+
+ int Accessible =
+ omp_target_is_accessible(Shared, Size, Device) &&
+ omp_target_is_accessible(&Shared[Size / 2], Size / 2, Device);
+
+ llvm_omp_target_free_shared(Shared, Device);
+
+ // CHECK: PASS
+ printf(Accessible ? "PASS\n" : "FAIL\n");
+}
diff --git a/offload/test/api/omp_device_managed_memory_ops.c b/offload/test/api/omp_device_managed_memory_ops.c
new file mode 100644
index 00000000000000..b60ef512d3a5b4
--- /dev/null
+++ b/offload/test/api/omp_device_managed_memory_ops.c
@@ -0,0 +1,68 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// Check the data movement API on device managed memory.
+
+#include <omp.h>
+#include <stdio.h>
+
+void *llvm_omp_target_alloc_shared(size_t, int);
+void llvm_omp_target_free_shared(void *, int);
+
+int main() {
+ const int N = 128;
+ const int Device = omp_get_default_device();
+ const int Host = omp_get_initial_device();
+
+ int *Shared = llvm_omp_target_alloc_shared(N * sizeof(int), Device);
+ int Buffer[N];
+ int Failures = 0;
+
+ if (!Shared) {
+ printf("FAIL: allocation\n");
+ return 1;
+ }
+
+ // The host can access the allocation directly.
+ for (int I = 0; I < N; ++I)
+ Shared[I] = I;
+
+ // The device can access the allocation directly.
+#pragma omp target teams distribute parallel for device(Device) \
+ is_device_ptr(Shared)
+ for (int I = 0; I < N; ++I)
+ Shared[I] += 1;
+
+ for (int I = 0; I < N; ++I)
+ Failures += (Shared[I] != I + 1);
+
+ // Filling the whole allocation.
+ if (omp_target_memset(Shared, 0, N * sizeof(int), Device) != Shared)
+ ++Failures;
+ for (int I = 0; I < N; ++I)
+ Failures += (Shared[I] != 0);
+
+ // Filling a subrange of the allocation.
+ if (omp_target_memset(&Shared[1], 0xFF, sizeof(int), Device) != &Shared[1])
+ ++Failures;
+ Failures += (Shared[0] != 0) + (Shared[1] != -1) + (Shared[2] != 0);
+
+ // Copying into and out of the allocation.
+ for (int I = 0; I < N; ++I)
+ Buffer[I] = 2 * I;
+ Failures += (omp_target_memcpy(Shared, Buffer, N * sizeof(int), 0, 0, Device,
+ Host) != 0);
+ for (int I = 0; I < N; ++I)
+ Buffer[I] = 0;
+ Failures += (omp_target_memcpy(Buffer, Shared, N * sizeof(int), 0, 0, Host,
+ Device) != 0);
+ for (int I = 0; I < N; ++I)
+ Failures += (Buffer[I] != 2 * I);
+
+ llvm_omp_target_free_shared(Shared, Device);
+
+ // CHECK: PASS
+ if (!Failures)
+ printf("PASS\n");
+ else
+ printf("FAIL: %d\n", Failures);
+}
diff --git a/openmp/docs/design/Runtimes.md b/openmp/docs/design/Runtimes.md
index 6fae6a84a549a6..251d0510939005 100644
--- a/openmp/docs/design/Runtimes.md
+++ b/openmp/docs/design/Runtimes.md
@@ -1335,6 +1335,7 @@ There are several environment variables to change the behavior of the plugins:
- `LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES`
- `LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS`
- `LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT`
+- `LIBOMPTARGET_AMDGPU_SHARED_ALLOC_SVM`
The environment variables `LIBOMPTARGET_STACK_SIZE` and
`LIBOMPTARGET_HEAP_SIZE` are described in
@@ -1425,6 +1426,17 @@ HSA wait state within the AMDGPU plugin. For the duration of this value
the HSA runtime may busy wait. This can reduce overall latency.
The default value is `2000000`.
+#### LIBOMPTARGET_AMDGPU_SHARED_ALLOC_SVM
+
+This environment variable controls whether shared (managed) allocations, e.g.,
+the ones performed by `omp_target_alloc_shared`, are served through the HSA
+shared virtual memory (SVM) interface in the AMDGPU plugin. SVM allocations are
+backed by ordinary system memory that, on systems with XNACK enabled, the driver
+can migrate between the host and the devices on demand. If SVM is disabled, or
+if it is not supported by the driver or the HSA runtime, shared allocations are
+served from a fine-grained host memory pool and therefore remain resident on the
+host. The default value is `true`.
+
(remote_offloading_plugin)=
### Remote Offloading Plugin:
>From 03d4d174d23a0705bf3e0f725d8b10a12f1ded58 Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Wed, 14 Oct 2026 04:47:43 -0500
Subject: [PATCH 2/2] fix braces
---
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index ff057d9501ab3b..95ba232cbeaada 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -2756,9 +2756,10 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
// See allocate() for the registration of host / shared memory as pinned
// memory.
- if (Kind == TARGET_ALLOC_HOST || Kind == TARGET_ALLOC_SHARED)
+ if (Kind == TARGET_ALLOC_HOST || Kind == TARGET_ALLOC_SHARED) {
if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr))
return Err;
+ }
AMDGPUMemoryPoolTy *MemoryPool = nullptr;
switch (Kind) {
More information about the Openmp-commits
mailing list