AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-422207 - Handle nonkernel nodes for graph opt

Browse Source 
- Support graph with different types of nodes with single
branch when DEBUG_CLR_GRAPH_PACKET_CAPTURE flag is enabled

Change-Id: I149a8629769cd0d5849ffefb04f1352668a685b6
This commit is contained in:
Anusha GodavarthySurya
2023-10-11 15:35:57 +00:00
committed by Saleel Kudchadker
parent 6926183974
commit 38d2c56784
4 changed files with 95 additions and 61 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_graph_internal.cpp
+59 -40
View File
@@ -494,51 +494,65 @@ hipError_t GraphExec::Init() {
hipError_t GraphExec::CaptureAQLPackets() {
hipError_t status = hipSuccess;
size_t KernArgSizeForGraph = 0;
bool GraphHasOnlyKerns = true;
// GPU packet capture is enabled for kernel nodes. Calculate the kernel arg size required for all
// graph kernel nodes to allocate
for (const auto& list : parallelLists_) {
hip::Stream* stream = GetAvailableStreams();
for (auto& node : list) {
node->SetStream(stream, this);
if (node->GetType() == hipGraphNodeTypeKernel) {
KernArgSizeForGraph += reinterpret_cast<hip::GraphKernelNode*>(node)->GetKerArgSize();
} else {
GraphHasOnlyKerns = false;
if (parallelLists_.size() == 1) {
size_t kernArgSizeForGraph = 0;
// GPU packet capture is enabled for kernel nodes. Calculate the kernel
// arg size required for all graph kernel nodes to allocate
for (const auto& list : parallelLists_) {
hip::Stream* stream = GetAvailableStreams();
for (auto& node : list) {
node->SetStream(stream, this);
if (node->GetType() == hipGraphNodeTypeKernel) {
kernArgSizeForGraph += reinterpret_cast<hip::GraphKernelNode*>(node)->GetKerArgSize();
}
}
}
}
auto device = g_devices[ihipGetDevice()]->devices()[0];
const auto& info = device->info();
// Enable allocating kerns on device memory if graph as only kernels. memcpy nodes require hdp
// flush. ToDo: Work on enabling device kern args later for all type of nodes for large bar
if (GraphHasOnlyKerns == true && info.largeBar_) {
kernarg_pool_graph_ = reinterpret_cast<address>(device->deviceLocalAlloc(KernArgSizeForGraph));
device_kernarg_pool_ = true;
} else {
kernarg_pool_graph_ = reinterpret_cast<address>(
device->hostAlloc(KernArgSizeForGraph, 0, amd::Device::MemorySegment::kKernArg));
}
auto device = g_devices[ihipGetDevice()]->devices()[0];
if (device->info().largeBar_) {
// Pad kernel argument buffer with sentinal size bytes to do a readback later
kernArgSizeForGraph += sizeof(int);
kernarg_pool_graph_ =
reinterpret_cast<address>(device->deviceLocalAlloc(kernArgSizeForGraph));
device_kernarg_pool_ = true;
} else {
kernarg_pool_graph_ = reinterpret_cast<address>(
device->hostAlloc(kernArgSizeForGraph, 0, amd::Device::MemorySegment::kKernArg));
}
if (kernarg_pool_graph_ == nullptr) {
return hipErrorMemoryAllocation;
}
kernarg_pool_size_graph_ = KernArgSizeForGraph;
if (kernarg_pool_graph_ == nullptr) {
return hipErrorMemoryAllocation;
}
kernarg_pool_size_graph_ = kernArgSizeForGraph;
for (auto& node : topoOrder_) {
if (node->GetType() == hipGraphNodeTypeKernel) {
auto kernelnode = reinterpret_cast<hip::GraphKernelNode*>(node);
status = node->CreateCommand(node->GetQueue());
// From the kernel pool allocate the kern arg size required for the current kernel node.
address kernArgOffset = allocKernArg(kernelnode->GetKernargSegmentByteSize(),
kernelnode->GetKernargSegmentAlignment());
if (kernArgOffset == nullptr) {
return hipErrorMemoryAllocation;
for (auto& node : topoOrder_) {
if (node->GetType() == hipGraphNodeTypeKernel) {
auto kernelNode = reinterpret_cast<hip::GraphKernelNode*>(node);
status = node->CreateCommand(node->GetQueue());
// From the kernel pool allocate the kern arg size required for the current kernel node.
address kernArgOffset = allocKernArg(kernelNode->GetKernargSegmentByteSize(),
kernelNode->GetKernargSegmentAlignment());
if (kernArgOffset == nullptr) {
return hipErrorMemoryAllocation;
}
// Form GPU packet capture for the kernel node.
kernelNode->CaptureAndFormPacket(kernArgOffset);
}
}
if (device_kernarg_pool_) {
// Write HDP_MEM_COHERENCY_FLUSH_CNTL reg to initiate flush read to HDP mem. Verify mem
// by readback of sentinal value at the tail end of the kernarg surface (allocated above)
// This needs to be done for PCIE connected devices only. HDP path is disabled for XGMI
// between CPU<->GPU
if (!device->isXgmi()) {
static int host_val = 1;
address dev_ptr = kernarg_pool_graph_ + kernarg_pool_size_graph_ - sizeof(int);
*dev_ptr = host_val;
*device->info().hdpMemFlushCntl = 1;
while (*dev_ptr != host_val);
host_val++;
}
// Enable GPU packet capture for the kernel node.
kernelnode->EnableCapturing(kernArgOffset);
}
}
return status;
@@ -561,9 +575,11 @@ hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
}
node->UpdateEventWaitLists(waitList);
}
std::vector<Node> rootNodes = clonedGraph->GetRootNodes();
ClPrint(amd::LOG_INFO, amd::LOG_CODE,
"[hipGraph] RootCommand get launched on stream (stream:%p)\n", stream);
for (auto& root : rootNodes) {
//If rootnode is launched on to the same stream dont add dependency
if (root->GetQueue() != stream) {
@@ -653,14 +669,16 @@ hipError_t GraphExec::Run(hipStream_t stream) {
} else {
repeatLaunch_ = true;
}
if (parallelLists_.size() == 1) {
if (device_kernarg_pool_) {
// If kernelArgs are in device memory flush the HDP.
// If kernelArgs are in device memory flush/invalidate L2
amd::Command* startCommand = nullptr;
startCommand = new amd::Marker(*hip_stream, false);
startCommand->enqueue();
startCommand->release();
}
for (int i = 0; i < topoOrder_.size(); i++) {
if (DEBUG_CLR_GRAPH_PACKET_CAPTURE && topoOrder_[i]->GetType() == hipGraphNodeTypeKernel) {
hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket());
@@ -670,6 +688,7 @@ hipError_t GraphExec::Run(hipStream_t stream) {
topoOrder_[i]->EnqueueCommands(stream);
}
}
if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
amd::Command* endCommand = nullptr;
endCommand = new amd::Marker(*hip_stream, false);
hipamd/src/hip_graph_internal.hpp
+12 -3
View File
@@ -463,6 +463,7 @@ struct Graph {
graphExeUserObjs.insert(userObj);
}
}
Graph* clone(std::unordered_map<Node, Node>& clonedNodes) const;
Graph* clone() const;
void GenerateDOT(std::ostream& fout, hipGraphDebugDotFlags flag) {
@@ -612,6 +613,7 @@ struct GraphExec {
}
return clonedNode;
}
address allocKernArg(size_t size, size_t alignment) {
assert(alignment != 0);
address result = nullptr;
@@ -622,12 +624,19 @@ struct GraphExec {
}
return result;
}
// check executable graphs validity
static bool isGraphExecValid(GraphExec* pGraphExec);
std::vector<Node>& GetNodes() { return topoOrder_; }
hip::Stream* GetAvailableStreams() { return parallel_streams_[currentQueueIndex_++]; }
hip::Stream* GetAvailableStreams() {
if (currentQueueIndex_ < parallel_streams_.size()) {
return parallel_streams_[currentQueueIndex_++];
}
return nullptr;
}
void ResetQueueIndex() { currentQueueIndex_ = 0; }
hipError_t Init();
hipError_t CreateStreams(uint32_t num_streams);
@@ -791,12 +800,12 @@ class GraphKernelNode : public GraphNode {
out << "];";
}
void EnableCapturing(address kernArgOffset) {
void CaptureAndFormPacket(address kernArgOffset) {
for (auto& command : commands_) {
reinterpret_cast<amd::NDRangeKernelCommand*>(command)->setCapturingState(
true, GetAqlPacket(), kernArgOffset);
// Enqueue command to capture GPU Packet. Packet is not sent to hardware queue.
// Enqueue command to capture GPU Packet. Packet is not sent to hardware queue.
command->submit(*(command->queue())->vdev());
command->release();
}
rocclr/device/device.hpp
+5
View File
@@ -1758,6 +1758,11 @@ class Device : public RuntimeObject {
return NULL;
}
virtual bool isXgmi() const {
ShouldNotCallThis();
return false;
}
virtual bool deviceAllowAccess(void* dst) const {
ShouldNotCallThis();
return true;
rocclr/device/rocm/rocvirtual.cpp
+19 -18
View File
@@ -521,7 +521,6 @@ std::vector<hsa_signal_t>& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi
// Validate all signals for the wait and skip already completed
for (uint32_t i = 0; i < external_signals_.size(); ++i) {
// Early signal status check
if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) {
const Settings& settings = gpu_.dev().settings();
// Actively wait on CPU to avoid extra overheads of signal tracking on GPU.
@@ -829,6 +828,13 @@ bool VirtualGPU::dispatchGenericAqlPacket(
// Check for queue full and wait if needed.
uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size);
uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_);
if (addSystemScope_) {
header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
addSystemScope_ = false;
}
auto expected_fence_state = extractAqlBits(header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE);
@@ -898,6 +904,10 @@ bool VirtualGPU::dispatchGenericAqlPacket(
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);
// Mark the flag indicating if a dispatch is outstanding.
// We are not waiting after every dispatch.
hasPendingDispatch_ = true;
// Wait on signal ?
if (blocking) {
LogInfo("Runtime reachead the AQL queue limit. SW is much ahead of HW. Blocking AQL queue!");
@@ -929,18 +939,14 @@ void VirtualGPU::dispatchBlockingWait() {
bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header,
uint16_t rest, bool blocking, bool capturing,
const uint8_t* aqlPacket) {
dispatchBlockingWait();
if (capturing == true) {
packet->header = header;
packet->setup = rest;
if (timestamp_ != nullptr) {
// Get active signal for current dispatch if profiling is necessary
packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
}
amd::Os::fastMemcpy(const_cast<uint8_t*>(aqlPacket), packet,
sizeof(hsa_kernel_dispatch_packet_t));
return true;
} else {
dispatchBlockingWait();
return dispatchGenericAqlPacket(packet, header, rest, blocking);
}
}
@@ -1053,12 +1059,16 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
}
inline bool VirtualGPU::dispatchAqlPacket(uint8_t* aqlpacket) {
dispatchBlockingWait();
auto packet = reinterpret_cast<hsa_kernel_dispatch_packet_t*>(aqlpacket);
// If rocprof tracing is enabled, store the correlation ID in the dispatch packet.
// The profiler can retrieve this correlation ID to attribute waves to specific dispatch
// locations.
if (activity_prof::IsEnabled(OP_ID_DISPATCH)) {
if (activity_prof::IsEnabled(OP_ID_DISPATCH) ||
(roc_device_.info().queueProperties_ & CL_QUEUE_PROFILING_ENABLE)) {
packet->reserved2 = activity_prof::correlation_id;
// Get active signal for current dispatch if profiling is necessary
packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
}
dispatchGenericAqlPacket(packet, packet->header, packet->setup, false);
return true;
@@ -3227,13 +3237,8 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
aqlHeaderWithOrder &= kAqlHeaderMask;
}
if (addSystemScope_ || (vcmd != nullptr &&
vcmd->getEventScope() == amd::Device::kCacheStateSystem)) {
aqlHeaderWithOrder &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
aqlHeaderWithOrder |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
addSystemScope_ = false;
if (vcmd != nullptr && vcmd->getEventScope() == amd::Device::kCacheStateSystem) {
addSystemScope_ = true;
}
// If profiling is enabled, store the correlation ID in the dispatch packet. The profiler can
@@ -3268,10 +3273,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
}
}
// Mark the flag indicating if a dispatch is outstanding.
// We are not waiting after every dispatch.
hasPendingDispatch_ = true;
// Output printf buffer
if (!printfDbg()->output(*this, printfEnabled, gpuKernel.printfInfo())) {
LogError("\nCould not print data from the printf buffer!");