diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 67866a8d4b..a0118ff408 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -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, diff --git a/projects/clr/hipamd/rocclr/CMakeLists.txt b/projects/clr/hipamd/rocclr/CMakeLists.txt index 252ac5abfa..21ce5b1c58 100755 --- a/projects/clr/hipamd/rocclr/CMakeLists.txt +++ b/projects/clr/hipamd/rocclr/CMakeLists.txt @@ -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 diff --git a/projects/clr/hipamd/rocclr/hip_hcc.def.in b/projects/clr/hipamd/rocclr/hip_hcc.def.in index 415e40cb50..da43f5f51d 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.def.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.def.in @@ -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 diff --git a/projects/clr/hipamd/rocclr/hip_hcc.map.in b/projects/clr/hipamd/rocclr/hip_hcc.map.in index 179a31c55a..e66d4be92d 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.map.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.map.in @@ -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; diff --git a/projects/clr/hipamd/rocclr/hip_hmm.cpp b/projects/clr/hipamd/rocclr/hip_hmm.cpp new file mode 100644 index 0000000000..5342c18349 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_hmm.cpp @@ -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 +#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(hipMemAdviseSetReadMostly) == + amd::MemoryAdvice::SetReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetReadMostly) == + amd::MemoryAdvice::UnsetReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseSetPreferredLocation) == + amd::MemoryAdvice::SetPreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetPreferredLocation) == + amd::MemoryAdvice::UnsetPreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseSetAccessedBy) == + amd::MemoryAdvice::SetAccessedBy, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetAccessedBy) == + amd::MemoryAdvice::UnsetAccessedBy, "Enum mismatch with ROCclr!"); + +static_assert(static_cast(hipMemRangeAttributeReadMostly) == + amd::MemRangeAttribute::ReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemRangeAttributePreferredLocation) == + amd::MemRangeAttribute::PreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemRangeAttributeAccessedBy) == + amd::MemRangeAttribute::AccessedBy, "Enum mismatch with ROCclr!"); +static_assert(static_cast(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(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(&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(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; +} \ No newline at end of file diff --git a/projects/clr/hipamd/rocclr/hip_memory.cpp b/projects/clr/hipamd/rocclr/hip_memory.cpp index 2854b0d7eb..796cc0e491 100755 --- a/projects/clr/hipamd/rocclr/hip_memory.cpp +++ b/projects/clr/hipamd/rocclr/hip_memory.cpp @@ -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(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);