SWDEV-240806 - Initial commit for hipGraph and stream capture infrastructure
On StreamBegincapture captures the parameters passed to APIs and respective node will be created and added to graph
All parameters are passed to STREAM_CAPTURE macro, it checks if stream in capture mode and redirects the call to the capture function and returns
Updated hipStream and hipEvent with capture parameters
Added handling for hipStreamBeginCapture & hipStreamEndCapture
Change-Id: Ic8926a7b4336c2cc81f0b3a9a224aa392c474134
[ROCm/clr commit: 8cc0e04239]
This commit is contained in:
Anusha Godavarthy Surya
committed by
Christophe Paquot
Christophe Paquot
parent
abc42e8c5c
commit
1fe1fe9361
@@ -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__
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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)
|
||||
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
@@ -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<hipGraphNode_t> 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<hipGraphNode_t> GetNodesPrevToRecorded() const { return nodesPrevToRecorded_; }
|
||||
/// Set last captured graph node before event record
|
||||
void SetNodesPrevToRecorded(std::vector<hipGraphNode_t>& graphNode) {
|
||||
nodesPrevToRecorded_ = graphNode;
|
||||
}
|
||||
|
||||
// IPC Events
|
||||
struct ihipIpcEvent_t {
|
||||
std::string ipc_name_;
|
||||
int ipc_fd_;
|
||||
|
||||
@@ -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<hipStream_t> g_captureStreams;
|
||||
std::unordered_map<amd::Command*, hipGraphExec_t> 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<size_t>(pNodeParams->gridDim.x) * pNodeParams->blockDim.x;
|
||||
size_t globalWorkSizeY = static_cast<size_t>(pNodeParams->gridDim.y) * pNodeParams->blockDim.y;
|
||||
size_t globalWorkSizeZ = static_cast<size_t>(pNodeParams->gridDim.z) * pNodeParams->blockDim.z;
|
||||
if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
|
||||
globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
|
||||
globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
|
||||
return hipErrorInvalidConfiguration;
|
||||
}
|
||||
status = ihipLaunchKernel_validate(
|
||||
func, static_cast<uint32_t>(globalWorkSizeX), static_cast<uint32_t>(globalWorkSizeY),
|
||||
static_cast<uint32_t>(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<hip::Stream*>(stream);
|
||||
hipKernelNodeParams nodeParams;
|
||||
nodeParams.func = const_cast<void*>(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<hip::Stream*>(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<hip::Stream*>(stream);
|
||||
hipGraphNode_t pGraphNode;
|
||||
hipGraph_t graph = nullptr;
|
||||
std::vector<hipGraphNode_t> 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<hip::Stream*>(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<hip::Stream*>(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<hip::Event*>(event);
|
||||
e->StartCapture(stream);
|
||||
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
|
||||
std::vector<hipGraphNode_t> 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<hip::Stream*>(stream);
|
||||
hip::Event* e = reinterpret_cast<hip::Event*>(event);
|
||||
|
||||
if (event == nullptr || stream == nullptr) {
|
||||
HIP_RETURN(hipErrorInvalidValue);
|
||||
}
|
||||
if (!s->IsOriginStream()) {
|
||||
s->SetCaptureGraph(reinterpret_cast<hip::Stream*>(e->GetCaptureStream())->GetCaptureGraph());
|
||||
s->SetCaptureMode(reinterpret_cast<hip::Stream*>(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<hip::Stream*>(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<hip::Stream*>(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<hip::Stream*>(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<std::vector<Node>> parallelLists;
|
||||
std::unordered_map<Node, std::vector<Node>> nodeWaitLists;
|
||||
graph->GetRunList(parallelLists, nodeWaitLists);
|
||||
std::vector<Node> 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));
|
||||
}
|
||||
@@ -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);
|
||||
@@ -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<amd::Command*>& commands, void* dst, int64_t value,
|
||||
size_t valueSize, size_t sizeBytes, amd::HostQueue* queue);
|
||||
|
||||
hipError_t ihipMemset3DCommand(std::vector<amd::Command*>& commands, hipPitchedPtr pitchedDevPtr,
|
||||
int value, hipExtent extent, amd::HostQueue* queue);
|
||||
@@ -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 <queue>
|
||||
|
||||
#define CASE_STRING(X, C) \
|
||||
case X: \
|
||||
case_string = #C; \
|
||||
break;
|
||||
const char* GetGraphNodeTypeString(uint32_t op) {
|
||||
const char* case_string;
|
||||
switch (static_cast<hipGraphNodeType>(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<Node> hipGraph::GetRootNodes() const {
|
||||
std::vector<Node> 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<Node> hipGraph::GetLeafNodes() const {
|
||||
std::vector<Node> 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<std::pair<Node, Node>> hipGraph::GetEdges() const {
|
||||
std::vector<std::pair<Node, Node>> 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<Node, bool>& visited,
|
||||
std::vector<Node>& singleList,
|
||||
std::vector<std::vector<Node>>& parallelLists,
|
||||
std::unordered_map<Node, std::vector<Node>>& 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<std::vector<Node>>& parallelLists,
|
||||
std::unordered_map<Node, std::vector<Node>>& dependencies) {
|
||||
std::vector<Node> singleList;
|
||||
|
||||
// Mark all the vertices as not visited
|
||||
std::unordered_map<Node, bool> 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<Node>& levelOrder) {
|
||||
std::vector<Node> roots = GetRootNodes();
|
||||
std::unordered_map<Node, bool> visited;
|
||||
std::queue<Node> 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<amd::Command*>(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;
|
||||
}
|
||||
@@ -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 <algorithm>
|
||||
#include <queue>
|
||||
#include <stack>
|
||||
#include <iostream>
|
||||
#include <unordered_map>
|
||||
#include <unordered_set>
|
||||
#include <vector>
|
||||
|
||||
#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<amd::Command*> 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<amd::Command*>& GetCommands() { return commands_; }
|
||||
hipGraphNodeType GetType() { return type_; }
|
||||
uint32_t GetLevel() { return level_; }
|
||||
void SetLevel(uint32_t level) { level_ = level; }
|
||||
};
|
||||
|
||||
class hipGraph {
|
||||
std::unordered_map<Node, size_t> nodeInDegree_; // count of in coming edges for every vertex
|
||||
std::unordered_map<Node, size_t> nodeOutDegree_; // count of outgoing edges for every vertex
|
||||
std::vector<Node> vertices_;
|
||||
std::unordered_map<Node, std::vector<Node>> 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<Node> GetRootNodes() const;
|
||||
/// Returns leaf nodes, all vertices with 0 out-degrees
|
||||
std::vector<Node> 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<Node> GetNodes() const { return vertices_; }
|
||||
/// returns all the edges in the graph
|
||||
std::vector<std::pair<Node, Node>> GetEdges() const;
|
||||
void GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
|
||||
std::vector<Node>& singleList, std::vector<std::vector<Node>>& parallelList,
|
||||
std::unordered_map<Node, std::vector<Node>>& dependencies);
|
||||
void GetRunList(std::vector<std::vector<Node>>& parallelList,
|
||||
std::unordered_map<Node, std::vector<Node>>& dependencies);
|
||||
hipError_t LevelOrder(std::vector<Node>& 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 T> 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 T> 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<std::vector<Node>> parallelLists_;
|
||||
std::vector<Node> levelOrder_;
|
||||
std::unordered_map<Node, std::vector<Node>> nodeWaitLists_;
|
||||
std::vector<amd::HostQueue*> parallelQueues_;
|
||||
static std::unordered_map<amd::Command*, hipGraphExec_t> activeGraphExec_;
|
||||
amd::Command::EventWaitList graphLastCmdWaitList_;
|
||||
amd::Command* lastEnqueuedGraphCmd_;
|
||||
std::atomic<bool> bExecPending_;
|
||||
amd::Command* rootCommand_;
|
||||
|
||||
public:
|
||||
hipGraphExec(std::vector<Node>& levelOrder, std::vector<std::vector<Node>>& lists,
|
||||
std::unordered_map<Node, std::vector<Node>>& 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);
|
||||
};
|
||||
@@ -275,4 +275,17 @@ __gnu_f2h_ieee
|
||||
hipExtStreamGetCUMask
|
||||
hipImportExternalMemory
|
||||
hipExternalMemoryGetMappedBuffer
|
||||
hipDestroyExternalMemory
|
||||
hipDestroyExternalMemory
|
||||
hipGraphCreate
|
||||
hipGraphDestroy
|
||||
hipGraphAddKernelNode
|
||||
hipGraphAddMemsetNode
|
||||
hipGraphAddMemcpyNode
|
||||
hipGraphAddMemcpyNode1D
|
||||
hipGraphAddHostNode
|
||||
hipGraphInstantiate
|
||||
hipGraphLaunch
|
||||
hipStreamIsCapturing
|
||||
hipStreamBeginCapture
|
||||
hipStreamEndCapture
|
||||
hipGraphExecDestroy
|
||||
|
||||
@@ -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*;
|
||||
|
||||
@@ -26,6 +26,9 @@
|
||||
#include "trace_helper.h"
|
||||
#include "utils/debug.hpp"
|
||||
#include "hip_formatting.hpp"
|
||||
|
||||
#include "hip_graph_capture.hpp"
|
||||
|
||||
#include <unordered_set>
|
||||
#include <thread>
|
||||
#include <stack>
|
||||
@@ -123,6 +126,19 @@ typedef struct ihipIpcEventHandle_st {
|
||||
} \
|
||||
} while (0);
|
||||
|
||||
#define STREAM_CAPTURE(name, stream, ...) \
|
||||
if (stream != nullptr && \
|
||||
reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus() == \
|
||||
hipStreamCaptureStatusActive) { \
|
||||
hipError_t status = capture##name(stream, ##__VA_ARGS__); \
|
||||
HIP_RETURN(status); \
|
||||
}
|
||||
|
||||
#define EVENT_CAPTURE(name, event, ...) \
|
||||
if (event != nullptr && reinterpret_cast<hip::Event*>(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<uint32_t> 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<hipGraphNode_t> lastCapturedNodes_;
|
||||
/// Derived streams/Paralell branches from the origin stream
|
||||
std::vector<hipStream_t> parallelCaptureStreams_;
|
||||
/// Capture events
|
||||
std::vector<hipEvent_t> captureEvents_;
|
||||
|
||||
public:
|
||||
Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false,
|
||||
const std::vector<uint32_t>& cuMask = {});
|
||||
const std::vector<uint32_t>& 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<hipGraphNode_t> 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<hipGraphNode_t> 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
|
||||
|
||||
@@ -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));
|
||||
}
|
||||
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -31,10 +31,16 @@ static std::unordered_set<hip::Stream*> streamSet;
|
||||
namespace hip {
|
||||
|
||||
// ================================================================================================
|
||||
Stream::Stream(hip::Device* dev, Priority p,
|
||||
unsigned int f, bool null_stream, const std::vector<uint32_t>& 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<uint32_t>& 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<hip::Event*>(event);
|
||||
e->EndCapture();
|
||||
}
|
||||
for (auto stream : parallelCaptureStreams_) {
|
||||
hip::Stream* s = reinterpret_cast<hip::Stream*>(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);
|
||||
}
|
||||
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
#include <chrono>
|
||||
#include <test_common.h>
|
||||
#include <vector>
|
||||
/* 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 <<<b, t, shmem>>>
|
||||
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 <<<b, t, shmem>>>
|
||||
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<hipGraphNode_t> 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();
|
||||
}
|
||||
Reference in New Issue
Block a user