llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Robert Imschweiler (ro-i) <details> <summary>Changes</summary> Claude assisted with this patch. --- Patch is 21.28 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/215801.diff 7 Files Affected: - (modified) offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp (+1) - (modified) offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h (+1) - (modified) offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h (+23-1) - (modified) offload/plugins-nextgen/amdgpu/src/rtl.cpp (+162-21) - (added) offload/test/api/amdgpu_managed_memory_accessible.c (+31) - (added) offload/test/api/omp_device_managed_memory_ops.c (+68) - (modified) openmp/docs/design/Runtimes.md (+13) ``````````diff diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp index 724a5711a6be5ed..fa384233327b701 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp @@ -69,6 +69,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) diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h index c303902ca3da7f4..108a8ca96350797 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 c736ec0759841dc..220b54901e67c12 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h @@ -125,6 +125,21 @@ 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_GLOBAL_FLAG = 0, + 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; @@ -163,7 +178,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 { @@ -174,6 +190,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 { @@ -211,6 +229,10 @@ 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_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 5b52849af75da14..a5313c0bf329ac1 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -46,6 +46,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" @@ -245,6 +246,28 @@ 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; +} } // namespace hsa_utils /// Utility class representing generic resource references to AMDGPU resources. @@ -2723,6 +2746,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: @@ -2732,21 +2761,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; @@ -2835,9 +2870,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"); @@ -3045,8 +3084,15 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { Error dataFillImpl(void *TgtPtr, const void *PatternPtr, int64_t PatternSize, int64_t Size, AsyncInfoWrapperTy &AsyncInfoWrapper) override { + // hsa_amd_memory_fill fills SVM on the host, which runtimes before ROCr + // commit 9ab39639dd ("fix CPU side hsa_amd_memory_fill", May 2026) get + // wrong for patterns wider than a byte. SVM thus takes the slow case below. + auto InfoOrErr = hsa_utils::getPointerInfo(TgtPtr); + if (!InfoOrErr) + return InfoOrErr.takeError(); + // Fast case, where we can use the 4 byte hsa_amd_memory_fill - if (Size % 4 == 0 && + if (!hsa_utils::isSVM(*InfoOrErr) && Size % 4 == 0 && (PatternSize == 4 || PatternSize == 2 || PatternSize == 1)) { uint32_t Pattern; if (PatternSize == 1) { @@ -3062,7 +3108,11 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { llvm_unreachable("Invalid pattern size"); } - if (hasPendingWorkImpl(AsyncInfoWrapper)) { + auto HasPendingWorkOrErr = hasPendingWorkImpl(AsyncInfoWrapper); + if (!HasPendingWorkOrErr) + return HasPendingWorkOrErr.takeError(); + + if (*HasPendingWorkOrErr) { AMDGPUStreamTy *Stream = nullptr; if (auto Err = getStream(AsyncInfoWrapper, Stream)) return Err; @@ -3072,14 +3122,14 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { uint32_t Pattern; int64_t Size; }; - auto *Args = new MemFillArgsTy{TgtPtr, Pattern, Size / 4}; + auto Args = std::make_unique<MemFillArgsTy>( + MemFillArgsTy{TgtPtr, Pattern, Size / 4}); auto Fill = [](void *Data) { - MemFillArgsTy *Args = reinterpret_cast<MemFillArgsTy *>(Data); - assert(Args && "Invalid arguments"); + std::unique_ptr<MemFillArgsTy> Args( + static_cast<MemFillArgsTy *>(Data)); auto Status = hsa_amd_memory_fill(Args->Dst, Args->Pattern, Args->Size); - delete Args; auto Err = Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n"); if (Err) { @@ -3090,7 +3140,11 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { // hsa_amd_memory_fill doesn't signal completion using a signal, so use // the existing host callback logic to handle that instead - return Stream->pushHostCallback(Fill, Args); + if (auto Err = Stream->pushHostCallback(Fill, Args.get())) + return Err; + + Args.release(); + return Plugin::success(); } // If there is no pending work, do the fill synchronously auto Status = hsa_amd_memory_fill(TgtPtr, Pattern, Size / 4); @@ -3514,6 +3568,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) @@ -4098,6 +4157,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy { if (auto Err = HostDevice->init()) return std::move(Err); + initSVM(HostAgents); + return NumDevices; } @@ -4185,6 +4246,59 @@ 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. + if (SVMAttrs.empty() || + Alignment > llvm::sys::Process::getPageSizeEstimate()) + return nullptr; + + // Reserve addresses that are not registered in the driver, so that the SVM + // attributes govern the memory backing them. + void *Ptr = nullptr; + hsa_status_t Status = hsa_amd_vmem_address_reserve( + &Ptr, Size, /*address=*/0, HSA_AMD_VMEM_ADDRESS_NO_REGISTER); + if (auto Err = + Plugin::check(Status, "error in hsa_amd_vmem_address_reserve: %s")) + return std::move(Err); + + Status = + hsa_amd_svm_attributes_set(Ptr, Size, 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, Size); + 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, @@ -4325,6 +4439,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(GenericDeviceTy &Device, @@ -4356,7 +4478,8 @@ 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. + // resolve to the same tracker entry. SVM only reports its host base address, + // which all the agents share. hsa_amd_pointer_info_t HsaInfo{}; HsaInfo.size = sizeof(hsa_amd_pointer_info_t); hsa_status_t Status = hsa_amd_pointer_info( @@ -4364,19 +4487,20 @@ AMDGPUPluginContextTy::getAllocInfo(const void *Ptr) { /*num_agents_accessible=*/nullptr, /*accessible=*/nullptr); if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s")) return std::move(Err); + void *Base = hsa_utils::isSVM(HsaInfo) ? HsaInfo.hostBaseAddress + : HsaInfo.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, HsaInfo.sizeInBytes}; } Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice, @@ -4588,6 +4712,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 000000000000000..9cf01aed56d83cb --- /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 000000000000000..b60ef512d3a5b4c --- /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 6fae6a84a549a62..bb97a2fba7c5680 100644 --- a/openmp/docs/design/Runtimes.md +++ b/openmp/docs/design/Runtimes.md @@ -1335,6 +1335,7 @... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/215801 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
