From e9961fedd8d669341b3da054fc08cbc8ab65d47d Mon Sep 17 00:00:00 2001 From: Sarbojit Sarkar Date: Tue, 5 Jul 2022 03:53:31 +0000 Subject: [PATCH] SWDEV-323441 - Phase-II : per thread default stream Change-Id: I3c796ddaebcf0223d7faf50c425c1674de215f9d --- .../hip/amd_detail/amd_hip_runtime_pt_api.h | 72 +++++++ hipamd/src/amdhip.def | 18 ++ hipamd/src/hip_graph.cpp | 150 +++++++++---- hipamd/src/hip_hcc.def.in | 19 +- hipamd/src/hip_hcc.map.in | 18 ++ hipamd/src/hip_memory.cpp | 198 +++++++++++++----- hipamd/src/hip_stream.cpp | 28 ++- 7 files changed, 406 insertions(+), 97 deletions(-) diff --git a/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h b/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h index 971af7ba92..435a8f616b 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h +++ b/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h @@ -25,6 +25,8 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H #define HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H +#if defined(__clang__) && defined(__HIP__) + /// hipStreamPerThread implementation #if defined(HIP_API_PER_THREAD_DEFAULT_STREAM) #define __HIP_STREAM_PER_THREAD @@ -46,6 +48,17 @@ THE SOFTWARE. #define hipMemset2D __HIP_API_SPT(hipMemset2D) #define hipMemset3D __HIP_API_SPT(hipMemset3D) #define hipMemcpyAsync __HIP_API_SPT(hipMemcpyAsync) + #define hipMemset3DAsync __HIP_API_SPT(hipMemset3DAsync) + #define hipMemset2DAsync __HIP_API_SPT(hipMemset2DAsync) + #define hipMemsetAsync __HIP_API_SPT(hipMemsetAsync) + #define hipMemcpy3DAsync __HIP_API_SPT(hipMemcpy3DAsync) + #define hipMemcpy2DAsync __HIP_API_SPT(hipMemcpy2DAsync) + #define hipMemcpyFromSymbolAsync __HIP_API_SPT(hipMemcpyFromSymbolAsync) + #define hipMemcpyToSymbolAsync __HIP_API_SPT(hipMemcpyToSymbolAsync) + #define hipMemcpyFromArray __HIP_API_SPT(hipMemcpyFromArray) + #define hipMemcpy2DToArray __HIP_API_SPT(hipMemcpy2DToArray) + #define hipMemcpy2DFromArrayAsync __HIP_API_SPT(hipMemcpy2DFromArrayAsync) + #define hipMemcpy2DToArrayAsync __HIP_API_SPT(hipMemcpy2DToArrayAsync) // Stream APIs #define hipStreamSynchronize __HIP_API_SPT(hipStreamSynchronize) @@ -53,6 +66,7 @@ THE SOFTWARE. #define hipStreamGetFlags __HIP_API_SPT(hipStreamGetFlags) #define hipStreamGetPriority __HIP_API_SPT(hipStreamGetPriority) #define hipStreamWaitEvent __HIP_API_SPT(hipStreamWaitEvent) + #define hipStreamAddCallback __HIP_API_SPT(hipStreamAddCallback) // Event APIs #define hipEventRecord __HIP_API_SPT(hipEventRecord) @@ -60,6 +74,14 @@ THE SOFTWARE. // Launch APIs #define hipLaunchKernel __HIP_API_SPT(hipLaunchKernel) #define hipLaunchCooperativeKernel __HIP_API_SPT(hipLaunchCooperativeKernel) + + // Graph APIs + #define hipGraphLaunch __HIP_API_SPT(hipGraphLaunch) + #define hipStreamBeginCapture __HIP_API_SPT(hipStreamBeginCapture) + #define hipStreamEndCapture __HIP_API_SPT(hipStreamEndCapture) + #define hipStreamIsCapturing __HIP_API_SPT(hipStreamIsCapturing) + #define hipStreamGetCaptureInfo __HIP_API_SPT(hipStreamGetCaptureInfo) + #define hipStreamGetCaptureInfo_v2 __HIP_API_SPT(hipStreamGetCaptureInfo_v2) #endif hipError_t hipMemcpy_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); @@ -83,13 +105,45 @@ hipError_t hipMemcpy3D_spt(const struct hipMemcpy3DParms* p); hipError_t hipMemset_spt(void* dst, int value, size_t sizeBytes); +hipError_t hipMemsetAsync_spt(void* dst, int value, size_t sizeBytes, hipStream_t stream); + hipError_t hipMemset2D_spt(void* dst, size_t pitch, int value, size_t width, size_t height); +hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value, + size_t width, size_t height, hipStream_t stream); + +hipError_t hipMemset3DAsync_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream); + hipError_t hipMemset3D_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ); hipError_t hipMemcpyAsync_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream); +hipError_t hipMemcpy3DAsync_spt(const hipMemcpy3DParms* p, hipStream_t stream); + +hipError_t hipMemcpy2DAsync_spt(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream); + +hipError_t hipMemcpyFromSymbolAsync_spt(void* dst, const void* symbol, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream); + +hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream); + +hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, + size_t count, hipMemcpyKind kind); + +hipError_t hipMemcpy2DToArray_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, + size_t spitch, size_t width, size_t height, hipMemcpyKind kind); + +hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_const_t src, + size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, + hipMemcpyKind kind, hipStream_t stream); + +hipError_t hipMemcpy2DToArrayAsync_spt(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); + hipError_t hipStreamQuery_spt(hipStream_t stream); hipError_t hipStreamSynchronize_spt(hipStream_t stream); @@ -100,6 +154,11 @@ hipError_t hipStreamWaitEvent_spt(hipStream_t stream, hipEvent_t event, unsigned hipError_t hipStreamGetFlags_spt(hipStream_t stream, unsigned int* flags); +hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, void* userData, + unsigned int flags); + +hipError_t hipEventRecord_spt(hipEvent_t event, hipStream_t stream); + hipError_t hipLaunchCooperativeKernel_spt(const void* f, dim3 gridDim, dim3 blockDim, void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream); @@ -114,4 +173,17 @@ hipError_t hipLaunchKernel_spt(const void* function_address, #ifdef __cplusplus } #endif // extern "C" + +hipError_t hipGraphLaunch_spt(hipGraphExec_t graphExec, hipStream_t stream); +hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mode); +hipError_t hipStreamEndCapture_spt(hipStream_t stream, hipGraph_t* pGraph); +hipError_t hipStreamIsCapturing_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus); +hipError_t hipStreamGetCaptureInfo_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId); +hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out, hipGraph_t* graph_out, + const hipGraphNode_t** dependencies_out, + size_t* numDependencies_out); + +#endif //defined(__clang__) && defined(__HIP__) #endif //HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H \ No newline at end of file diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index 2c51b50002..46e3efb458 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -406,3 +406,21 @@ hipStreamWriteValue64 hipStreamWaitValue32 hipStreamWaitValue64 hipGetStreamDeviceId +hipGraphLaunch_spt +hipStreamBeginCapture_spt +hipStreamEndCapture_spt +hipStreamIsCapturing_spt +hipStreamGetCaptureInfo_spt +hipStreamGetCaptureInfo_v2_spt +hipStreamAddCallback_spt +hipMemsetAsync_spt +hipMemset2DAsync_spt +hipMemset3DAsync_spt +hipMemcpy3DAsync_spt +hipMemcpy2DAsync_spt +hipMemcpyFromSymbolAsync_spt +hipMemcpyToSymbolAsync_spt +hipMemcpyFromArray_spt +hipMemcpy2DToArray_spt +hipMemcpy2DFromArrayAsync_spt +hipMemcpy2DToArrayAsync_spt diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index a845ec0774..41cb9bfbb5 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -726,17 +726,27 @@ hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& return hipSuccess; } -hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { - HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); +hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { if (pCaptureStatus == nullptr || !hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } if (stream == nullptr) { *pCaptureStatus = hipStreamCaptureStatusNone; } else { *pCaptureStatus = reinterpret_cast(stream)->GetCaptureStatus(); } - HIP_RETURN(hipSuccess); + return hipSuccess; +} + +hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { + HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); + HIP_RETURN(hipStreamIsCapturing_common(stream, pCaptureStatus)); +} + +hipError_t hipStreamIsCapturing_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { + HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipStreamIsCapturing_common(stream, pCaptureStatus)); } hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) { @@ -754,22 +764,22 @@ hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) { HIP_RETURN_DURATION(hipSuccess); } -hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) { - HIP_INIT_API(hipStreamBeginCapture, stream, mode); +hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode) { if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } // capture cannot be initiated on legacy stream if (stream == nullptr) { - HIP_RETURN(hipErrorStreamCaptureUnsupported); + return hipErrorStreamCaptureUnsupported; } - if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) { - HIP_RETURN(hipErrorInvalidValue); + if (mode < hipStreamCaptureModeGlobal || + mode > hipStreamCaptureModeRelaxed) { + return hipErrorInvalidValue; } hip::Stream* s = reinterpret_cast(stream); // It can be initiated if the stream is not already in capture mode if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) { - HIP_RETURN(hipErrorIllegalState); + return hipErrorIllegalState; } s->SetCaptureGraph(new ihipGraph()); @@ -782,35 +792,45 @@ hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) amd::ScopedLock lock(g_captureStreamsLock); g_captureStreams.push_back(s); } - HIP_RETURN_DURATION(hipSuccess); + return hipSuccess; } -hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { - HIP_INIT_API(hipStreamEndCapture, stream, pGraph); +hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) { + HIP_INIT_API(hipStreamBeginCapture, stream, mode); + HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode)); +} + +hipError_t hipStreamBeginCapture_spt(hipStream_t stream, hipStreamCaptureMode mode) { + HIP_INIT_API(hipStreamBeginCapture, stream, mode); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipStreamBeginCapture_common(stream, mode)); +} + +hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) { if (pGraph == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } if (stream == nullptr) { - HIP_RETURN(hipErrorIllegalState); + return hipErrorIllegalState; } if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); + return hipErrorContextIsDestroyed; } hip::Stream* s = reinterpret_cast(stream); // Capture status must be active before endCapture can be initiated if (s->GetCaptureStatus() == hipStreamCaptureStatusNone) { - HIP_RETURN(hipErrorIllegalState); + return hipErrorIllegalState; } // Capture must be ended on the same stream in which it was initiated if (!s->IsOriginStream()) { - HIP_RETURN(hipErrorStreamCaptureUnmatched); + return hipErrorStreamCaptureUnmatched; } // If mode is not hipStreamCaptureModeRelaxed, hipStreamEndCapture must be called on the stream // from the same thread const auto& it = std::find(l_captureStreams.begin(), l_captureStreams.end(), s); if (s->GetCaptureMode() != hipStreamCaptureModeRelaxed) { if (it == l_captureStreams.end()) { - HIP_RETURN(hipErrorStreamCaptureWrongThread); + return hipErrorStreamCaptureWrongThread; } l_captureStreams.erase(it); } @@ -821,7 +841,7 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { // If capture was invalidated, due to a violation of the rules of stream capture if (s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) { *pGraph = nullptr; - HIP_RETURN(hipErrorStreamCaptureInvalidated); + return hipErrorStreamCaptureInvalidated; } // check if all parallel streams have joined // Nodes that are removed from the dependency set via API hipStreamUpdateCaptureDependencies do @@ -838,12 +858,23 @@ hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { } } if (foundInRemovedDep == false) { - HIP_RETURN(hipErrorStreamCaptureUnjoined); + return hipErrorStreamCaptureUnjoined; } } *pGraph = s->GetCaptureGraph(); // end capture on all streams/events part of graph capture - HIP_RETURN_DURATION(s->EndCapture()); + return s->EndCapture(); +} + +hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { + HIP_INIT_API(hipStreamEndCapture, stream, pGraph); + HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph)); +} + +hipError_t hipStreamEndCapture_spt(hipStream_t stream, hipGraph_t* pGraph) { + HIP_INIT_API(hipStreamEndCapture, stream, pGraph); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipStreamEndCapture_common(stream, pGraph)); } hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) { @@ -1038,12 +1069,22 @@ hipError_t ihipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { return graphExec->Run(stream); } +hipError_t hipGraphLaunch_common( hipGraphExec_t graphExec, hipStream_t stream) { + if (graphExec == nullptr || !hip::isValid(stream) || !hipGraphExec::isGraphExecValid(graphExec)) { + return hipErrorInvalidValue; + } + return ihipGraphLaunch(graphExec, stream); +} + hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { HIP_INIT_API(hipGraphLaunch, graphExec, stream); - if (graphExec == nullptr || !hip::isValid(stream) || !hipGraphExec::isGraphExecValid(graphExec)) { - HIP_RETURN(hipErrorInvalidValue); - } - HIP_RETURN_DURATION(ihipGraphLaunch(graphExec, stream)); + HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream)); +} + +hipError_t hipGraphLaunch_spt(hipGraphExec_t graphExec, hipStream_t stream) { + HIP_INIT_API(hipGraphLaunch, graphExec, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipGraphLaunch_common(graphExec, stream)); } hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) { @@ -1310,37 +1351,47 @@ hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGra HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(childGraph)); } -hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, +hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, unsigned long long* pId) { - HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId); if (pCaptureStatus == nullptr || !hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } if (stream == nullptr) { - HIP_RETURN(hipErrorUnknown); + return hipErrorUnknown; } hip::Stream* s = reinterpret_cast(stream); *pCaptureStatus = s->GetCaptureStatus(); if (*pCaptureStatus == hipStreamCaptureStatusActive && pId != nullptr) { *pId = s->GetCaptureID(); } - HIP_RETURN(hipSuccess); + return hipSuccess; } -hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, +hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId) { + HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId); + HIP_RETURN(hipStreamGetCaptureInfo_common(stream, pCaptureStatus, pId)); +} + +hipError_t hipStreamGetCaptureInfo_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus, + unsigned long long* pId) { + HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipStreamGetCaptureInfo_common(stream, pCaptureStatus, pId)); +} + +hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, unsigned long long* id_out, hipGraph_t* graph_out, const hipGraphNode_t** dependencies_out, size_t* numDependencies_out) { - HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out, - dependencies_out, numDependencies_out); if (captureStatus_out == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } if (stream == nullptr) { - HIP_RETURN(hipErrorUnknown); + return hipErrorUnknown; } if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } hip::Stream* s = reinterpret_cast(stream); *captureStatus_out = s->GetCaptureStatus(); @@ -1358,7 +1409,28 @@ hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus *numDependencies_out = s->GetLastCapturedNodes().size(); } } - HIP_RETURN(hipSuccess); + return hipSuccess; +} + +hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out, hipGraph_t* graph_out, + const hipGraphNode_t** dependencies_out, + size_t* numDependencies_out) { + HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out, + dependencies_out, numDependencies_out); + HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out, + dependencies_out, numDependencies_out)); +} + +hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out, + unsigned long long* id_out, hipGraph_t* graph_out, + const hipGraphNode_t** dependencies_out, + size_t* numDependencies_out) { + HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out, + dependencies_out, numDependencies_out); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out, + dependencies_out, numDependencies_out)); } hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies, diff --git a/hipamd/src/hip_hcc.def.in b/hipamd/src/hip_hcc.def.in index 27c11b0c33..2493ae9c12 100644 --- a/hipamd/src/hip_hcc.def.in +++ b/hipamd/src/hip_hcc.def.in @@ -421,4 +421,21 @@ hipStreamWriteValue64 hipStreamWaitValue32 hipStreamWaitValue64 hipDeviceSetLimit -hipGetStreamDeviceId \ No newline at end of file +hipGetStreamDeviceId +hipGraphLaunch_spt +hipStreamBeginCapture_spt +hipStreamEndCapture_spt +hipStreamIsCapturing_spt +hipStreamGetCaptureInfo_spt +hipStreamGetCaptureInfo_v2_spt +hipStreamAddCallback_spt +hipMemsetAsync_spt +hipMemset2DAsync_spt +hipMemset3DAsync_spt +hipMemcpy3DAsync_spt +hipMemcpy2DAsync_spt +hipMemcpyFromSymbolAsync_spt +hipMemcpyToSymbolAsync_spt +hipMemcpyFromArray_spt +hipMemcpy2DToArray_spt +hipMemcpy2DFromArrayAsync_spt diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 759f604ca2..c1490b4456 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -475,6 +475,24 @@ global: hipDeviceSetLimit; hiprtcGetBitcode; hiprtcGetBitcodeSize; + hipGraphLaunch_spt; + hipStreamBeginCapture_spt; + hipStreamEndCapture_spt; + hipStreamIsCapturing_spt; + hipStreamGetCaptureInfo_spt; + hipStreamGetCaptureInfo_v2_spt; + hipStreamAddCallback_spt; + hipMemsetAsync_spt; + hipMemset2DAsync_spt; + hipMemset3DAsync_spt; + hipMemcpy3DAsync_spt; + hipMemcpy2DAsync_spt; + hipMemcpyFromSymbolAsync_spt; + hipMemcpyToSymbolAsync_spt; + hipMemcpyFromArray_spt; + hipMemcpy2DToArray_spt; + hipMemcpy2DFromArrayAsync_spt; + hipMemcpy2DToArrayAsync_spt; local: *; } hip_5.2; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index ad3116ba65..0537949a87 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1261,14 +1261,12 @@ hipError_t hipMemcpyFromSymbol_spt(void* dst, const void* symbol, size_t sizeByt getPerThreadDefaultStream())); } -hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, +hipError_t hipMemcpyToSymbolAsync_common(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream); - STREAM_CAPTURE(hipMemcpyToSymbolAsync, stream, symbol, src, sizeBytes, offset, kind); - if (kind != hipMemcpyHostToDevice && kind != hipMemcpyDeviceToDevice) { - HIP_RETURN(hipErrorInvalidMemcpyDirection); + if (kind != hipMemcpyHostToDevice) { + return hipErrorInvalidMemcpyDirection; } size_t sym_size = 0; @@ -1279,29 +1277,53 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si return status; } /* Copy memory from source to destination address */ - HIP_RETURN_DURATION(hipMemcpyAsync(device_ptr, src, sizeBytes, kind, stream)); + return hipMemcpyAsync(device_ptr, src, sizeBytes, kind, stream); +} + +hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream); + HIP_RETURN_DURATION(hipMemcpyToSymbolAsync_common(symbol, src, sizeBytes, offset, kind, stream)); +} + +hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, sizeBytes, offset, kind, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpyToSymbolAsync_common(symbol, src, sizeBytes, offset, kind, stream)); +} + +hipError_t hipMemcpyFromSymbolAsync_common(void* dst, const void* symbol, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream) { + STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind); + + if (kind != hipMemcpyDeviceToHost) { + return hipErrorInvalidMemcpyDirection; + } + + size_t sym_size = 0; + hipDeviceptr_t device_ptr = nullptr; + + hipError_t status = ihipMemcpySymbol_validate(symbol, sizeBytes, offset, sym_size, device_ptr); + if (status != hipSuccess) { + return status; + } + + /* Copy memory from source to destination address */ + return hipMemcpyAsync(dst, device_ptr, sizeBytes, kind, stream); } hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyFromSymbolAsync, symbol, dst, sizeBytes, offset, kind, stream); + HIP_INIT_API(hipMemcpyFromSymbolAsync, dst, symbol, sizeBytes, offset, kind, stream); + HIP_RETURN_DURATION(hipMemcpyFromSymbolAsync_common(dst, symbol, sizeBytes, offset, kind, stream)); +} - STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind); - - if (kind != hipMemcpyDeviceToHost && kind != hipMemcpyDeviceToDevice) { - HIP_RETURN(hipErrorInvalidMemcpyDirection); - } - - size_t sym_size = 0; - hipDeviceptr_t device_ptr = nullptr; - - hipError_t status = ihipMemcpySymbol_validate(symbol, sizeBytes, offset, sym_size, device_ptr); - if (status != hipSuccess) { - return status; - } - - /* Copy memory from source to destination address */ - HIP_RETURN_DURATION(hipMemcpyAsync(dst, device_ptr, sizeBytes, kind, stream)); +hipError_t hipMemcpyFromSymbolAsync_spt(void* dst, const void* symbol, size_t sizeBytes, + size_t offset, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpyFromSymbolAsync, dst, symbol, sizeBytes, offset, kind, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpyFromSymbolAsync_common(dst, symbol, sizeBytes, offset, kind, stream)); } hipError_t hipMemcpyHtoD(hipDeviceptr_t dstDevice, @@ -2158,11 +2180,23 @@ hipError_t hipMemcpy2D_spt(void* dst, size_t dpitch, const void* src, size_t spi getPerThreadDefaultStream())); } +hipError_t hipMemcpy2DAsync_common(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream) { + STREAM_CAPTURE(hipMemcpy2DAsync, stream, dst, dpitch, src, spitch, width, height, kind); + return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, stream, true); +} + hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream); - STREAM_CAPTURE(hipMemcpy2DAsync, stream, dst, dpitch, src, spitch, width, height, kind); - HIP_RETURN_DURATION(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, stream, true)); + HIP_RETURN_DURATION(hipMemcpy2DAsync_common(dst, dpitch, src, spitch, width, height, kind, stream)); +} + +hipError_t hipMemcpy2DAsync_spt(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpy2DAsync, dst, dpitch, src, spitch, width, height, kind, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpy2DAsync_common(dst, dpitch, src, spitch, width, height, kind, stream)); } hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { @@ -2259,11 +2293,11 @@ hipError_t ihipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray_const_t src, return ihipMemcpyParam2D(&desc, stream, isAsync); } -hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind); +hipError_t hipMemcpyFromArray_common(void* dst, hipArray_const_t src, size_t wOffsetSrc, + size_t hOffset, size_t count, hipMemcpyKind kind, hipStream_t stream) { CHECK_STREAM_CAPTURING(); if (src == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } const size_t arrayHeight = (src->height != 0) ? src->height : 1; @@ -2271,7 +2305,18 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc const size_t height = (count / src->width) / hip::getElementSize(src); - HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, nullptr)); + return ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, stream); +} + +hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind); + HIP_RETURN_DURATION(hipMemcpyFromArray_common(dst, src, wOffsetSrc, hOffset, count, kind, nullptr)); +} + +hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind) { + HIP_INIT_API(hipMemcpyFromArray, dst, src, wOffsetSrc, hOffset, count, kind); + HIP_RETURN_DURATION(hipMemcpyFromArray_common(dst, src, wOffsetSrc, hOffset, count, kind, + getPerThreadDefaultStream())); } hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin, @@ -2429,12 +2474,20 @@ hipError_t hipMemcpy3D_spt(const hipMemcpy3DParms* p) { HIP_RETURN_DURATION(hipMemcpy3D_common(p, getPerThreadDefaultStream())); } +hipError_t hipMemcpy3DAsync_common(const hipMemcpy3DParms* p, hipStream_t stream) { + STREAM_CAPTURE(hipMemcpy3DAsync, stream, p); + return ihipMemcpy3D(p, stream, true); +} + hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) { HIP_INIT_API(hipMemcpy3DAsync, p, stream); + HIP_RETURN_DURATION(hipMemcpy3DAsync_common(p, stream)); +} - STREAM_CAPTURE(hipMemcpy3DAsync, stream, p); - - HIP_RETURN_DURATION(ihipMemcpy3D(p, stream, true)); +hipError_t hipMemcpy3DAsync_spt(const hipMemcpy3DParms* p, hipStream_t stream) { + HIP_INIT_API(hipMemcpy3DAsync, p, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpy3DAsync_common(p, stream)); } hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy) { @@ -2578,11 +2631,21 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { HIP_RETURN(hipMemset_common(dst, value, sizeBytes)); } -hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { - HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream); +hipError_t hipMemsetAsync_common(void* dst, int value, size_t sizeBytes, hipStream_t stream) { size_t valueSize = sizeof(int8_t); STREAM_CAPTURE(hipMemsetAsync, stream, dst, value, valueSize, sizeBytes); - HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true)); + return ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true); +} + +hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { + HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream); + HIP_RETURN(hipMemsetAsync_common(dst, value, sizeBytes, stream)); +} + +hipError_t hipMemsetAsync_spt(void* dst, int value, size_t sizeBytes, hipStream_t stream) { + HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipMemsetAsync_common(dst, value, sizeBytes, stream)); } hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) { @@ -2727,14 +2790,24 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t HIP_RETURN(hipMemset2D_common(dst, pitch, value, width, height)); } +hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value, + size_t width, size_t height, hipStream_t stream) { + STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height); + + return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true); +} hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream) { HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream); + HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream)); +} - STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height); - - HIP_RETURN(ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true)); +hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value, + size_t width, size_t height, hipStream_t stream) { + HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream)); } hipError_t hipMemset3D_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream=nullptr) { @@ -2753,12 +2826,20 @@ hipError_t hipMemset3D_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent ext HIP_RETURN(hipMemset3D_common(pitchedDevPtr, value, extent,stream)); } +hipError_t hipMemset3DAsync_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) { + STREAM_CAPTURE(hipMemset3DAsync, stream, pitchedDevPtr, value, extent); + return ihipMemset3D(pitchedDevPtr, value, extent, stream, true); +} + hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) { HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, extent, stream); + HIP_RETURN(hipMemset3DAsync_common(pitchedDevPtr, value, extent, stream)); +} - STREAM_CAPTURE(hipMemset3DAsync, stream, pitchedDevPtr, value, extent); - - HIP_RETURN(ihipMemset3D(pitchedDevPtr, value, extent, stream, true)); +hipError_t hipMemset3DAsync_spt(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) { + HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, extent, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipMemset3DAsync_common(pitchedDevPtr, value, extent, stream)); } hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, @@ -3236,16 +3317,24 @@ hipError_t hipMemcpy2DFromArray_spt(void* dst, size_t dpitch, hipArray_const_t s HIP_RETURN_DURATION(hipMemcpy2DFromArray_common(dst, dpitch, src, wOffsetSrc, hOffset, width, height, kind, stream)); } -hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream); +hipError_t hipMemcpy2DFromArrayAsync_common(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { STREAM_CAPTURE(hipMemcpy2DFromArrayAsync, stream, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind); - if (dpitch == 0) { - HIP_RETURN(hipErrorInvalidPitchValue); + return hipErrorInvalidPitchValue; } + return ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true); +} - HIP_RETURN_DURATION(ihipMemcpy2DFromArray(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true)); +hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream); + HIP_RETURN_DURATION(hipMemcpy2DFromArrayAsync_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream)); +} + +hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { + HIP_INIT_API(hipMemcpy2DFromArrayAsync, dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpy2DFromArrayAsync_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream)); } hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, hipMemcpyKind kind, hipStream_t stream) { @@ -3264,15 +3353,24 @@ hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffs 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); +hipError_t hipMemcpy2DToArrayAsync_common(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) { STREAM_CAPTURE(hipMemcpy2DToArrayAsync, stream, dst, wOffset, hOffset, src, spitch, width, height, kind); if (spitch == 0) { - HIP_RETURN(hipErrorInvalidPitchValue); + return hipErrorInvalidPitchValue; } + return ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true); +} - HIP_RETURN_DURATION(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, 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); + HIP_RETURN_DURATION(hipMemcpy2DToArrayAsync_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream)); +} + +hipError_t hipMemcpy2DToArrayAsync_spt(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); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN_DURATION(hipMemcpy2DToArrayAsync_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream)); } hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind, hipStream_t stream) { diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 53ba7a1b95..7c915d6979 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -574,16 +574,15 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) { } // ================================================================================================ -hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, +hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { - HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); //flags - Reserved for future use, must be 0 if (callback == nullptr || flags != 0) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); + return hipErrorContextIsDestroyed; } amd::HostQueue* hostQueue = hip::getQueue(stream); @@ -594,7 +593,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback } amd::Command* command = new amd::Marker(*hostQueue, !kMarkerDisableFlush, eventWaitList); if (command == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } StreamCallback* cbo = new StreamCallback(stream, callback, userData, command); @@ -617,14 +616,29 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback eventWaitList.push_back(command); amd::Command* block_command = new amd::Marker(*hostQueue, !kMarkerDisableFlush, eventWaitList); if (block_command == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } block_command->enqueue(); block_command->release(); // Release the callback marker command->release(); - HIP_RETURN(hipSuccess); + return hipSuccess; +} + +// ================================================================================================ +hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, + unsigned int flags) { + HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); + HIP_RETURN(hipStreamAddCallback_common(stream, callback, userData, flags)); +} + +// ================================================================================================ +hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, void* userData, + unsigned int flags) { + HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipStreamAddCallback_common(stream, callback, userData, flags)); } // ================================================================================================