https://github.com/ro-i updated https://github.com/llvm/llvm-project/pull/215801
>From e98b79611494ed3c1c0e846630f9ee48faec9abd Mon Sep 17 00:00:00 2001 From: Robert Imschweiler <[email protected]> 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 6f8d771047505..6e5edd2d8fcd0 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 c303902ca3da7..108a8ca963507 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 baf6a171d9892..738c79c9fdc08 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h @@ -136,6 +136,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; @@ -174,7 +188,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 { @@ -185,6 +200,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 { @@ -222,6 +239,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 710e6ac272c1b..09b7398e0da37 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. @@ -2716,6 +2753,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: @@ -2725,21 +2768,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; @@ -2828,9 +2877,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"); @@ -3041,6 +3094,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); @@ -3064,14 +3122,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"); @@ -3086,7 +3145,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"); } @@ -3507,6 +3566,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) @@ -4002,12 +4066,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; @@ -4091,6 +4157,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy { if (auto Err = HostDevice->init()) return std::move(Err); + initSVM(HostAgents); + return NumDevices; } @@ -4178,6 +4246,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, @@ -4318,6 +4442,14 @@ struct AMDGPUPluginTy final : public GenericPluginTy { /// The device representing all HSA host agents. AMDHostDeviceTy *HostDevice; + + /// 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( @@ -4328,7 +4460,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; } @@ -4349,27 +4481,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, @@ -4581,6 +4711,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 0000000000000..9cf01aed56d83 --- /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 0000000000000..b60ef512d3a5b --- /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 6fae6a84a549a..251d051093900 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 b9aeaf321aba0cf59a6f26c6c32e7018689ec906 Mon Sep 17 00:00:00 2001 From: Robert Imschweiler <[email protected]> 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 09b7398e0da37..640edc55915e4 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -2755,9 +2755,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) { _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
