AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Initial support for HIP managed memory

Browse Source 
- Call the new ROCclr interfaces for HMM

Change-Id: I2cd1bf438f712a9e9e328340e7d0c025257ca6c1


[ROCm/hip commit: f4211c3905]
This commit is contained in:
German Andryeyev
2020-06-13 00:52:36 -04:00
parent 969f1651fb
commit 61c65cf0b7
6 changed files with 377 additions and 45 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+134 -9
View File
@@ -178,8 +178,10 @@ enum hipLimit_t {
0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
///< allocation.
#define hipMemAttachGlobal 0x0
#define hipMemAttachHost 0x1
#define hipMemAttachGlobal 0x01 ///< Memory can be accessed by any stream on any device
#define hipMemAttachHost 0x02 ///< Memory cannot be accessed by any stream on any device
#define hipMemAttachSingle 0x04 ///< Memory can only be accessed by a single stream on
///< the associated device
#define hipDeviceMallocDefault 0x0
#define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device.
@@ -217,6 +219,41 @@ enum hipLimit_t {
#define hipCooperativeLaunchMultiDeviceNoPreSync 0x01
#define hipCooperativeLaunchMultiDeviceNoPostSync 0x02
#define hipCpuDeviceId ((int)-1)
#define hipInvalidDeviceId ((int)-2)
/*
* @brief HIP Memory Advise values
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemoryAdvise {
hipMemAdviseSetReadMostly = 1, ///< Data will mostly be read and only occassionally
///< be written to
hipMemAdviseUnsetReadMostly = 2, ///< Undo the effect of hipMemAdviseSetReadMostly
hipMemAdviseSetPreferredLocation = 3, ///< Set the preferred location for the data as
///< the specified device
hipMemAdviseUnsetPreferredLocation = 4, ///< Clear the preferred location for the data
hipMemAdviseSetAccessedBy = 5, ///< Data will be accessed by the specified device,
///< so prevent page faults as much as possible
hipMemAdviseUnsetAccessedBy = 6 ///< Let the Unified Memory subsystem decide on
///< the page faulting policy for the specified device
} hipMemoryAdvise;
/*
* @brief HIP range attributes
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemRangeAttribute {
hipMemRangeAttributeReadMostly = 1, ///< Whether the range will mostly be read and
///< only occassionally be written to
hipMemRangeAttributePreferredLocation = 2, ///< The preferred location of the range
hipMemRangeAttributeAccessedBy = 3, ///< Memory range has cudaMemAdviseSetAccessedBy
///< set for specified device
hipMemRangeAttributeLastPrefetchLocation = 4,///< The last location to which the range was prefetched
} hipMemRangeAttribute;
/*
* @brief hipJitOption
* @enum
@@ -1180,15 +1217,18 @@ hipError_t hipMemAllocHost(void** ptr, size_t size);
hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
/**
* @brief Allocates memory that will be automatically managed by the Unified Memory system.
* @brief Allocates memory that will be automatically managed by AMD HMM.
*
* @param[out] ptr Pointer to the allocated managed memory
* @param[in] size Requested memory size
* @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost
* @param [out] dev_ptr - pointer to allocated device memory
* @param [in] size - requested allocation size in bytes
* @param [in] flags - must be either hipMemAttachGlobal or hipMemAttachHost
* (defaults to hipMemAttachGlobal)
*
* @return #hipSuccess, #hipErrorOutOfMemory
* @returns #hipSuccess, #hipErrorMemoryAllocation, #hipErrorNotSupported, #hipErrorInvalidValue
*/
hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
hipError_t hipMallocManaged(void** dev_ptr,
size_t size,
unsigned int flags __dparm(hipMemAttachGlobal));
/**
* @brief Allocate device accessible page locked host memory [Deprecated]
@@ -3369,7 +3409,6 @@ hipError_t __hipPopCallConfiguration(dim3 *gridDim,
* @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
*
*/
hipError_t hipLaunchKernel(const void* function_address,
dim3 numBlocks,
dim3 dimBlocks,
@@ -3377,6 +3416,92 @@ hipError_t hipLaunchKernel(const void* function_address,
size_t sharedMemBytes __dparm(0),
hipStream_t stream __dparm(0));
/**
* @brief Prefetches memory to the specified destination device using AMD HMM.
*
* @param [in] dev_ptr pointer to be prefetched
* @param [in] count size in bytes for prefetching
* @param [in] device destination device to prefetch to
* @param [in] stream stream to enqueue prefetch operation
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemPrefetchAsync(const void* dev_ptr,
size_t count,
int device,
hipStream_t stream __dparm(0));
/**
* @brief Advise about the usage of a given memory range to AMD HMM.
*
* @param [in] dev_ptr pointer to memory to set the advice for
* @param [in] count size in bytes of the memory range
* @param [in] advice advice to be applied for the specified memory range
* @param [in] device device to apply the advice for
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemAdvise(const void* dev_ptr,
size_t count,
hipMemoryAdvise advice,
int device);
/**
* @brief Query an attribute of a given memory range in AMD HMM.
*
* @param [in/out] data a pointer to a memory location where the result of each
* attribute query will be written to
* @param [in] data_size the size of data
* @param [in] attribute the attribute to query
* @param [in] dev_ptr start of the range to query
* @param [in] count size of the range to query
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemRangeGetAttribute(void* data,
size_t data_size,
hipMemRangeAttribute attribute,
const void* dev_ptr,
size_t count);
/**
* @brief Query attributes of a given memory range in AMD HMM.
*
* @param [in/out] data a two-dimensional array containing pointers to memory locations
* where the result of each attribute query will be written to
* @param [in] data_sizes an array, containing the sizes of each result
* @param [in] attributes the attribute to query
* @param [in] num_attributes an array of attributes to query (numAttributes and the number
* of attributes in this array should match)
* @param [in] dev_ptr start of the range to query
* @param [in] count size of the range to query
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemRangeGetAttributes(void** data,
size_t* data_sizes,
hipMemRangeAttribute* attributes,
size_t num_attributes,
const void* dev_ptr,
size_t count);
/**
* @brief Attach memory to a stream asynchronously in AMD HMM.
*
* @param [in] stream - stream in which to enqueue the attach operation
* @param [in] dev_ptr - pointer to memory (must be a pointer to managed memory or
* to a valid host-accessible region of system-allocated memory)
* @param [in] length - length of memory (defaults to zero)
* @param [in] flags - must be one of cudaMemAttachGlobal, cudaMemAttachHost or
* cudaMemAttachSingle (defaults to cudaMemAttachSingle)
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipStreamAttachMemAsync(hipStream_t stream,
hipDeviceptr_t* dev_ptr,
size_t length __dparm(0),
unsigned int flags __dparm(hipMemAttachSingle));
#if __HIP_ROCclr__ || !defined(__HCC__)
//TODO: Move this to hip_ext.h
hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,
projects/hip/rocclr/CMakeLists.txt
+1
View File
@@ -81,6 +81,7 @@ add_library(hip64 OBJECT
hip_error.cpp
hip_event.cpp
hip_global.cpp
hip_hmm.cpp
hip_memory.cpp
hip_module.cpp
hip_peer.cpp
projects/hip/rocclr/hip_hcc.def.in
+5
View File
@@ -80,6 +80,7 @@ hipMallocManaged
hipArrayCreate
hipArray3DCreate
hipMallocArray
hipMemAdvise
hipMemAllocPitch
hipMallocPitch
hipMemcpy
@@ -111,7 +112,10 @@ hipMemGetAddressRange
hipGetSymbolAddress
hipGetSymbolSize
hipMemGetInfo
hipMemPrefetchAsync
hipMemPtrGetInfo
hipMemRangeGetAttribute
hipMemRangeGetAttributes
hipMemset
hipMemsetAsync
hipMemsetD8
@@ -154,6 +158,7 @@ hipGetDeviceFlags
hipSetDevice
hipSetDeviceFlags
hipStreamAddCallback
hipStreamAttachMemAsync
hipStreamCreate
hipStreamCreateWithFlags
hipStreamCreateWithPriority
projects/hip/rocclr/hip_hcc.map.in
+5
View File
@@ -62,6 +62,7 @@ global:
hipGetErrorName;
hipGetErrorString;
hipGetLastError;
hipMemAdvise;
hipMemAllocHost;
hipHostAlloc;
hipHostFree;
@@ -112,7 +113,10 @@ global:
hipGetSymbolAddress;
hipGetSymbolSize;
hipMemGetInfo;
hipMemPrefetchAsync;
hipMemPtrGetInfo;
hipMemRangeGetAttribute;
hipMemRangeGetAttributes;
hipMemset;
hipMemsetAsync;
hipMemsetD8;
@@ -154,6 +158,7 @@ global:
hipSetDevice;
hipSetDeviceFlags;
hipStreamAddCallback;
hipStreamAttachMemAsync;
hipStreamCreate;
hipStreamCreateWithFlags;
hipStreamCreateWithPriority;
projects/hip/rocclr/hip_hmm.cpp
+215
View File
@@ -0,0 +1,215 @@
/* Copyright (c) 2020-present Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_conversions.hpp"
#include "platform/context.hpp"
#include "platform/command.hpp"
#include "platform/memory.hpp"
// Forward declaraiton of a static function
static hipError_t ihipMallocManaged(void** ptr, size_t size);
// Make sure HIP defines match ROCclr to avoid double conversion
static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!");
static_assert(hipInvalidDeviceId == amd::InvalidDeviceId,
"Invalid device ID mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetReadMostly) ==
amd::MemoryAdvice::SetReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetReadMostly) ==
amd::MemoryAdvice::UnsetReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetPreferredLocation) ==
amd::MemoryAdvice::SetPreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetPreferredLocation) ==
amd::MemoryAdvice::UnsetPreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetAccessedBy) ==
amd::MemoryAdvice::SetAccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetAccessedBy) ==
amd::MemoryAdvice::UnsetAccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeReadMostly) ==
amd::MemRangeAttribute::ReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributePreferredLocation) ==
amd::MemRangeAttribute::PreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeAccessedBy) ==
amd::MemRangeAttribute::AccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeLastPrefetchLocation) ==
amd::MemRangeAttribute::LastPrefetchLocation, "Enum mismatch with ROCclr!");
// ================================================================================================
hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) {
HIP_INIT_API(hipMallocManaged, dev_ptr, size, flags);
if ((dev_ptr == nullptr) || (flags != hipMemAttachGlobal)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMallocManaged(dev_ptr, size), *dev_ptr);
}
// ================================================================================================
hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
hipStream_t stream) {
HIP_INIT_API(hipMemPrefetchAsync, dev_ptr, count, device, stream);
if ((dev_ptr == nullptr) || (count == 0) || (stream == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = nullptr;
bool cpu_access = (device == hipCpuDeviceId) ? true : false;
// Pick the specified stream or Null one from the provided device
if (stream != nullptr) {
queue = hip::getQueue(stream);
} else {
if (!cpu_access) {
queue = g_devices[device]->NullStream();
} else {
queue = hip::getCurrentDevice()->NullStream();
}
}
if (queue == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Command::EventWaitList waitList;
amd::SvmPrefetchAsyncCommand* command =
new amd::SvmPrefetchAsyncCommand(*queue, waitList, dev_ptr, count, cpu_access);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
if (!command->validateMemory()) {
delete command;
HIP_RETURN(hipErrorInvalidValue);
}
command->enqueue();
command->release();
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) {
HIP_INIT_API(hipMemAdvise, dev_ptr, count, advice, device);
if ((dev_ptr == nullptr) || (count == 0) || (device >= g_devices.size())) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* dev = g_devices[device]->devices()[0];
// Set the allocation attributes in AMD HMM
if (!dev->SetSvmAttributes(dev_ptr, count, static_cast<amd::MemoryAdvice>(advice))) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute,
const void* dev_ptr, size_t count) {
HIP_INIT_API(hipMemRangeGetAttribute, data, data_size, attribute, dev_ptr, count);
if ((data == nullptr) || (data_size == 0) || (dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Shouldn't matter for which device the interface is called
amd::Device* dev = g_devices[0]->devices()[0];
// Get the allocation attribute from AMD HMM
if (!dev->GetSvmAttributes(&data, &data_size, reinterpret_cast<int*>(&attribute), 1,
dev_ptr, count)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipMemRangeGetAttributes(void** data, size_t* data_sizes,
hipMemRangeAttribute* attributes, size_t num_attributes,
const void* dev_ptr, size_t count) {
HIP_INIT_API(hipMemRangeGetAttributes, data, data_sizes,
attributes, num_attributes, dev_ptr, count);
if ((data == nullptr) || (data_sizes == nullptr) || (attributes == nullptr) ||
(num_attributes == 0) || (dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Shouldn't matter for which device the interface is called
amd::Device* dev = g_devices[0]->devices()[0];
// Get the allocation attributes from AMD HMM
if (!dev->GetSvmAttributes(data, data_sizes, reinterpret_cast<int*>(attributes),
num_attributes, dev_ptr, count)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipStreamAttachMemAsync(hipStream_t stream, hipDeviceptr_t* dev_ptr,
size_t length, unsigned int flags) {
HIP_INIT_API(hipStreamAttachMemAsync, stream, dev_ptr, length, flags);
if ((stream == nullptr) || (dev_ptr == nullptr) || (length == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Unclear what should be done for this interface in AMD HMM, since it's generic SVM alloc
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
static hipError_t ihipMallocManaged(void** ptr, size_t size) {
if (size == 0) {
*ptr = nullptr;
return hipSuccess;
} else if (ptr == nullptr) {
return hipErrorInvalidValue;
}
assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid");
amd::Context& ctx = *hip::host_device->asContext();
const amd::Device& dev = *ctx.devices()[0];
// For now limit to the max allocation size on the device.
// The apps should be able to go over theCould you limit allocation in the future
if (dev.info().maxMemAllocSize_ < size) {
return hipErrorMemoryAllocation;
}
// Allocate SVM fine grain buffer with the forced host pointer, avoiding explicit memory
// allocation in the device driver
*ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR,
size, dev.info().memBaseAddrAlign_);
if (*ptr == nullptr) {
return hipErrorMemoryAllocation;
}
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMallocManaged ptr=0x%zx", getpid(),
std::this_thread::get_id(), *ptr);
return hipSuccess;
}
projects/hip/rocclr/hip_memory.cpp
+17 -36
View File
@@ -105,6 +105,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
return hipSuccess;
}
// ================================================================================================
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
amd::HostQueue& queue, bool isAsync = false) {
if (sizeBytes == 0) {
@@ -154,36 +155,26 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
*srcMemory->asBuffer(), sOffset, sizeBytes, dst);
isAsync = false;
} else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
if (queueDevice != srcMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
if ((kind == hipMemcpyDeviceToDevice) &&
// Check if the queue device doesn't match the device on any memory object. Hence
// it's a P2P transfer, because the app has requested access to another GPU
(srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0])) {
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
command->release();
return hipSuccess;
}
if (queueDevice != dstMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
// Make sure runtime has valid memory for the command execution. P2P access
// requires page table mapping on the current device to another GPU memory
if (!static_cast<amd::CopyMemoryP2PCommand*>(command)->validateMemory()) {
delete command;
return hipErrorInvalidValue;
}
command->release();
return hipSuccess;
} else {
command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
}
command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
}
if (command == nullptr) {
return hipErrorOutOfMemory;
}
@@ -201,6 +192,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
return hipSuccess;
}
// ================================================================================================
hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags);
@@ -245,17 +237,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr);
}
hipError_t hipMallocManaged(void** devPtr, size_t size,
unsigned int flags) {
HIP_INIT_API(hipMallocManaged, devPtr, size, flags);
if (flags != hipMemAttachGlobal) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *devPtr);
}
hipError_t hipFree(void* ptr) {
HIP_INIT_API(hipFree, ptr);