diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 3c6792e884..ba561472e9 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -177,6 +177,8 @@ enum hipLimit_t { 0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific ///< allocation. +#define hipDeviceMallocDefault 0x0 +#define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. //! Flags that can be used with hipHostRegister #define hipHostRegisterDefault 0x0 ///< Memory is Mapped and Portable @@ -1055,6 +1057,22 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void */ hipError_t hipMalloc(void** ptr, size_t size); +/** + * @brief Allocate memory on the default accelerator + * + * @param[out] ptr Pointer to the allocated memory + * @param[in] size Requested memory size + * @param[in] flags Type of memory allocation + * + * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. + * + * @return #hipSuccess, #hipErrorMemoryAllocation, #hipErrorInvalidValue (bad context, null *ptr) + * + * @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray, + * hipHostFree, hipHostMalloc + */ +hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags); + /** * @brief Allocate pinned host memory [Deprecated] * diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 8a291f3ea4..2955db5e37 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -2117,8 +2117,9 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, } } -void ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind, +bool ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind, bool resolveOn) { + bool retStatus = true; ihipCtx_t* ctx = this->getCtx(); const ihipDevice_t* device = ctx->getDevice(); @@ -2163,11 +2164,20 @@ void ihipStream_t::locked_copy2DSync(void* dst, const void* src, size_t width, s sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy); printPointerInfo(DB_COPY, " dst", dst, dstPtrInfo); printPointerInfo(DB_COPY, " src", src, srcPtrInfo); - +#if (__hcc_workweek__ >= 19101) + if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, + copyDevice ? ©Device->getDevice()->_acc : nullptr, + forceUnpinnedCopy)) { + tprintf(DB_COPY,"locked_copy2DSync failed to use SDMA\n"); + retStatus = false; + } +#else crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); +#endif } + return retStatus; } void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes) { @@ -2340,8 +2350,9 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes } } -void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind) +bool ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind) { + bool retStatus = true; const ihipCtx_t* ctx = this->getCtx(); if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) { @@ -2376,13 +2387,26 @@ void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, try { if (HIP_FORCE_SYNC_COPY) { +#if (__hcc_workweek__ >= 19101) + if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, + ©Device->getDevice()->_acc, + forceUnpinnedCopy)){ + tprintf(DB_COPY,"locked_copy2DASync with HIP_FORCE_SYNC_COPY failed to use SDMA\n"); + retStatus = false; + } +#else crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, ©Device->getDevice()->_acc, forceUnpinnedCopy); +#endif } else { - crit->_av.copy2d_async_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, + const auto& future = crit->_av.copy2d_async_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, ©Device->getDevice()->_acc); + if(!future.valid()) { + tprintf(DB_COPY, "locked_copy2DAsync failed to use SDMA\n"); + retStatus = false; + } } } catch (Kalmar::runtime_exception) { throw ihipException(hipErrorRuntimeOther); @@ -2397,10 +2421,20 @@ void ihipStream_t::locked_copy2DAsync(void* dst, const void* src, size_t width, } else { //Do sync 2D copy LockedAccessor_StreamCrit_t crit(_criticalData); +#if (__hcc_workweek__ >= 19101) + if(!crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, + copyDevice ? ©Device->getDevice()->_acc : nullptr, + forceUnpinnedCopy)){ + tprintf(DB_COPY, "locked_copy2DAsync Sync copy failed to use SDMA\n"); + retStatus = false; + } +#else crit->_av.copy2d_ext(src, dst, width, height, srcPitch, dstPitch, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); +#endif } + return retStatus; } //------------------------------------------------------------------------------------------------- diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 95a65b55c9..fb1588cc54 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -529,12 +529,12 @@ class ihipStream_t { void locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true); - void locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind, + bool locked_copy2DSync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind, bool resolveOn = true); void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); - void locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind); + bool locked_copy2DAsync(void* dst, const void* src, size_t width, size_t height, size_t srcPitch, size_t dstPitch, unsigned kind); void lockedSymbolCopySync(hc::accelerator& acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind); diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 9eebcfb28c..d540ab782b 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -264,6 +264,42 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { return ihipLogStatus(hip_status); } +hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) { + HIP_INIT_SPECIAL_API(hipExtMallocWithFlags, (TRACE_MEM), ptr, sizeBytes, flags); + HIP_SET_DEVICE(); + +#if (__hcc_workweek__ >= 19115) + hipError_t hip_status = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + // return NULL pointer when malloc size is 0 + if (sizeBytes == 0) { + *ptr = NULL; + hip_status = hipSuccess; + } else if ((ctx == nullptr) || (ptr == nullptr)) { + hip_status = hipErrorInvalidValue; + } else { + unsigned amFlags = 0; + if (flags & hipDeviceMallocFinegrained) { + amFlags = amDeviceFinegrained; + } else if (flags != hipDeviceMallocDefault) { + hip_status = hipErrorInvalidValue; + return ihipLogStatus(hip_status); + } + auto device = ctx->getWriteableDevice(); + *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/, + amFlags /*amFlags*/, 0 /*hipFlags*/, 0); + + if (sizeBytes && (*ptr == NULL)) { + hip_status = hipErrorMemoryAllocation; + } + } +#else + hipError_t hip_status = hipErrorMemoryAllocation; +#endif + + return ihipLogStatus(hip_status); +} hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); @@ -1573,7 +1609,10 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch stream->locked_copySync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind); } else { - stream->locked_copy2DSync(dst, src, width, height, spitch, dpitch, kind); + if(!stream->locked_copy2DSync(dst, src, width, height, spitch, dpitch, kind)){ + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + stream->locked_wait(); + } } } catch (ihipException& ex) { e = ex._code; @@ -1621,7 +1660,9 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind, stream); } else{ - stream->locked_copy2DAsync(dst, src, width, height, spitch, dpitch, kind); + if(!stream->locked_copy2DAsync(dst, src, width, height, spitch, dpitch, kind)){ + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + } } } catch (ihipException& ex) { e = ex._code;