From 2bc6661737ac27f25e022dce3bb8b53ea5389a4f Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras Date: Mon, 4 Sep 2023 13:35:08 +0100 Subject: [PATCH] SWDEV-420222 - Removed implementation of hipMemcpyArray functions hipMemcpyArrayToArray, hipMemcpyFromArrayAsync, and hipMemcpyToArrayAsync are deprecated in cuda and are missing the headers in hip_runtime_api.h. Removed their implementation from hip_memory.cpp. Change-Id: I9720aec6241515c56c66b7e90a37b2ed53347eb2 --- hipamd/src/hip_graph.cpp | 76 -------------------------------- hipamd/src/hip_graph_capture.hpp | 8 ---- hipamd/src/hip_memory.cpp | 38 ---------------- 3 files changed, 122 deletions(-) diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 511f28c030..2706135163 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -400,44 +400,6 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz return hipSuccess; } -hipError_t capturehipMemcpyFromArrayAsync(hipStream_t& stream, void*& dst, hipArray_const_t& src, - size_t& wOffsetSrc, size_t& hOffsetSrc, size_t& count, - hipMemcpyKind& kind) { - ClPrint(amd::LOG_INFO, amd::LOG_API, - "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream); - if (src == nullptr || dst == nullptr) { - return hipErrorInvalidValue; - } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } - hip::Stream* s = reinterpret_cast(stream); - hip::GraphNode* pGraphNode; - hipMemcpy3DParms p = {}; - memset(&p, 0, sizeof(p)); - p.srcPos = {wOffsetSrc, hOffsetSrc, 0}; - p.kind = kind; - p.srcPtr.ptr = nullptr; - p.srcArray = const_cast(src); - - p.kind = kind; - p.dstPtr.ptr = dst; - p.dstArray = nullptr; // Ignored. - p.dstPtr.pitch = 0; - const size_t arrayHeight = (src->height != 0) ? src->height : 1; - const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / src->width) / hip::getElementSize(src); - p.extent = {widthInBytes / hip::getElementSize(p.srcArray), height, 1}; - hipError_t status = - ihipGraphAddMemcpyNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), - s->GetLastCapturedNodes().size(), &p); - if (status != hipSuccess) { - return status; - } - s->SetLastCapturedNode(pGraphNode); - return hipSuccess; -} - hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, size_t& wOffset, size_t& hOffset, const void*& src, size_t& spitch, size_t& width, size_t& height, hipMemcpyKind& kind) { @@ -473,44 +435,6 @@ hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, s return hipSuccess; } -hipError_t capturehipMemcpyToArrayAsync(hipStream_t& stream, hipArray_t& dst, size_t& wOffset, - size_t& hOffset, const void*& src, size_t& count, - hipMemcpyKind& kind) { - ClPrint(amd::LOG_INFO, amd::LOG_API, - "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream); - if (src == nullptr || dst == nullptr) { - return hipErrorInvalidValue; - } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } - hip::Stream* s = reinterpret_cast(stream); - hip::GraphNode* pGraphNode; - hipMemcpy3DParms p = {}; - memset(&p, 0, sizeof(p)); - p.dstPos = {wOffset, hOffset, 0}; - p.kind = kind; - p.dstPtr.ptr = nullptr; - p.dstArray = dst; // Ignored. - - p.kind = kind; - p.srcPtr.ptr = const_cast(src); - p.srcArray = nullptr; // Ignored. - p.srcPtr.pitch = 0; - const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; - const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / dst->width) / hip::getElementSize(dst); - p.extent = {widthInBytes / hip::getElementSize(p.dstArray), height, 1}; - hipError_t status = - ihipGraphAddMemcpyNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), - s->GetLastCapturedNodes().size(), &p); - if (status != hipSuccess) { - return status; - } - s->SetLastCapturedNode(pGraphNode); - return hipSuccess; -} - hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D*& pCopy) { ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream); diff --git a/hipamd/src/hip_graph_capture.hpp b/hipamd/src/hip_graph_capture.hpp index 443b19147d..9506ae99cd 100644 --- a/hipamd/src/hip_graph_capture.hpp +++ b/hipamd/src/hip_graph_capture.hpp @@ -52,18 +52,10 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz size_t& hOffsetSrc, size_t& width, size_t& height, hipMemcpyKind& kind); -hipError_t capturehipMemcpyFromArrayAsync(hipStream_t& stream, void*& dst, hipArray_const_t& src, - size_t& wOffsetSrc, size_t& hOffsetSrc, size_t& count, - hipMemcpyKind& kind); - hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, size_t& wOffset, size_t& hOffset, const void*& src, size_t& spitch, size_t& width, size_t& height, hipMemcpyKind& kind); -hipError_t capturehipMemcpyToArrayAsync(hipStream_t& stream, hipArray_t& dst, size_t& wOffset, - size_t& hOffset, const void*& src, size_t& count, - hipMemcpyKind& kind); - hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray*& srcArray, size_t& srcOffset, size_t& ByteCount); diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 3bee6c2721..949e45a7d1 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -3801,12 +3801,6 @@ hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOf HIP_RETURN_DURATION(ihipMemcpy2DArrayToArray(dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind, nullptr)); } -hipError_t hipMemcpyArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyArrayToArray, dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind); - CHECK_STREAM_CAPTURING(); - HIP_RETURN_DURATION(ihipMemcpy2DArrayToArray(dst, wOffsetDst, hOffsetDst, src, wOffsetSrc, hOffsetSrc, width, height, kind, nullptr)); -} - hipError_t hipMemcpy2DFromArray_common(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream=nullptr, @@ -3853,22 +3847,6 @@ hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_cons HIP_RETURN_DURATION(hipMemcpy2DFromArray_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true)); } -hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyFromArrayAsync, dst, src, wOffsetSrc, hOffsetSrc, count, kind, stream); - STREAM_CAPTURE(hipMemcpyFromArrayAsync, stream, dst, src, wOffsetSrc, hOffsetSrc, count, kind); - - if (src == nullptr) { - HIP_RETURN(hipErrorInvalidValue); - } - - const size_t arrayHeight = (src->height != 0) ? src->height : 1; - const size_t widthInBytes = count / arrayHeight; - - const size_t height = (count / src->width) / hip::getElementSize(src); - - HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffsetSrc, widthInBytes, height, kind, stream, true)); -} - hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind, stream); STREAM_CAPTURE(hipMemcpy2DToArrayAsync, stream, dst, wOffset, hOffset, src, spitch, width, height, @@ -3884,22 +3862,6 @@ hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOf HIP_RETURN_DURATION(hipMemcpy2DToArray_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true)); } -hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyToArrayAsync, dst, wOffset, hOffset, src, count, kind); - STREAM_CAPTURE(hipMemcpyToArrayAsync, stream, dst, wOffset, hOffset, src, count, kind); - - if (dst == nullptr) { - HIP_RETURN(hipErrorInvalidValue); - } - - const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; - const size_t widthInBytes = count / arrayHeight; - - const size_t height = (count / dst->width) / hip::getElementSize(dst); - - HIP_RETURN_DURATION(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, widthInBytes, height, kind, stream, true)); -} - hipError_t hipMemcpyAtoA(hipArray* dstArray, size_t dstOffset, hipArray* srcArray,