AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-392732 - Initial commit for graph doorbell optimization(AQL Buffering)

Kaynağa Gözat 
Change-Id: I451725006c54c249dc530c55d2af2a31594bf49b
Bu işleme şunda yer alıyor:
Anusha GodavarthySurya
2023-06-17 13:48:05 +00:00
işlemeyi yapan: Anusha Godavarthy Surya
ebeveyn 84559d0d6f
işleme b0e6f99ad7
7 değiştirilmiş dosya ile 149 ekleme ve 19 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
hipamd/src/hip_graph_internal.cpp
+14 -2
Dosyayı Görüntüle
@@ -611,8 +611,20 @@ hipError_t hipGraphExec::Run(hipStream_t stream) {
cmd->enqueue();
cmd->release();
}
for (auto& node : topoOrder_) {
node->EnqueueCommands(stream);
for (int i = 0; i < topoOrder_.size(); i++) {
if (DEBUG_CLR_GRAPH_ENABLE_BUFFERING) {
// Enable buffering for graph with single branch
if (parallelLists_.size() == 1) {
// Peep through the next node. If current and next node are kernel then enable AQL
// buffering
if (((i + 1) != topoOrder_.size()) &&
topoOrder_[i]->GetType() == hipGraphNodeTypeKernel &&
topoOrder_[i + 1]->GetType() == hipGraphNodeTypeKernel) {
topoOrder_[i]->EnableBuffering();
}
}
}
topoOrder_[i]->EnqueueCommands(stream);
}
if (endCommand != nullptr) {
endCommand->enqueue();
hipamd/src/hip_graph_internal.hpp
+5
Dosyayı Görüntüle
@@ -326,6 +326,11 @@ struct hipGraphNode : public hipGraphNodeDOTAttribute {
command->release();
}
}
virtual void EnableBuffering() {
for (auto& command : commands_) {
command->setBufferingState(true);
}
}
ihipGraph* GetParentGraph() { return parentGraph_; }
virtual ihipGraph* GetChildGraph() { return nullptr; }
void SetParentGraph(ihipGraph* graph) { parentGraph_ = graph; }
rocclr/device/rocm/rocvirtual.cpp
+105 -9
Dosyayı Görüntüle
@@ -810,6 +810,74 @@ static inline void packet_store_release(uint32_t* packet, uint16_t header, uint1
__atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
}
bool VirtualGPU::dispatchAqlBuffer() {
size_t size = aqlBuffer_.size();
if (size > 0) {
const uint32_t queueSize = gpu_queue_->size;
const uint32_t queueMask = queueSize - 1;
const uint32_t sw_queue_size = queueMask;
// 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_);
while (index - hsa_queue_load_read_index_scacquire(gpu_queue_) >= sw_queue_size - size) {
amd::Os::yield();
}
if (timestamp_ != nullptr) {
for (uint i = 0; i < size; i++) {
// Get active signal for current dispatch if profiling is necessary
aqlBuffer_[i].completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
}
}
for (uint i = 0; i < size; i++) {
ClPrint(
amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx, Dispatch AQL Buffer Header = "
"0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
"setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, "
"group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx",
gpu_queue_->base_address, aqlBuffer_[i].header,
extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_BARRIER,
HSA_PACKET_HEADER_WIDTH_BARRIER),
extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
extractAqlBits(aqlBuffer_[i].header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE),
aqlBuffer_[i].setup, (aqlBuffer_[i]).grid_size_x, (aqlBuffer_[i]).grid_size_y,
(aqlBuffer_[i]).grid_size_z, (aqlBuffer_[i]).workgroup_size_x,
(aqlBuffer_[i]).workgroup_size_y, (aqlBuffer_[i]).workgroup_size_z,
(aqlBuffer_[i]).private_segment_size, (aqlBuffer_[i]).group_segment_size,
(aqlBuffer_[i]).kernel_object, (aqlBuffer_[i]).kernarg_address,
(aqlBuffer_[i]).completion_signal);
}
uint16_t firstPacketHeader = aqlBuffer_.front().header;
aqlBuffer_.front().header = kInvalidAql;
hsa_kernel_dispatch_packet_t* aql_loc =
&((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[index & queueMask];
size_t size_before_wrap = queueSize - (index % queueSize);
if (size_before_wrap < size) {
amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size_before_wrap);
hsa_kernel_dispatch_packet_t* aql_loc_0 =
&((hsa_kernel_dispatch_packet_t*)(gpu_queue_->base_address))[0];
amd::Os::fastMemcpy(aql_loc_0, &aqlBuffer_[size_before_wrap],
sizeof(hsa_kernel_dispatch_packet_t) * (size - size_before_wrap));
} else {
amd::Os::fastMemcpy(aql_loc, &aqlBuffer_[0], sizeof(hsa_kernel_dispatch_packet_t) * size);
}
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), firstPacketHeader,
aqlBuffer_.front().setup);
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
aqlBuffer_.clear();
}
return true;
}
// ================================================================================================
template <typename AqlPacket>
bool VirtualGPU::dispatchGenericAqlPacket(
@@ -919,13 +987,36 @@ void VirtualGPU::dispatchBlockingWait() {
}
// ================================================================================================
bool VirtualGPU::dispatchAqlPacket(
hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) {
bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header,
uint16_t rest, bool blocking, bool buffering) {
static size_t initialAQLBufferSize = 1;
dispatchBlockingWait();
return dispatchGenericAqlPacket(packet, header, rest, blocking);
if (buffering == true) {
aqlBuffer_.push_back(*packet);
aqlBuffer_.back().header = header;
aqlBuffer_.back().setup = rest;
if (!(aqlBuffer_.size() >=
initialAQLBufferSize)) { // Buffer maximum of AQL Buffer size packets once it
// exceeds send them for dispatch
return true;
}
} else if (!aqlBuffer_.empty()) { // If buffering is disabled and AQLBuffer is not empty then
// make sure current packet is added to the buffer for dispatch
aqlBuffer_.push_back(*packet);
aqlBuffer_.back().header = header;
aqlBuffer_.back().setup = rest;
}
if (aqlBuffer_.empty()) {
return dispatchGenericAqlPacket(packet, header, rest, blocking);
} else {
ClPrint(amd::LOG_DEBUG, amd::LOG_CODE, "Dispath AQL Buffer size:%ld", aqlBuffer_.size());
// Increment buffer size ^2 until DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE
if (initialAQLBufferSize < DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE) {
initialAQLBufferSize = initialAQLBufferSize << 1;
}
return dispatchAqlBuffer();
}
}
// ================================================================================================
bool VirtualGPU::dispatchAqlPacket(
hsa_barrier_and_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) {
@@ -1355,6 +1446,11 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) {
//! That means the app didn't call clFlush/clFinish for very long time.
// Reset the signal for the barrier packet
hsa_signal_silent_store_relaxed(kernarg_pool_signal_[active_chunk_], kInitSignalValueOne);
// dispatch any buffered AQL packets
bool status = dispatchAqlBuffer();
if (!status) {
LogError("dispatch Aql Buffer failed!");
}
// Dispatch a barrier packet into the queue
dispatchBarrierPacket(kBarrierPacketHeader, true, kernarg_pool_signal_[active_chunk_]);
// Get the next chunk
@@ -3171,10 +3267,10 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes,
}
// Dispatch the packet
if (!dispatchAqlPacket(
&dispatchPacket, aqlHeaderWithOrder,
(sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
GPU_FLUSH_ON_EXECUTION)) {
if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder,
(sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
GPU_FLUSH_ON_EXECUTION,
(vcmd != nullptr) ? vcmd->getBufferingState() : false)) {
return false;
}
}
rocclr/device/rocm/rocvirtual.hpp
+4 -1
Dosyayı Görüntüle
@@ -424,13 +424,15 @@ class VirtualGPU : public device::VirtualDevice {
//! Dispatches a barrier with blocking HSA signals
void dispatchBlockingWait();
bool dispatchAqlBuffer();
bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, uint16_t header,
uint16_t rest, bool blocking = true);
uint16_t rest, bool blocking = true, bool buffering = false);
bool dispatchAqlPacket(hsa_barrier_and_packet_t* packet, uint16_t header,
uint16_t rest, bool blocking = true);
template <typename AqlPacket> bool dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header,
uint16_t rest, bool blocking,
size_t size = 1);
void dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal = false,
hsa_signal_t signal = hsa_signal_t{0});
bool dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion,
@@ -564,5 +566,6 @@ class VirtualGPU : public device::VirtualDevice {
int fence_state_; //!< Fence scope
//!< kUnknown/kFlushedToDevice/kFlushedToSystem
bool fence_dirty_; //!< Fence modified flag
std::vector<hsa_kernel_dispatch_packet_t> aqlBuffer_; //!< AQL packet buffer for graphs
};
}
rocclr/platform/command.cpp
+7 -4
Dosyayı Görüntüle
@@ -310,15 +310,18 @@ bool Event::notifyCmdQueue(bool cpu_wait) {
const Event::EventWaitList Event::nullWaitList(0);
// ================================================================================================
Command::Command(HostQueue& queue, cl_command_type type,
const EventWaitList& eventWaitList, uint32_t commandWaitBits, const Event* waitingEvent)
: Event(queue, activity_prof::IsEnabled(activity_prof::OperationId(type)) ||
queue.properties().test(CL_QUEUE_PROFILING_ENABLE) || Agent::shouldPostEventEvents()),
Command::Command(HostQueue& queue, cl_command_type type, const EventWaitList& eventWaitList,
uint32_t commandWaitBits, const Event* waitingEvent)
: Event(queue,
activity_prof::IsEnabled(activity_prof::OperationId(type)) ||
queue.properties().test(CL_QUEUE_PROFILING_ENABLE) ||
Agent::shouldPostEventEvents()),
queue_(&queue),
next_(nullptr),
type_(type),
data_(nullptr),
waitingEvent_(waitingEvent),
buffering_(false),
eventWaitList_(eventWaitList),
commandWaitBits_(commandWaitBits) {
// Retain the commands from the event wait list.
rocclr/platform/command.hpp
+9 -1
Dosyayı Görüntüle
@@ -270,6 +270,7 @@ class Command : public Event {
cl_command_type type_; //!< This command's OpenCL type.
void* data_;
const Event* waitingEvent_; //!< Waiting event associated with the marker
bool buffering_; //!< Flag to enable/disable AQL buffering
protected:
bool cpu_wait_ = false; //!< If true, then the command was issued for CPU/GPU sync
@@ -278,7 +279,7 @@ class Command : public Event {
EventWaitList eventWaitList_;
//! Force await completion of previous command
//! 0x1 - wait before enqueue, 0x2 - wait after, 0x3 - wait both.
//! 0x1 - wait before enqueue, 0x2 - wait after, 0x3 - wait both
uint32_t commandWaitBits_;
//! Construct a new command of the given OpenCL type.
@@ -293,6 +294,7 @@ class Command : public Event {
type_(type),
data_(nullptr),
waitingEvent_(nullptr),
buffering_(false),
eventWaitList_(nullWaitList),
commandWaitBits_(0) {}
@@ -307,6 +309,12 @@ class Command : public Event {
}
public:
//! Returns AQL buffer state
bool getBufferingState() const { return buffering_; }
//! Sets AQL buffer state
void setBufferingState(bool state) { buffering_ = state; }
//! Return the queue this command is enqueued into.
HostQueue* queue() const { return queue_; }
rocclr/utils/flags.hpp
+5 -2
Dosyayı Görüntüle
@@ -287,8 +287,11 @@ release(size_t, HIP_INITIAL_DM_SIZE, 8 * Mi, \
release(bool, HIP_FORCE_DEV_KERNARG, 0, \
"Force device mem for kernel args.") \
release(bool, DEBUG_CLR_USE_SDMA_QUERY, 0, \
"Use SDMA query API to make copy decisions.")
"Use SDMA query API to make copy decisions.") \
release(uint, DEBUG_CLR_GRAPH_MAX_AQL_BUFFER_SIZE, 32, \
"Size of AQL buffering queue") \
release(bool, DEBUG_CLR_GRAPH_ENABLE_BUFFERING, false, \
"Enable/Disable graph AQL buffering")
namespace amd {
extern bool IS_HIP;