diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h index 5cd74f687e..79bc84a84c 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_runtime_api.h @@ -3886,6 +3886,130 @@ const char* hipKernelNameRef(const hipFunction_t f); const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream); int hipGetStreamDeviceId(hipStream_t stream); +#ifdef __cplusplus +/** + * An opaque value that represents a hip graph + */ +class hipGraph; +typedef hipGraph* hipGraph_t; + +/** + * An opaque value that represents a hip graph node + */ +class hipGraphNode; +typedef hipGraphNode* hipGraphNode_t; + +/** + * An opaque value that represents a hip graph Exec + */ +class hipGraphExec; +typedef hipGraphExec* hipGraphExec_t; +typedef enum hipGraphNodeType { + hipGraphNodeTypeKernel = 1, ///< GPU kernel node + hipGraphNodeTypeMemcpy = 2, ///< Memcpy 3D node + hipGraphNodeTypeMemset = 3, ///< Memset 1D node + hipGraphNodeTypeHost = 4, ///< Host (executable) node + hipGraphNodeTypeGraph = 5, ///< Node which executes an embedded graph + hipGraphNodeTypeEmpty = 6, ///< Empty (no-op) node + hipGraphNodeTypeWaitEvent = 7, ///< External event wait node + hipGraphNodeTypeEventRecord = 8, ///< External event record node + hipGraphNodeTypeMemcpy1D = 9, ///< Memcpy 1D node + hipGraphNodeTypeMemcpyFromSymbol = 10, ///< MemcpyFromSymbol node + hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node + hipGraphNodeTypeCount +} hipGraphNodeType; + +typedef void (*hipHostFn_t)(void* userData); +typedef struct hipHostNodeParams { + hipHostFn_t fn; + void* userData; +} hipHostNodeParams; + +typedef struct hipKernelNodeParams { + dim3 blockDim; + void** extra; + void* func; + dim3 gridDim; + void** kernelParams; + unsigned int sharedMemBytes; +} hipKernelNodeParams; + +typedef struct hipMemsetParams { + void* dst; + unsigned int elementSize; + size_t height; + size_t pitch; + unsigned int value; + size_t width; +} hipMemsetParams; + +enum hipGraphExecUpdateResult { + hipGraphExecUpdateSuccess = 0x0, ///< The update succeeded + hipGraphExecUpdateError = 0x1, ///< The update failed for an unexpected reason which is described + ///< in the return value of the function + hipGraphExecUpdateErrorTopologyChanged = 0x2, ///< The update failed because the topology changed + hipGraphExecUpdateErrorNodeTypeChanged = 0x3, ///< The update failed because a node type changed + hipGraphExecUpdateErrorFunctionChanged = + 0x4, ///< The update failed because the function of a kernel node changed + hipGraphExecUpdateErrorParametersChanged = + 0x5, ///< The update failed because the parameters changed in a way that is not supported + hipGraphExecUpdateErrorNotSupported = + 0x6, ///< The update failed because something about the node is not supported + hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7 +}; + +enum hipStreamCaptureMode { + hipStreamCaptureModeGlobal = 0, + hipStreamCaptureModeThreadLocal, + hipStreamCaptureModeRelaxed +}; + +enum hipStreamCaptureStatus { + hipStreamCaptureStatusNone = 0, ///< Stream is not capturing + hipStreamCaptureStatusActive, ///< Stream is actively capturing + hipStreamCaptureStatusInvalidated ///< Stream is part of a capture sequence that has been + ///< invalidated, but not terminated +}; + +hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode); + +hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph); + +// Creates a graph. +hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags); + +// Destroys a graph. +hipError_t hipGraphDestroy(hipGraph_t graph); + +// Destroys an executable graph. +hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec); + +// Creates an executable graph from a graph. +hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, + hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize); + +// Launches an executable graph in a stream. +hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream); + +// Creates a kernel execution node and adds it to a graph. +hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipKernelNodeParams* pNodeParams); + +// Creates a memcpy node and adds it to a graph. +hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemcpy3DParms* pCopyParams); + +// Creates a memset node and adds it to a graph. +hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemsetParams* pMemsetParams); +#endif +// doxygen end graph API +/** + * @} + */ #ifdef __cplusplus } /* extern "c" */ #endif @@ -4085,6 +4209,11 @@ static inline hipError_t hipUnbindTexture( return hipUnbindTexture(&tex); } +// doxygen end Texture +/** + * @} + */ + #endif // __cplusplus #ifdef __GNUC__ diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index d00cfaf54f..7cd26bc1e4 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -280,6 +280,28 @@ typedef enum __HIP_NODISCARD hipError_t { ///< that was launched via cooperative launch APIs exceeds the maximum number of ///< allowed blocks for the current device hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented + hipErrorStreamCaptureUnsupported = 900, ///< The operation is not permitted when the stream + ///< is capturing. + hipErrorStreamCaptureInvalidated = 901, ///< The current capture sequence on the stream + ///< has been invalidated due to a previous error. + hipErrorStreamCaptureMerge = 902, ///< The operation would have resulted in a merge of + ///< two independent capture sequences. + hipErrorStreamCaptureUnmatched = 903, ///< The capture was not initiated in this stream. + hipErrorStreamCaptureUnjoined = 904, ///< The capture sequence contains a fork that was not + ///< joined to the primary stream. + hipErrorStreamCaptureIsolation = 905, ///< A dependency would have been created which crosses + ///< the capture sequence boundary. Only implicit + ///< in-stream ordering dependencies are allowed + ///< to cross the boundary + hipErrorStreamCaptureImplicit = 906, ///< The operation would have resulted in a disallowed + ///< implicit dependency on a current capture sequence + ///< from hipStreamLegacy. + hipErrorCapturedEvent = 907, ///< The operation is not permitted on an event which was last + ///< recorded in a capturing stream. + hipErrorStreamCaptureWrongThread = 908, ///< A stream capture sequence not initiated with + ///< the hipStreamCaptureModeRelaxed argument to + ///< hipStreamBeginCapture was passed to + ///< hipStreamEndCapture in a different thread. hipErrorUnknown = 999, //< Unknown error. // HSA Runtime Error Codes start here. hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen diff --git a/projects/clr/hipamd/rocclr/CMakeLists.txt b/projects/clr/hipamd/rocclr/CMakeLists.txt index 28177b7c75..63229359d0 100755 --- a/projects/clr/hipamd/rocclr/CMakeLists.txt +++ b/projects/clr/hipamd/rocclr/CMakeLists.txt @@ -142,6 +142,8 @@ add_library(hip64 OBJECT cl_gl.cpp cl_lqdflash_amd.cpp fixme.cpp + hip_graph.cpp + hip_graph_internal.cpp ) set_target_properties(hip64 PROPERTIES POSITION_INDEPENDENT_CODE ON) diff --git a/projects/clr/hipamd/rocclr/hip_event.cpp b/projects/clr/hipamd/rocclr/hip_event.cpp index e5043bc763..05143bbda2 100755 --- a/projects/clr/hipamd/rocclr/hip_event.cpp +++ b/projects/clr/hipamd/rocclr/hip_event.cpp @@ -321,6 +321,8 @@ bool createIpcEventShmemIfNeeded(hip::Event::ihipIpcEvent_t& ipc_evt) { hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(hipEventRecord, event, stream); + STREAM_CAPTURE(hipEventRecord, stream, event); + if (event == nullptr) { HIP_RETURN(hipErrorInvalidHandle); } diff --git a/projects/clr/hipamd/rocclr/hip_event.hpp b/projects/clr/hipamd/rocclr/hip_event.hpp index e63ee84be8..94b71c3fcc 100644 --- a/projects/clr/hipamd/rocclr/hip_event.hpp +++ b/projects/clr/hipamd/rocclr/hip_event.hpp @@ -64,10 +64,18 @@ typedef struct ihipIpcEventShmem_s { } ihipIpcEventShmem_t; class Event { -public: + /// event recorded on stream where capture is active + bool onCapture_; + /// capture stream where event is recorded + hipStream_t captureStream_; + /// Previous captured nodes before event record + std::vector nodesPrevToRecorded_; + + public: Event(unsigned int flags) : flags(flags), lock_("hipEvent_t", true), event_(nullptr), recorded_(false) { // No need to init event_ here as addMarker does that + onCapture_ = false; device_id_ = hip::getCurrentDevice()->deviceId(); // Created in current device ctx } @@ -89,7 +97,30 @@ public: const int deviceId() { return device_id_; } void setDeviceId(int id) { device_id_ = id; } - //IPC Events + /// End capture on this event + void EndCapture() { + onCapture_ = false; + captureStream_ = nullptr; + } + /// Start capture when waited on this event + void StartCapture(hipStream_t stream) { + onCapture_ = true; + captureStream_ = stream; + } + /// Get capture status of the graph + bool GetCaptureStatus() { return onCapture_; } + /// Get capture stream where event is recorded + hipStream_t GetCaptureStream() { return captureStream_; } + /// Set capture stream where event is recorded + void SetCaptureStream(hipStream_t stream) { captureStream_ = stream; } + /// Returns previous captured nodes before event record + std::vector GetNodesPrevToRecorded() const { return nodesPrevToRecorded_; } + /// Set last captured graph node before event record + void SetNodesPrevToRecorded(std::vector& graphNode) { + nodesPrevToRecorded_ = graphNode; + } + + // IPC Events struct ihipIpcEvent_t { std::string ipc_name_; int ipc_fd_; diff --git a/projects/clr/hipamd/rocclr/hip_graph.cpp b/projects/clr/hipamd/rocclr/hip_graph.cpp new file mode 100644 index 0000000000..a8f5df5100 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_graph.cpp @@ -0,0 +1,415 @@ +/* Copyright (c) 2021-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_graph_internal.hpp" +#include "platform/command.hpp" +#include "hip_conversions.hpp" +#include "hip_platform.hpp" +#include "hip_event.hpp" + +thread_local std::vector g_captureStreams; +std::unordered_map hipGraphExec::activeGraphExec_; + +hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipKernelNodeParams* pNodeParams) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) { + return hipErrorInvalidValue; + } + hipFunction_t func = nullptr; + hipError_t status = + PlatformState::instance().getStatFunc(&func, pNodeParams->func, ihipGetDevice()); + if ((status != hipSuccess) || (func == nullptr)) { + return hipErrorInvalidDeviceFunction; + } + size_t globalWorkSizeX = static_cast(pNodeParams->gridDim.x) * pNodeParams->blockDim.x; + size_t globalWorkSizeY = static_cast(pNodeParams->gridDim.y) * pNodeParams->blockDim.y; + size_t globalWorkSizeZ = static_cast(pNodeParams->gridDim.z) * pNodeParams->blockDim.z; + if (globalWorkSizeX > std::numeric_limits::max() || + globalWorkSizeY > std::numeric_limits::max() || + globalWorkSizeZ > std::numeric_limits::max()) { + return hipErrorInvalidConfiguration; + } + status = ihipLaunchKernel_validate( + func, static_cast(globalWorkSizeX), static_cast(globalWorkSizeY), + static_cast(globalWorkSizeZ), pNodeParams->blockDim.x, pNodeParams->blockDim.y, + pNodeParams->blockDim.z, pNodeParams->sharedMemBytes, pNodeParams->kernelParams, + pNodeParams->extra, ihipGetDevice(), 0); + if (status != hipSuccess) { + return status; + } + *pGraphNode = new hipGraphKernelNode(pNodeParams, func); + if (numDependencies == 0) { + graph->AddNode(*pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + return hipSuccess; +} + +hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemcpy3DParms* pCopyParams) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) { + return hipErrorInvalidValue; + } + ihipMemcpy3D_validate(pCopyParams); + *pGraphNode = new hipGraphMemcpyNode(pCopyParams); + if (numDependencies == 0) { + graph->AddNode(*pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + return hipSuccess; +} + +hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemsetParams* pMemsetParams) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || pMemsetParams == nullptr) { + return hipErrorInvalidValue; + } + if (pMemsetParams->height == 1) { + ihipMemset_validate(pMemsetParams->dst, pMemsetParams->value, pMemsetParams->elementSize, + pMemsetParams->width * pMemsetParams->elementSize); + } else { + auto sizeBytes = pMemsetParams->width * pMemsetParams->height * 1; + ihipMemset3D_validate( + {pMemsetParams->dst, pMemsetParams->pitch, pMemsetParams->width, pMemsetParams->height}, + pMemsetParams->value, {pMemsetParams->width, pMemsetParams->height, 1}, sizeBytes); + } + + *pGraphNode = new hipGraphMemsetNode(pMemsetParams); + if (numDependencies == 0) { + graph->AddNode(*pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + return hipSuccess; +} + +hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim, + dim3& blockDim, void**& args, size_t& sharedMemBytes) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node kernel launch on stream : %p", stream); + hip::Stream* s = reinterpret_cast(stream); + hipKernelNodeParams nodeParams; + nodeParams.func = const_cast(hostFunction); + nodeParams.blockDim = blockDim; + nodeParams.extra = nullptr; + nodeParams.gridDim = gridDim; + nodeParams.kernelParams = args; + nodeParams.sharedMemBytes = sharedMemBytes; + + hipGraphNode_t pGraphNode; + hipError_t status = + ihipGraphAddKernelNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &nodeParams); + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + return hipSuccess; +} + +hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& p) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy3D on stream : %p", + stream); + hip::Stream* s = reinterpret_cast(stream); + hipGraphNode_t pGraphNode; + 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 capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& src, + size_t& sizeBytes, hipMemcpyKind& kind) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memcpy1D on stream : %p", + stream); + hip::Stream* s = reinterpret_cast(stream); + hipGraphNode_t pGraphNode; + hipGraph_t graph = nullptr; + std::vector pDependencies = s->GetLastCapturedNodes(); + size_t numDependencies = s->GetLastCapturedNodes().size(); + graph = s->GetCaptureGraph(); + ihipMemcpy_validate(dst, src, sizeBytes, kind); + pGraphNode = new hipGraphMemcpyNode1D(dst, src, sizeBytes, kind); + if (numDependencies == 0) { + graph->AddNode(pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(pDependencies[i], pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + s->SetLastCapturedNode(pGraphNode); + return hipSuccess; +} + +hipError_t capturehipMemcpyFromSymbolAsync(hipStream_t& stream, void*& dst, const void*& symbol, + size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node MemcpyFromSymbolNode on stream : %p", stream); + return hipSuccess; +} + +hipError_t capturehipMemcpyToSymbolAsync(hipStream_t& stream, const void*& symbol, const void*& src, + size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node MemcpyToSymbolNode on stream : %p", stream); + return hipSuccess; +} + +hipError_t capturehipMemsetAsync(hipStream_t& stream, void*& dst, int& value, size_t& valueSize, + size_t& sizeBytes) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset1D on stream : %p", + stream); + + hipMemsetParams memsetParams = {0}; + memsetParams.dst = dst; + memsetParams.value = value; + memsetParams.elementSize = valueSize; + memsetParams.width = sizeBytes / valueSize; + memsetParams.height = 1; + + hip::Stream* s = reinterpret_cast(stream); + hipGraphNode_t pGraphNode; + hipError_t status = + ihipGraphAddMemsetNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &memsetParams); + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + return hipSuccess; +} + +hipError_t capturehipMemset2DAsync(hipStream_t& stream, void*& dst, size_t& pitch, int& value, + size_t& width, size_t& height) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset2D on stream : %p", + stream); + hipMemsetParams memsetParams = {0}; + + memsetParams.dst = dst; + memsetParams.value = value; + memsetParams.width = width; + memsetParams.height = height; + memsetParams.pitch = pitch; + hip::Stream* s = reinterpret_cast(stream); + hipGraphNode_t pGraphNode; + hipError_t status = + ihipGraphAddMemsetNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size(), &memsetParams); + if (status != hipSuccess) { + return status; + } + s->SetLastCapturedNode(pGraphNode); + return hipSuccess; +} + +hipError_t capturehipMemset3DAsync(hipStream_t& stream, hipPitchedPtr& pitchedDevPtr, int& value, + hipExtent& extent) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset3D on stream : %p", + stream); + return hipSuccess; +} + +hipError_t capturehipEventRecord(hipStream_t& stream, hipEvent_t& event) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node EventRecord on stream : %p, Event %p", stream, event); + if (event == nullptr) { + HIP_RETURN(hipErrorInvalidHandle); + } + hip::Event* e = reinterpret_cast(event); + e->StartCapture(stream); + hip::Stream* s = reinterpret_cast(stream); + std::vector lastCapturedNodes = s->GetLastCapturedNodes(); + if (!lastCapturedNodes.empty()) { + e->SetNodesPrevToRecorded(lastCapturedNodes); + } + return hipSuccess; +} + +hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, unsigned int& flags) { + ClPrint(amd::LOG_INFO, amd::LOG_API, + "[hipGraph] current capture node StreamWaitEvent on stream : %p, Event %p", stream, + event); + hip::Stream* s = reinterpret_cast(stream); + hip::Event* e = reinterpret_cast(event); + + if (event == nullptr || stream == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + if (!s->IsOriginStream()) { + s->SetCaptureGraph(reinterpret_cast(e->GetCaptureStream())->GetCaptureGraph()); + s->SetCaptureMode(reinterpret_cast(e->GetCaptureStream())->GetCaptureMode()); + s->SetParentStream(e->GetCaptureStream()); + } + s->AddCrossCapturedNode(e->GetNodesPrevToRecorded()); + g_captureStreams.push_back(stream); + return hipSuccess; +} + +hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus** pCaptureStatus) { + HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); + if (stream == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hipStreamCaptureStatus captureStatus = reinterpret_cast(stream)->GetCaptureStatus(); + *pCaptureStatus = &captureStatus; + HIP_RETURN(hipSuccess); +} + +hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) { + HIP_INIT_API(hipStreamBeginCapture, stream, mode); + hip::Stream* s = reinterpret_cast(stream); + // capture cannot be initiated on legacy stream + // It can be initiated if the stream is not already in capture mode + if (stream == nullptr || s->GetCaptureStatus() == hipStreamCaptureStatusActive) { + HIP_RETURN(hipErrorInvalidValue); + } + s->SetCaptureGraph(new hipGraph()); + s->SetCaptureMode(mode); + s->SetOriginStream(); + g_captureStreams.push_back(stream); + HIP_RETURN_DURATION(hipSuccess); +} + +hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph) { + HIP_INIT_API(hipStreamEndCapture, stream, pGraph); + hip::Stream* s = reinterpret_cast(stream); + // Capture must be ended on the same stream in which it was initiated + if (!s->IsOriginStream()) { + HIP_RETURN(hipErrorStreamCaptureUnmatched); + } + // If mode is not hipStreamCaptureModeRelaxed, hipStreamEndCapture must be called on the stream + // from the same thread + if (s->GetCaptureMode() != hipStreamCaptureModeRelaxed && + std::find(g_captureStreams.begin(), g_captureStreams.end(), stream) == + g_captureStreams.end()) { + HIP_RETURN(hipErrorStreamCaptureWrongThread); + } + // If capture was invalidated, due to a violation of the rules of stream capture + if (s->GetCaptureStatus() == hipStreamCaptureStatusInvalidated) { + *pGraph = nullptr; + HIP_RETURN(hipErrorStreamCaptureInvalidated); + } + // check if all parallel streams have joined + if (s->GetCaptureGraph()->GetLeafNodeCount() != 1) { + return hipErrorStreamCaptureUnjoined; + } + *pGraph = s->GetCaptureGraph(); + // end capture on all streams/events part of graph capture + HIP_RETURN_DURATION(s->EndCapture()); +} + +hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags) { + HIP_INIT_API(hipGraphCreate, pGraph, flags); + *pGraph = new hipGraph(); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphDestroy(hipGraph_t graph) { + HIP_INIT_API(hipGraphDestroy, graph); + delete graph; + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipKernelNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphAddKernelNode, pGraphNode, graph, pDependencies, numDependencies, + pNodeParams); + HIP_RETURN_DURATION( + ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams);); +} + +hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemcpy3DParms* pCopyParams) { + HIP_INIT_API(hipGraphAddMemcpyNode, pGraphNode, graph, pDependencies, numDependencies, + pCopyParams); + + HIP_RETURN_DURATION( + ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams);); +} + +hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipMemsetParams* pMemsetParams) { + HIP_INIT_API(hipGraphAddMemsetNode, pGraphNode, graph, pDependencies, numDependencies, + pMemsetParams); + + HIP_RETURN_DURATION( + ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams);); +} + +hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, + hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize) { + std::vector> parallelLists; + std::unordered_map> nodeWaitLists; + graph->GetRunList(parallelLists, nodeWaitLists); + std::vector levelOrder; + graph->LevelOrder(levelOrder); + *pGraphExec = new hipGraphExec(levelOrder, parallelLists, nodeWaitLists); + if (*pGraphExec != nullptr) { + return (*pGraphExec)->Init(); + } else { + return hipErrorOutOfMemory; + } +} + +hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, + hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize) { + HIP_INIT_API(hipGraphInstantiate, pGraphExec, graph); + HIP_RETURN_DURATION(ihipGraphInstantiate(pGraphExec, graph, pErrorNode, pLogBuffer, bufferSize)); +} + +hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) { + HIP_INIT_API(hipGraphExecDestroy, pGraphExec); + delete pGraphExec; + HIP_RETURN(hipSuccess); +} + +hipError_t ihipGraphlaunch(hipGraphExec_t graphExec, hipStream_t stream) { + return graphExec->Run(stream); +} + +hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { + HIP_INIT_API(hipGraphLaunch, graphExec, stream); + HIP_RETURN_DURATION(ihipGraphlaunch(graphExec, stream)); +} diff --git a/projects/clr/hipamd/rocclr/hip_graph_capture.hpp b/projects/clr/hipamd/rocclr/hip_graph_capture.hpp new file mode 100644 index 0000000000..16abfdeef7 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_graph_capture.hpp @@ -0,0 +1,48 @@ +/* Copyright (c) 2021-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. */ + +#pragma once +// forward declaration of capture methods +hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim, + dim3& blockDim, void**& args, size_t& sharedMemBytes); + +hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& p); + +hipError_t capturehipMemcpyAsync(hipStream_t& stream, void*& dst, const void*& src, + size_t& sizeBytes, hipMemcpyKind& kind); + +hipError_t capturehipMemcpyFromSymbolAsync(hipStream_t& stream, void*& dst, const void*& symbol, + size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind); + +hipError_t capturehipMemcpyToSymbolAsync(hipStream_t& stream, const void*& symbol, const void*& src, + size_t& sizeBytes, size_t& offset, hipMemcpyKind& kind); + +hipError_t capturehipMemsetAsync(hipStream_t& stream, void*& dst, int& value, size_t& valueSize, + size_t& sizeBytes); + +hipError_t capturehipMemset2DAsync(hipStream_t& stream, void*& dst, size_t& pitch, int& value, + size_t& width, size_t& height); + +hipError_t capturehipMemset3DAsync(hipStream_t& stream, hipPitchedPtr& pitchedDevPtr, int& value, + hipExtent& extent); + +hipError_t capturehipEventRecord(hipStream_t& stream, hipEvent_t& event); + +hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, unsigned int& flags); diff --git a/projects/clr/hipamd/rocclr/hip_graph_helper.hpp b/projects/clr/hipamd/rocclr/hip_graph_helper.hpp new file mode 100644 index 0000000000..b28c14e811 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_graph_helper.hpp @@ -0,0 +1,35 @@ +hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p); + +hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); + +hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes, + hipMemcpyKind kind, amd::HostQueue& queue); + +hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, + uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, + uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, + uint32_t sharedMemBytes, void** kernelParams, void** extra, + int deviceId, uint32_t params); + +hipError_t ihipMemset_validate(void* dst, int64_t value, size_t valueSize, size_t sizeBytes); + +hipError_t ihipMemset3D_validate(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, + size_t sizeBytes); + +hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f, + uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, + uint32_t globalWorkSizeZ, uint32_t blockDimX, uint32_t blockDimY, + uint32_t blockDimZ, uint32_t sharedMemBytes, + amd::HostQueue* queue, void** kernelParams, void** extra, + hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, + uint32_t params, uint32_t gridId, uint32_t numGrids, + uint64_t prevGridSum, uint64_t allGridSum, uint32_t firstDevice); + +hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p, + amd::HostQueue* queue); + +hipError_t ihipMemsetCommand(std::vector& commands, void* dst, int64_t value, + size_t valueSize, size_t sizeBytes, amd::HostQueue* queue); + +hipError_t ihipMemset3DCommand(std::vector& commands, hipPitchedPtr pitchedDevPtr, + int value, hipExtent extent, amd::HostQueue* queue); diff --git a/projects/clr/hipamd/rocclr/hip_graph_internal.cpp b/projects/clr/hipamd/rocclr/hip_graph_internal.cpp new file mode 100644 index 0000000000..75eb6f6726 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_graph_internal.cpp @@ -0,0 +1,364 @@ +/* Copyright (c) 2021-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_graph_internal.hpp" +#include + +#define CASE_STRING(X, C) \ + case X: \ + case_string = #C; \ + break; +const char* GetGraphNodeTypeString(uint32_t op) { + const char* case_string; + switch (static_cast(op)) { + CASE_STRING(hipGraphNodeTypeKernel, KernelNode) + CASE_STRING(hipGraphNodeTypeMemcpy, Memcpy3DNode) + CASE_STRING(hipGraphNodeTypeMemset, MemsetNode) + CASE_STRING(hipGraphNodeTypeHost, HostNode) + CASE_STRING(hipGraphNodeTypeGraph, GraphNode) + CASE_STRING(hipGraphNodeTypeEmpty, EmptyNode) + CASE_STRING(hipGraphNodeTypeWaitEvent, WaitEventNode) + CASE_STRING(hipGraphNodeTypeEventRecord, EventRecordNode) + CASE_STRING(hipGraphNodeTypeMemcpy1D, Memcpy1DNode) + CASE_STRING(hipGraphNodeTypeMemcpyFromSymbol, MemcpyFromSymbolNode) + CASE_STRING(hipGraphNodeTypeMemcpyToSymbol, MemcpyToSymbolNode) + default: + case_string = "Unknown node type"; + }; + return case_string; +}; + +hipError_t hipGraph::AddNode(const Node& node) { + vertices_.emplace_back(node); + nodeOutDegree_[node] = 0; + nodeInDegree_[node] = 0; + node->SetLevel(0); + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n", + GetGraphNodeTypeString(node->GetType()), node); + return hipSuccess; +} + +hipError_t hipGraph::AddEdge(const Node& parentNode, const Node& childNode) { + // if vertice doesn't exist, add it to the graph + if (std::find(vertices_.begin(), vertices_.end(), parentNode) == vertices_.end()) { + AddNode(parentNode); + } + if (std::find(vertices_.begin(), vertices_.end(), childNode) == vertices_.end()) { + AddNode(childNode); + } + // Check if edge already exists + auto connectedEdges = edges_.find(parentNode); + if (connectedEdges != edges_.end()) { + if (std::find(connectedEdges->second.begin(), connectedEdges->second.end(), childNode) != + connectedEdges->second.end()) { + return hipSuccess; + } + connectedEdges->second.emplace_back(childNode); + } else { + edges_[parentNode] = {childNode}; + } + nodeOutDegree_[parentNode]++; + nodeInDegree_[childNode]++; + childNode->SetLevel(std::max(childNode->GetLevel(), parentNode->GetLevel() + 1)); + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add edge btwn %s(%p) - %s(%p)\n", + GetGraphNodeTypeString(parentNode->GetType()), parentNode, + GetGraphNodeTypeString(childNode->GetType()), childNode); + return hipSuccess; +} + +// root nodes are all vertices with 0 in-degrees +std::vector hipGraph::GetRootNodes() const { + std::vector roots; + for (auto entry : vertices_) { + if (nodeInDegree_.at(entry) == 0) { + roots.push_back(entry); + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] root node: %s(%p)\n", + GetGraphNodeTypeString(entry->GetType()), entry); + } + } + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "\n"); + return roots; +} + +// leaf nodes are all vertices with 0 out-degrees +std::vector hipGraph::GetLeafNodes() const { + std::vector leafNodes; + for (auto entry : vertices_) { + if (nodeOutDegree_.at(entry) == 0) { + leafNodes.push_back(entry); + } + } + return leafNodes; +} + +size_t hipGraph::GetLeafNodeCount() const { + int numLeafNodes = 0; + for (auto entry : vertices_) { + if (nodeOutDegree_.at(entry) == 0) { + numLeafNodes++; + } + } + return numLeafNodes; +} + +std::vector> hipGraph::GetEdges() const { + std::vector> edges; + for (const auto& i : edges_) { + for (const auto& j : i.second) { + edges.push_back(std::make_pair(i.first, j)); + } + } + return edges; +} + +void hipGraph::GetRunListUtil(Node v, std::unordered_map& visited, + std::vector& singleList, + std::vector>& parallelLists, + std::unordered_map>& dependencies) { + // Mark the current node as visited. + visited[v] = true; + singleList.push_back(v); + // Recurse for all the vertices adjacent to this vertex + for (auto& adjNode : edges_[v]) { + if (!visited[adjNode]) { + // For the parallel list nodes add parent as the dependency + if (singleList.empty()) { + ClPrint(amd::LOG_INFO, amd::LOG_CODE, + "[hipGraph] For %s(%p)- add parent as dependency %s(%p)\n", + GetGraphNodeTypeString(adjNode->GetType()), adjNode, + GetGraphNodeTypeString(v->GetType()), v); + dependencies[adjNode].push_back(v); + } + GetRunListUtil(adjNode, visited, singleList, parallelLists, dependencies); + } else { + for (auto& list : parallelLists) { + // Merge singleList when adjNode matches with the first element of the list in existing + // lists + if (adjNode == list[0]) { + for (auto k = singleList.rbegin(); k != singleList.rend(); ++k) { + list.insert(list.begin(), *k); + } + singleList.erase(singleList.begin(), singleList.end()); + } + } + // If the list cannot be merged with the existing list add as dependancy + if (!singleList.empty()) { + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] For %s(%p)- add dependency %s(%p)\n", + GetGraphNodeTypeString(adjNode->GetType()), adjNode, + GetGraphNodeTypeString(v->GetType()), v); + dependencies[adjNode].push_back(v); + } + } + } + if (!singleList.empty()) { + parallelLists.push_back(singleList); + singleList.erase(singleList.begin(), singleList.end()); + } +} +// The function to do Topological Sort. +// It uses recursive GetRunListUtil() +void hipGraph::GetRunList(std::vector>& parallelLists, + std::unordered_map>& dependencies) { + std::vector singleList; + + // Mark all the vertices as not visited + std::unordered_map visited; + for (auto node : vertices_) visited[node] = false; + + // Call the recursive helper function for all vertices one by one + for (auto node : vertices_) { + if (visited[node] == false) { + GetRunListUtil(node, visited, singleList, parallelLists, dependencies); + } + } + for (size_t i = 0; i < parallelLists.size(); i++) { + for (size_t j = 0; j < parallelLists[i].size(); j++) { + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] list %d - %s(%p)\n", i + 1, + GetGraphNodeTypeString(parallelLists[i][j]->GetType()), parallelLists[i][j]); + } + } +} + +hipError_t hipGraph::LevelOrder(std::vector& levelOrder) { + std::vector roots = GetRootNodes(); + std::unordered_map visited; + std::queue q; + for (auto it = roots.begin(); it != roots.end(); it++) { + q.push(*it); + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] %s(%p) level:%d \n", + GetGraphNodeTypeString((*it)->GetType()), *it, (*it)->GetLevel()); + } + while (!q.empty()) { + Node& node = q.front(); + q.pop(); + levelOrder.push_back(node); + for (const auto& i : edges_[node]) { + if (visited.find(i) == visited.end() && i->GetLevel() == (node->GetLevel() + 1)) { + q.push(i); + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] %s(%p) level:%d \n", + GetGraphNodeTypeString(i->GetType()), i, i->GetLevel()); + visited[i] = true; + } + } + } + return hipSuccess; +} + +hipError_t hipGraphExec::CreateQueues() { + parallelQueues_.reserve(parallelLists_.size()); + for (size_t i = 0; i < parallelLists_.size(); i++) { + amd::HostQueue* queue; + cl_command_queue_properties properties = + (callbacks_table.is_enabled() || HIP_FORCE_QUEUE_PROFILING) ? CL_QUEUE_PROFILING_ENABLE : 0; + queue = new amd::HostQueue(*hip::getCurrentDevice()->asContext(), + *hip::getCurrentDevice()->devices()[0], properties); + + bool result = (queue != nullptr) ? queue->create() : false; + // Create a host queue + if (result) { + parallelQueues_.push_back(queue); + } else { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to create host queue\n"); + return hipErrorOutOfMemory; + } + } + return hipSuccess; +} + +hipError_t hipGraphExec::FillCommands() { + // Create commands + int i = 0; + hipError_t status; + for (const auto& list : parallelLists_) { + for (auto& node : list) { + status = node->CreateCommand(parallelQueues_[i]); + if (status != hipSuccess) return status; + } + i++; + } + // Add waitlists for all the commands + for (auto& node : levelOrder_) { + auto nodeWaitList = nodeWaitLists_.find(node); + if (nodeWaitList != nodeWaitLists_.end()) { + amd::Command::EventWaitList waitList; + for (auto depNode : nodeWaitList->second) { + for (auto command : depNode->GetCommands()) { + waitList.push_back(command); + } + } + for (auto command : nodeWaitList->first->GetCommands()) { + command->updateEventWaitList(waitList); + } + } + } + return status; +} + + +hipError_t hipGraphExec::Init() { + hipError_t status; + status = CreateQueues(); + if (status != hipSuccess) { + return status; + } + status = FillCommands(); + if (status != hipSuccess) { + return status; + } + rootCommand_ = nullptr; + /// stream should execute next command after graph finishes + /// Add marker to the stream that waits for all the last commands in parallel queues of graph + for (auto& singleList : parallelLists_) { + graphLastCmdWaitList_.push_back(singleList.back()->GetCommands().back()); + } + return status; +} + +void hipGraphExec::ResetGraph(cl_event event, cl_int command_exec_status, void* user_data) { + ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Inside resetGraph!\n"); + hipGraphExec_t graphExec = + hipGraphExec::activeGraphExec_[reinterpret_cast(user_data)]; + if (graphExec != nullptr) { + for (auto& node : graphExec->levelOrder_) { + for (auto& command : node->GetCommands()) { + command->resetStatus(CL_INT_MAX); + } + } + graphExec->rootCommand_->resetStatus(CL_INT_MAX); + graphExec->bExecPending_.store(false); + } else { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] graphExec is nullptr during resetGraph!\n"); + } +} + +hipError_t hipGraphExec::UpdateGraphToWaitOnRoot() { + for (auto& singleList : parallelLists_) { + amd::Command::EventWaitList waitList; + waitList.push_back(rootCommand_); + if (!singleList.empty()) { + auto commands = singleList[0]->GetCommands(); + if (!commands.empty()) { + commands[0]->updateEventWaitList(waitList); + } + } + } + return hipSuccess; +} + +hipError_t hipGraphExec::Run(hipStream_t stream) { + if (bExecPending_.load() == true) { + ClPrint( + amd::LOG_INFO, amd::LOG_CODE, + "[hipGraph] Same graph launched while previous one is active, wait for it to finish!\n"); + lastEnqueuedGraphCmd_->awaitCompletion(); + } + amd::HostQueue* queue = hip::getQueue(stream); + if (queue == nullptr) { + return hipErrorInvalidResourceHandle; + } + if (rootCommand_ == nullptr || rootCommand_->queue() != queue) { + if (rootCommand_ != nullptr) { + rootCommand_->release(); + } + rootCommand_ = new amd::Marker(*queue, false, {}); + UpdateGraphToWaitOnRoot(); + } + rootCommand_->enqueue(); + for (auto& node : levelOrder_) { + for (auto& command : node->GetCommands()) { + command->enqueue(); + } + } + + amd::Command* command = new amd::Marker(*queue, false, graphLastCmdWaitList_); + if (command == nullptr) { + return hipErrorOutOfMemory; + } + amd::Event& event = command->event(); + if (!event.setCallback(CL_COMPLETE, hipGraphExec::ResetGraph, command)) { + return hipErrorInvalidHandle; + } + hipGraphExec::activeGraphExec_[command] = this; + lastEnqueuedGraphCmd_ = command; + bExecPending_.store(true); + command->enqueue(); + command->release(); + return hipSuccess; +} \ No newline at end of file diff --git a/projects/clr/hipamd/rocclr/hip_graph_internal.hpp b/projects/clr/hipamd/rocclr/hip_graph_internal.hpp new file mode 100644 index 0000000000..8c034176f7 --- /dev/null +++ b/projects/clr/hipamd/rocclr/hip_graph_internal.hpp @@ -0,0 +1,355 @@ +/* Copyright (c) 2021-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. */ + +#pragma once +#include +#include +#include +#include +#include +#include +#include + +#include "hip/hip_runtime.h" +#include "hip_internal.hpp" +#include "hip_graph_helper.hpp" + +typedef hipGraphNode* Node; + +class hipGraphNode { + protected: + uint32_t level_; + hipGraphNodeType type_; + std::vector commands_; + bool visited_; + + public: + hipGraphNode(hipGraphNodeType type) { + type_ = type; + level_ = 0; + visited_ = false; + } + virtual ~hipGraphNode() { + for (auto command : commands_) { + delete command; + } + } + virtual hipError_t CreateCommand(amd::HostQueue* queue) { return hipSuccess; } + std::vector& GetCommands() { return commands_; } + hipGraphNodeType GetType() { return type_; } + uint32_t GetLevel() { return level_; } + void SetLevel(uint32_t level) { level_ = level; } +}; + +class hipGraph { + std::unordered_map nodeInDegree_; // count of in coming edges for every vertex + std::unordered_map nodeOutDegree_; // count of outgoing edges for every vertex + std::vector vertices_; + std::unordered_map> edges_; + + public: + hipGraph() {} + ~hipGraph(){}; + /// add node to the graph + hipError_t AddNode(const Node& node); + /// add edge to the graph + hipError_t AddEdge(const Node& parentNode, const Node& childNode); + /// Returns root nodes, all vertices with 0 in-degrees + std::vector GetRootNodes() const; + /// Returns leaf nodes, all vertices with 0 out-degrees + std::vector GetLeafNodes() const; + /// Returns number of leaf nodes + size_t GetLeafNodeCount() const; + /// Returns total numbers of nodes in the graph + size_t GetNodeCount() const { return vertices_.size(); } + /// returns all the nodes in the graph + std::vector GetNodes() const { return vertices_; } + /// returns all the edges in the graph + std::vector> GetEdges() const; + void GetRunListUtil(Node v, std::unordered_map& visited, + std::vector& singleList, std::vector>& parallelList, + std::unordered_map>& dependencies); + void GetRunList(std::vector>& parallelList, + std::unordered_map>& dependencies); + hipError_t LevelOrder(std::vector& levelOrder); +}; + +class hipGraphKernelNode : public hipGraphNode { + hipKernelNodeParams* pKernelParams_; + hipFunction_t func_; + + public: + hipGraphKernelNode(const hipKernelNodeParams* pNodeParams, const hipFunction_t func) + : hipGraphNode(hipGraphNodeTypeKernel) { + pKernelParams_ = new hipKernelNodeParams(*pNodeParams); + func_ = func; + } + ~hipGraphKernelNode() { delete pKernelParams_; } + + hipError_t CreateCommand(amd::HostQueue* queue) { + commands_.reserve(1); + amd::Command* command; + hipError_t status = ihipLaunchKernelCommand( + command, func_, pKernelParams_->gridDim.x * pKernelParams_->blockDim.x, + pKernelParams_->gridDim.y * pKernelParams_->blockDim.y, + pKernelParams_->gridDim.z * pKernelParams_->blockDim.z, pKernelParams_->blockDim.x, + pKernelParams_->blockDim.y, pKernelParams_->blockDim.z, pKernelParams_->sharedMemBytes, + queue, pKernelParams_->kernelParams, pKernelParams_->extra, nullptr, nullptr, 0, 0, 0, 0, 0, + 0, 0); + commands_.emplace_back(command); + return status; + } + + void GetParams(hipKernelNodeParams* params) { + std::memcpy(params, pKernelParams_, sizeof(hipKernelNodeParams)); + } + void SetParams(hipKernelNodeParams* params) { + std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams)); + } +}; + +class hipGraphMemcpyNode : public hipGraphNode { + hipMemcpy3DParms* pCopyParams_; + + public: + hipGraphMemcpyNode(const hipMemcpy3DParms* pCopyParams) : hipGraphNode(hipGraphNodeTypeMemcpy) { + pCopyParams_ = new hipMemcpy3DParms(*pCopyParams); + } + ~hipGraphMemcpyNode() { delete pCopyParams_; } + + hipError_t CreateCommand(amd::HostQueue* queue) { + commands_.reserve(1); + amd::Command* command; + hipError_t status = ihipMemcpy3DCommand(command, pCopyParams_, queue); + commands_.emplace_back(command); + return status; + } + + void GetParams(hipMemcpy3DParms* params) { + std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms)); + } + void SetParams(hipMemcpy3DParms* params) { + std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms)); + } +}; + + +class hipGraphMemcpyNode1D : public hipGraphNode { + void* dst_; + const void* src_; + size_t count_; + hipMemcpyKind kind_; + + public: + hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind) + : hipGraphNode(hipGraphNodeTypeMemcpy1D), dst_(dst), src_(src), count_(count), kind_(kind) {} + ~hipGraphMemcpyNode1D() {} + + hipError_t CreateCommand(amd::HostQueue* queue) { + commands_.reserve(1); + amd::Command* command = nullptr; + hipError_t status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *queue); + commands_.emplace_back(command); + return status; + } + + void SetParams(void* dst, const void* src, size_t count, hipMemcpyKind kind) { + dst_ = dst; + src_ = src; + count_ = count; + kind_ = kind; + } +}; + +template class hipGraphMemcpyNodeFromSymbol : public hipGraphNode { + void* dst_; + const T& symbol_; + size_t count_; + size_t offset_; + hipMemcpyKind kind_; + + public: + hipGraphMemcpyNodeFromSymbol(void* dst, const void* symbol, size_t count, size_t offset, + hipMemcpyKind kind) + : hipGraphNode(hipGraphNodeTypeMemcpyFromSymbol), + dst_(dst), + symbol_(symbol), + count_(count), + offset_(offset), + kind_(kind) {} + ~hipGraphMemcpyNodeFromSymbol() {} + + hipError_t CreateCommand(amd::HostQueue* queue); + + void SetParams(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) { + dst_ = dst; + symbol_ = symbol; + count_ = count; + offset_ = offset; + kind_ = kind; + } +}; + +template class hipGraphMemcpyNodeToSymbol : public hipGraphNode { + const T& symbol_; + const void* src_; + size_t count_; + size_t offset_; + hipMemcpyKind kind_; + + public: + hipGraphMemcpyNodeToSymbol(const T& symbol, void* src, size_t count, size_t offset, + hipMemcpyKind kind) + : hipGraphNode(hipGraphNodeTypeMemcpyToSymbol), + symbol_(symbol), + src_(src), + count_(count), + offset_(offset), + kind_(kind) {} + ~hipGraphMemcpyNodeToSymbol() {} + + hipError_t CreateCommand(amd::HostQueue* queue); + + void SetParams(const T& symbol, void* src, size_t count, size_t offset, hipMemcpyKind kind) { + symbol_ = symbol; + src_ = src; + count_ = count; + offset_ = offset; + kind_ = kind; + } +}; + +class hipGraphMemsetNode : public hipGraphNode { + hipMemsetParams* pMemsetParams_; + + public: + hipGraphMemsetNode(const hipMemsetParams* pMemsetParams) : hipGraphNode(hipGraphNodeTypeMemset) { + pMemsetParams_ = new hipMemsetParams(*pMemsetParams); + } + ~hipGraphMemsetNode() { delete pMemsetParams_; } + + hipError_t CreateCommand(amd::HostQueue* queue) { + if (pMemsetParams_->height == 1) { + return ihipMemsetCommand(commands_, pMemsetParams_->dst, pMemsetParams_->value, + pMemsetParams_->elementSize, + pMemsetParams_->width * pMemsetParams_->elementSize, queue); + } else { + return ihipMemset3DCommand(commands_, + {pMemsetParams_->dst, pMemsetParams_->pitch, pMemsetParams_->width, + pMemsetParams_->height}, + pMemsetParams_->elementSize, + {pMemsetParams_->width, pMemsetParams_->height, 1}, queue); + } + return hipSuccess; + } + + void GetParams(hipMemsetParams* params) { + std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams)); + } + void SetParams(hipMemsetParams* params) { + std::memcpy(pMemsetParams_, params, sizeof(hipMemsetParams)); + } +}; + +class hipGraphEventRecordNode : public hipGraphNode { + hipEvent_t event_; + + public: + hipGraphEventRecordNode(hipEvent_t event) + : hipGraphNode(hipGraphNodeTypeEventRecord), event_(event) {} + ~hipGraphEventRecordNode() {} + + hipError_t CreateCommand(amd::HostQueue* queue); + + void GetParams(hipEvent_t* event) { *event = event_; } + void SetParams(hipEvent_t event) { event_ = event; } +}; + +class hipGraphEventWaitNode : public hipGraphNode { + hipEvent_t event_; + + public: + hipGraphEventWaitNode(hipEvent_t event) + : hipGraphNode(hipGraphNodeTypeWaitEvent), event_(event) {} + ~hipGraphEventWaitNode() {} + + hipError_t CreateCommand(amd::HostQueue* queue); + + void GetParams(hipEvent_t* event) { *event = event_; } + void SetParams(hipEvent_t event) { event_ = event; } +}; + +class hipGraphHostNode : public hipGraphNode { + hipHostNodeParams* pNodeParams_; + + public: + hipGraphHostNode(const hipHostNodeParams* pNodeParams) : hipGraphNode(hipGraphNodeTypeHost) { + pNodeParams_ = new hipHostNodeParams(*pNodeParams); + } + ~hipGraphHostNode() { delete pNodeParams_; } + + hipError_t CreateCommand(amd::HostQueue* queue); + + void GetParams(hipHostNodeParams* params) { + std::memcpy(params, pNodeParams_, sizeof(hipHostNodeParams)); + } + void SetParams(hipHostNodeParams* params) { + std::memcpy(pNodeParams_, params, sizeof(hipHostNodeParams)); + } +}; + +class hipGraphExec { + std::vector> parallelLists_; + std::vector levelOrder_; + std::unordered_map> nodeWaitLists_; + std::vector parallelQueues_; + static std::unordered_map activeGraphExec_; + amd::Command::EventWaitList graphLastCmdWaitList_; + amd::Command* lastEnqueuedGraphCmd_; + std::atomic bExecPending_; + amd::Command* rootCommand_; + + public: + hipGraphExec(std::vector& levelOrder, std::vector>& lists, + std::unordered_map>& nodeWaitLists) + : parallelLists_(lists), + levelOrder_(levelOrder), + nodeWaitLists_(nodeWaitLists), + lastEnqueuedGraphCmd_(nullptr), + rootCommand_(nullptr) { + bExecPending_.store(false); + } + + ~hipGraphExec() { + for (auto queue : parallelQueues_) { + queue->release(); + } + for (auto node : levelOrder_) { + delete node; + } + } + + hipError_t CreateQueues(); + hipError_t FillCommands(); + hipError_t Init(); + hipError_t UpdateGraphToWaitOnRoot(); + hipError_t Run(hipStream_t stream); + static void ResetGraph(cl_event event, cl_int command_exec_status, void* user_data); +}; diff --git a/projects/clr/hipamd/rocclr/hip_hcc.def.in b/projects/clr/hipamd/rocclr/hip_hcc.def.in index eb18971dda..fd150837d8 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.def.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.def.in @@ -275,4 +275,17 @@ __gnu_f2h_ieee hipExtStreamGetCUMask hipImportExternalMemory hipExternalMemoryGetMappedBuffer -hipDestroyExternalMemory \ No newline at end of file +hipDestroyExternalMemory +hipGraphCreate +hipGraphDestroy +hipGraphAddKernelNode +hipGraphAddMemsetNode +hipGraphAddMemcpyNode +hipGraphAddMemcpyNode1D +hipGraphAddHostNode +hipGraphInstantiate +hipGraphLaunch +hipStreamIsCapturing +hipStreamBeginCapture +hipStreamEndCapture +hipGraphExecDestroy diff --git a/projects/clr/hipamd/rocclr/hip_hcc.map.in b/projects/clr/hipamd/rocclr/hip_hcc.map.in index cf433b18f5..8c3ac57554 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.map.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.map.in @@ -270,6 +270,19 @@ global: hipImportExternalMemory; hipExternalMemoryGetMappedBuffer; hipDestroyExternalMemory; + hipGraphCreate; + hipGraphDestroy; + hipGraphAddKernelNode; + hipGraphAddMemsetNode; + hipGraphAddMemcpyNode; + hipGraphAddMemcpyNode1D; + hipGraphAddHostNode; + hipGraphInstantiate; + hipGraphLaunch; + hipStreamIsCapturing; + hipStreamBeginCapture; + hipStreamEndCapture; + hipGraphExecDestroy; extern "C++" { hip_impl::hipLaunchKernelGGLImpl*; hip_impl::demangle*; diff --git a/projects/clr/hipamd/rocclr/hip_internal.hpp b/projects/clr/hipamd/rocclr/hip_internal.hpp index f624d93eca..577039ded2 100755 --- a/projects/clr/hipamd/rocclr/hip_internal.hpp +++ b/projects/clr/hipamd/rocclr/hip_internal.hpp @@ -26,6 +26,9 @@ #include "trace_helper.h" #include "utils/debug.hpp" #include "hip_formatting.hpp" + +#include "hip_graph_capture.hpp" + #include #include #include @@ -123,6 +126,19 @@ typedef struct ihipIpcEventHandle_st { } \ } while (0); +#define STREAM_CAPTURE(name, stream, ...) \ + if (stream != nullptr && \ + reinterpret_cast(stream)->GetCaptureStatus() == \ + hipStreamCaptureStatusActive) { \ + hipError_t status = capture##name(stream, ##__VA_ARGS__); \ + HIP_RETURN(status); \ + } + +#define EVENT_CAPTURE(name, event, ...) \ + if (event != nullptr && reinterpret_cast(event)->GetCaptureStatus() == true) { \ + hipError_t status = capture##name(event, ##__VA_ARGS__); \ + HIP_RETURN(status); \ + } namespace hc { class accelerator; @@ -134,7 +150,8 @@ namespace hip { class Stream { public: - enum Priority : int {High = -1, Normal = 0, Low = 1}; + enum Priority : int { High = -1, Normal = 0, Low = 1 }; + private: amd::HostQueue* queue_; mutable amd::Monitor lock_; @@ -144,11 +161,31 @@ namespace hip { bool null_; const std::vector cuMask_; + /// Stream capture related parameters + + /// Current capture status of the stream + hipStreamCaptureStatus captureStatus_; + /// Graph that is constructed with capture + hipGraph_t pCaptureGraph_; + /// Based on mode stream capture places restrictions on API calls that can be made within or + /// concurrently + hipStreamCaptureMode captureMode_; + bool originStream_; + /// Origin sream has no parent. Parent stream for the derived captured streams with event + /// dependencies + hipStream_t parentStream_; + /// Last graph node captured in the stream + std::vector lastCapturedNodes_; + /// Derived streams/Paralell branches from the origin stream + std::vector parallelCaptureStreams_; + /// Capture events + std::vector captureEvents_; + public: Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false, - const std::vector& cuMask = {}); + const std::vector& cuMask = {}, + hipStreamCaptureStatus captureStatus = hipStreamCaptureStatusNone); ~Stream(); - /// Creates the hip stream object, including AMD host queue bool Create(); @@ -173,6 +210,44 @@ namespace hip { /// Sync all non-blocking streams static void syncNonBlockingStreams(); + + /// Returns capture status of the current stream + hipStreamCaptureStatus GetCaptureStatus() const { return captureStatus_; } + /// Returns capture mode of the current stream + hipStreamCaptureMode GetCaptureMode() const { return captureMode_; } + /// Returns if stream is origin stream + bool IsOriginStream() const { return originStream_; } + void SetOriginStream() { originStream_ = true; } + /// Returns captured graph + hipGraph_t GetCaptureGraph() const { return pCaptureGraph_; } + /// Returns last captured graph node + std::vector GetLastCapturedNodes() const { return lastCapturedNodes_; } + /// Set last captured graph node + void SetLastCapturedNode(hipGraphNode_t graphNode) { + lastCapturedNodes_.clear(); + lastCapturedNodes_.push_back(graphNode); + } + /// Append captured node via the wait event cross stream + void AddCrossCapturedNode(std::vector graphNodes) { + for (auto node : graphNodes) { + lastCapturedNodes_.push_back(node); + } + } + /// Set graph that is being captured + void SetCaptureGraph(hipGraph_t pGraph) { + pCaptureGraph_ = pGraph; + captureStatus_ = hipStreamCaptureStatusActive; + } + /// reset capture parameters + hipError_t EndCapture(); + /// Set capture status + void SetCaptureStatus(hipStreamCaptureStatus captureStatus) { captureStatus_ = captureStatus; } + /// Set capture mode + void SetCaptureMode(hipStreamCaptureMode captureMode) { captureMode_ = captureMode; } + /// Set parent stream + void SetParentStream(hipStream_t parentStream) { parentStream_ = parentStream; } + /// Get parent stream + hipStream_t GetParentStream() { return parentStream_; } }; /// HIP Device class diff --git a/projects/clr/hipamd/rocclr/hip_memory.cpp b/projects/clr/hipamd/rocclr/hip_memory.cpp index 18705e2932..df5f83e727 100755 --- a/projects/clr/hipamd/rocclr/hip_memory.cpp +++ b/projects/clr/hipamd/rocclr/hip_memory.cpp @@ -193,7 +193,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipSuccess; } -inline hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, +hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { if (dst == nullptr || src == nullptr) { return hipErrorInvalidValue; @@ -980,6 +980,8 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si 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); + size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -995,6 +997,8 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBy size_t offset, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpyFromSymbolAsync, symbol, dst, sizeBytes, offset, kind, stream); + STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind); + size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; @@ -1035,6 +1039,8 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpyAsync, dst, src, sizeBytes, kind, stream); + STREAM_CAPTURE(hipMemcpyAsync, stream, dst, src, sizeBytes, kind); + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN_DURATION(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true)); @@ -1927,6 +1933,8 @@ hipError_t hipMemcpy3D(const hipMemcpy3DParms* p) { hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) { HIP_INIT_API(hipMemcpy3DAsync, p, stream); + STREAM_CAPTURE(hipMemcpy3DAsync, stream, p); + HIP_RETURN_DURATION(ihipMemcpy3D(p, stream, true)); } @@ -1961,7 +1969,7 @@ hipError_t packFillMemoryCommand(amd::Command*& command, amd::Memory* memory, si return hipSuccess; } -inline hipError_t ihipMemset_validate(void* dst, int64_t value, size_t valueSize, +hipError_t ihipMemset_validate(void* dst, int64_t value, size_t valueSize, size_t sizeBytes) { if (sizeBytes == 0) { // Skip if nothing needs filling. @@ -2075,6 +2083,8 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, 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)); } @@ -2118,7 +2128,7 @@ hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, HIP_RETURN(ihipMemset(dst, value, sizeof(int32_t), count * sizeof(int32_t), stream, true)); } -inline hipError_t ihipMemset3D_validate(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, +hipError_t ihipMemset3D_validate(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, size_t sizeBytes) { size_t offset = 0; amd::Memory* memory = getMemoryObject(pitchedDevPtr.ptr, offset); @@ -2200,6 +2210,8 @@ 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); + STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height); + HIP_RETURN(ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true)); } @@ -2212,6 +2224,8 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) { HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, extent, stream); + STREAM_CAPTURE(hipMemset3DAsync, stream, pitchedDevPtr, value, extent); + HIP_RETURN(ihipMemset3D(pitchedDevPtr, value, extent, stream, true)); } diff --git a/projects/clr/hipamd/rocclr/hip_module.cpp b/projects/clr/hipamd/rocclr/hip_module.cpp index 08ef410670..464473b1a8 100755 --- a/projects/clr/hipamd/rocclr/hip_module.cpp +++ b/projects/clr/hipamd/rocclr/hip_module.cpp @@ -209,7 +209,7 @@ hipError_t hipFuncSetSharedMemConfig ( const void* func, hipSharedMemConfig conf HIP_RETURN(hipSuccess); } -inline hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, +hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, uint32_t sharedMemBytes, @@ -484,7 +484,9 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, hipStream_t stream) { HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); - HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, nullptr, nullptr, 0)); + STREAM_CAPTURE(hipLaunchKernel, stream, hostFunction, gridDim, blockDim, args, sharedMemBytes); + HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, + nullptr, nullptr, 0)); } extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, diff --git a/projects/clr/hipamd/rocclr/hip_stream.cpp b/projects/clr/hipamd/rocclr/hip_stream.cpp index 50ef6960a9..e265b376c2 100755 --- a/projects/clr/hipamd/rocclr/hip_stream.cpp +++ b/projects/clr/hipamd/rocclr/hip_stream.cpp @@ -31,10 +31,16 @@ static std::unordered_set streamSet; namespace hip { // ================================================================================================ -Stream::Stream(hip::Device* dev, Priority p, - unsigned int f, bool null_stream, const std::vector& cuMask) - : queue_(nullptr), lock_("Stream Callback lock"), device_(dev), - priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {} +Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream, + const std::vector& cuMask, hipStreamCaptureStatus captureStatus) + : queue_(nullptr), + lock_("Stream Callback lock"), + device_(dev), + priority_(p), + flags_(f), + null_(null_stream), + cuMask_(cuMask), + captureStatus_(captureStatus) {} // ================================================================================================ Stream::~Stream() { @@ -47,6 +53,25 @@ Stream::~Stream() { } } +hipError_t Stream::EndCapture() { + for (auto event : captureEvents_) { + hip::Event* e = reinterpret_cast(event); + e->EndCapture(); + } + for (auto stream : parallelCaptureStreams_) { + hip::Stream* s = reinterpret_cast(stream); + s->EndCapture(); + } + captureStatus_ = hipStreamCaptureStatusNone; + pCaptureGraph_ = nullptr; + originStream_ = false; + parentStream_ = nullptr; + lastCapturedNodes_.clear(); + parallelCaptureStreams_.clear(); + captureEvents_.clear(); + + return hipSuccess; +} // ================================================================================================ bool Stream::Create() { // Enable queue profiling if a profiler is attached which sets the callback_table flag @@ -345,6 +370,8 @@ void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_d hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(hipStreamWaitEvent, stream, event, flags); + EVENT_CAPTURE(hipStreamWaitEvent, event, stream, flags); + if (event == nullptr) { HIP_RETURN(hipErrorInvalidHandle); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/graph/hipGraph.cpp b/projects/clr/hipamd/tests/src/runtimeApi/graph/hipGraph.cpp new file mode 100644 index 0000000000..a964ec07ff --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/graph/hipGraph.cpp @@ -0,0 +1,274 @@ +/* Copyright (c) 2021-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 +#include +#include +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia + * TEST: %t EXCLUDE_HIP_PLATFORM all + * HIT_END + + */ +#define THREADS_PER_BLOCK 512 +#define GRAPH_LAUNCH_ITERATIONS 3 + +__global__ void reduce(float* d_in, double* d_out, size_t inputSize, size_t outputSize) { + // sdata is allocated in the kernel call: 3rd arg to <<>> + int myId = threadIdx.x + blockDim.x * blockIdx.x; + int tid = threadIdx.x; + + // do reduction in global mem + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + d_in[myId] += d_in[myId + s]; + } + __syncthreads(); // make sure all adds at one stage are done! + } + + // only thread 0 writes result for this block back to global mem + if (tid == 0) { + int blkx = blockIdx.x; + d_out[blockIdx.x] = d_in[myId]; + } +} +__global__ void reduceFinal(double* d_in, double* d_out, size_t inputSize) { + // sdata is allocated in the kernel call: 3rd arg to <<>> + int myId = threadIdx.x + blockDim.x * blockIdx.x; + int tid = threadIdx.x; + + // do reduction in global mem + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + d_in[myId] += d_in[myId + s]; + } + __syncthreads(); // make sure all adds at one stage are done! + } + + // only thread 0 writes result for this block back to global mem + if (tid == 0) { + *d_out = d_in[myId]; + } +} + +void init_input(float* a, size_t size) { + for (size_t i = 0; i < size; i++) a[i] = (rand() & 0xFF) / (float)RAND_MAX; +} + +bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d, + double* result_d, size_t inputSize, size_t numOfBlocks) { + hipStream_t stream1, stream2, stream3, streamForGraph; + hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2; + hipGraph_t graph; + double result_h = 0.0; + + HIPCHECK(hipStreamCreate(&stream1)); + HIPCHECK(hipStreamCreate(&stream2)); + HIPCHECK(hipStreamCreate(&stream3)); + HIPCHECK(hipStreamCreate(&streamForGraph)); + + HIPCHECK(hipEventCreate(&forkStreamEvent)); + HIPCHECK(hipEventCreate(&memsetEvent1)); + HIPCHECK(hipEventCreate(&memsetEvent2)); + + HIPCHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + + HIPCHECK(hipEventRecord(forkStreamEvent, stream1)); + HIPCHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIPCHECK(hipStreamWaitEvent(stream3, forkStreamEvent, 0)); + + HIPCHECK( + hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(float) * inputSize, hipMemcpyDefault, stream1)); + + HIPCHECK(hipMemsetAsync(outputVec_d, 0, sizeof(double) * numOfBlocks, stream2)); + + HIPCHECK(hipEventRecord(memsetEvent1, stream2)); + + HIPCHECK(hipMemsetAsync(result_d, 0, sizeof(double), stream3)); + HIPCHECK(hipEventRecord(memsetEvent2, stream3)); + + HIPCHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0)); + + hipLaunchKernelGGL(reduce, dim3(inputSize / THREADS_PER_BLOCK, 1, 1), + dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, inputVec_d, outputVec_d, inputSize, + numOfBlocks); + HIPCHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0)); + + hipLaunchKernelGGL(reduceFinal, dim3(1, 1, 1), dim3(THREADS_PER_BLOCK, 1, 1), 0, stream1, + outputVec_d, result_d, numOfBlocks); + HIPCHECK(hipMemcpyAsync(&result_h, result_d, sizeof(double), hipMemcpyDefault, stream1)); + + HIPCHECK(hipStreamEndCapture(stream1, &graph)); + + hipGraphExec_t graphExec; + HIPCHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { + HIPCHECK(hipGraphLaunch(graphExec, streamForGraph)); + } + HIPCHECK(hipStreamSynchronize(streamForGraph)); + HIPCHECK(hipGraphExecDestroy(graphExec)); + HIPCHECK(hipGraphDestroy(graph)); + HIPCHECK(hipStreamDestroy(stream1)); + HIPCHECK(hipStreamDestroy(stream2)); + HIPCHECK(hipStreamDestroy(streamForGraph)); + double result_h_cpu = 0.0; + + for (int i = 0; i < inputSize; i++) { + result_h_cpu += inputVec_h[i]; + } + if (result_h_cpu != result_h) { + printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h); + return false; + } + return true; +} + +bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, + size_t inputSize, size_t numOfBlocks) { + hipStream_t streamForGraph; + hipGraph_t graph; + std::vector nodeDependencies; + hipGraphNode_t memcpyNode, kernelNode, memsetNode; + double result_h = 0.0; + HIPCHECK(hipStreamCreate(&streamForGraph)); + hipKernelNodeParams kernelNodeParams = {0}; + hipMemcpy3DParms memcpyParams = {0}; + hipMemsetParams memsetParams = {0}; + memcpyParams.srcArray = NULL; + memcpyParams.srcPos = make_hipPos(0, 0, 0); + memcpyParams.srcPtr = make_hipPitchedPtr(inputVec_h, sizeof(float) * inputSize, inputSize, 1); + memcpyParams.dstArray = NULL; + memcpyParams.dstPos = make_hipPos(0, 0, 0); + memcpyParams.dstPtr = make_hipPitchedPtr(inputVec_d, sizeof(float) * inputSize, inputSize, 1); + memcpyParams.extent = make_hipExtent(sizeof(float) * inputSize, 1, 1); + memcpyParams.kind = hipMemcpyHostToDevice; + memsetParams.dst = (void*)outputVec_d; + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(float); // elementSize can be max 4 bytes + memsetParams.width = numOfBlocks * 2; + memsetParams.height = 1; + HIPCHECK(hipGraphCreate(&graph, 0)); + HIPCHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, NULL, 0, &memcpyParams)); + HIPCHECK(hipGraphAddMemsetNode(&memsetNode, graph, NULL, 0, &memsetParams)); + nodeDependencies.push_back(memsetNode); + nodeDependencies.push_back(memcpyNode); + void* kernelArgs[4] = {(void*)&inputVec_d, (void*)&outputVec_d, &inputSize, &numOfBlocks}; + kernelNodeParams.func = (void*)reduce; + kernelNodeParams.gridDim = dim3(inputSize / THREADS_PER_BLOCK, 1, 1); + kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = (void**)kernelArgs; + kernelNodeParams.extra = NULL; + HIPCHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &kernelNodeParams)); + nodeDependencies.clear(); + nodeDependencies.push_back(kernelNode); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = result_d; + memsetParams.value = 0; + memsetParams.elementSize = sizeof(float); + memsetParams.width = 2; + memsetParams.height = 1; + HIPCHECK(hipGraphAddMemsetNode(&memsetNode, graph, NULL, 0, &memsetParams)); + nodeDependencies.push_back(memsetNode); + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = (void*)reduceFinal; + kernelNodeParams.gridDim = dim3(1, 1, 1); + kernelNodeParams.blockDim = dim3(THREADS_PER_BLOCK, 1, 1); + kernelNodeParams.sharedMemBytes = 0; + void* kernelArgs2[3] = {(void*)&outputVec_d, (void*)&result_d, &numOfBlocks}; + kernelNodeParams.kernelParams = kernelArgs2; + kernelNodeParams.extra = NULL; + HIPCHECK(hipGraphAddKernelNode(&kernelNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &kernelNodeParams)); + nodeDependencies.clear(); + nodeDependencies.push_back(kernelNode); + memset(&memcpyParams, 0, sizeof(memcpyParams)); + memcpyParams.srcArray = NULL; + memcpyParams.srcPos = make_hipPos(0, 0, 0); + memcpyParams.srcPtr = make_hipPitchedPtr(result_d, sizeof(double), 1, 1); + memcpyParams.dstArray = NULL; + memcpyParams.dstPos = make_hipPos(0, 0, 0); + memcpyParams.dstPtr = make_hipPitchedPtr(&result_h, sizeof(double), 1, 1); + memcpyParams.extent = make_hipExtent(sizeof(double), 1, 1); + memcpyParams.kind = hipMemcpyDeviceToHost; + HIPCHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &memcpyParams)); + nodeDependencies.clear(); + nodeDependencies.push_back(memcpyNode); + hipGraphNode_t hostNode; + + hipGraphExec_t graphExec; + HIPCHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + for (int i = 0; i < GRAPH_LAUNCH_ITERATIONS; i++) { + HIPCHECK(hipGraphLaunch(graphExec, streamForGraph)); + } + HIPCHECK(hipStreamSynchronize(streamForGraph)); + + HIPCHECK(hipGraphExecDestroy(graphExec)); + HIPCHECK(hipGraphDestroy(graph)); + HIPCHECK(hipStreamDestroy(streamForGraph)); + + double result_h_cpu = 0.0; + + for (int i = 0; i < inputSize; i++) { + result_h_cpu += inputVec_h[i]; + } + if (result_h_cpu != result_h) { + printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h); + return false; + } + return true; +} + +int main(int argc, char** argv) { + size_t size = 1 << 12; // number of elements to reduce + size_t maxBlocks = 512; + // This will pick the best possible CUDA capable device + int devID = 1; // TODO: implement: findCudaDevice(argc, (const char**)argv); based of max GFLOPS + // incase of multiple devic + hipSetDevice(0); // + printf("%zu elements\n", size); + printf("threads per block = %d\n", THREADS_PER_BLOCK); + printf("Graph Launch iterations = %d\n", GRAPH_LAUNCH_ITERATIONS); + float *inputVec_d = NULL, *inputVec_h = NULL; + double *outputVec_d = NULL, *result_d; + inputVec_h = (float*)malloc(sizeof(float) * size); + HIPCHECK(hipMalloc(&inputVec_d, sizeof(float) * size)); + HIPCHECK(hipMalloc(&outputVec_d, sizeof(double) * maxBlocks)); + HIPCHECK(hipMalloc(&result_d, sizeof(double))); + init_input(inputVec_h, size); + bool status1 = hipGraphsManual(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); + bool status2 = hipGraphsUsingStreamCapture(inputVec_h, inputVec_d, outputVec_d, result_d, size, maxBlocks); + HIPCHECK(hipFree(inputVec_d)); + HIPCHECK(hipFree(outputVec_d)); + HIPCHECK(hipFree(result_d)); + if(!status1) { + failed("Failed during hip Graph Manual\n"); + } + if(!status2) { + failed("Failed during hip Graphs during stream capture\n"); + } + passed(); +}