AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

SWDEV-371123 - Refactor dispatchBarrierValuePacket

Procházet zdrojové kódy 
Change-Id: I846bd3f60dd8db125e2ca5475e9d25fdd567922b


[ROCm/clr commit: 10ecf8f159]
Tento commit je obsažen v:
Saleel Kudchadker
2022-11-30 18:35:33 -08:00
rodič 7eaabcd8bd
revize f14b550e87
2 změnil soubory, kde provedl 94 přidání a 63 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+87 -61
Zobrazit soubor
@@ -1031,6 +1031,59 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
barrier_packet_.dep_signal[4] = hsa_signal_t{};
}
// ================================================================================================
void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, hsa_signal_t signal,
hsa_signal_value_t value, hsa_signal_value_t mask,
hsa_signal_condition32_t cond, bool skipTs,
hsa_signal_t completionSignal) {
hsa_amd_barrier_value_packet_t barrier_value_packet_ = {0};
uint16_t rest = HSA_AMD_PACKET_TYPE_BARRIER_VALUE;
const uint32_t queueSize = gpu_queue_->size;
const uint32_t queueMask = queueSize - 1;
barrier_value_packet_.signal = signal;
barrier_value_packet_.value = value;
barrier_value_packet_.mask = mask;
barrier_value_packet_.cond = cond;
if (completionSignal.handle == 0) {
// Get active signal for current dispatch if profiling is necessary
barrier_value_packet_.completion_signal =
Barriers().ActiveSignal(kInitSignalValueOne, skipTs ? nullptr : timestamp_);
} else {
// Attach external signal to the packet
barrier_value_packet_.completion_signal = completionSignal;
}
uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1);
while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask);
hsa_amd_barrier_value_packet_t* aql_loc = &(reinterpret_cast<hsa_amd_barrier_value_packet_t*>(
gpu_queue_->base_address))[index & queueMask];
*aql_loc = barrier_value_packet_;
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), packetHeader, rest);
auto cache_state = extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE);
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x ",
"(type=%d, barrier=%d, acquire=%d, release=%d), "
"completion_signal=0x%zx value = 0x%llx mask = 0x%llx cond: %d (GTE: %d EQ: %d NE: %d)",
gpu_queue_, packetHeader, rest,
extractAqlBits(packetHeader, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
cache_state,
barrier_value_packet_.completion_signal,
barrier_value_packet_.value,
barrier_value_packet_.mask,
barrier_value_packet_.cond,
HSA_SIGNAL_CONDITION_GTE, HSA_SIGNAL_CONDITION_EQ, HSA_SIGNAL_CONDITION_NE);
}
// ================================================================================================
void VirtualGPU::ResetQueueStates() {
// Release all transfer buffers on this command queue
@@ -1199,7 +1252,7 @@ bool VirtualGPU::create() {
return false;
}
// Initialize barrier packet.
// Initialize barrier and barrier value packets
memset(&barrier_packet_, 0, sizeof(barrier_packet_));
barrier_packet_.header = kInvalidAql;
@@ -1882,6 +1935,7 @@ void VirtualGPU::submitSvmCopyMemory(amd::SvmCopyMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -1978,6 +2032,7 @@ void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2015,6 +2070,7 @@ void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2054,6 +2110,7 @@ void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2153,6 +2210,7 @@ void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2243,6 +2301,7 @@ void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const void* pattern,
size_t patternSize, const amd::Coord3D& surface,
const amd::Coord3D& origin, const amd::Coord3D& size,
@@ -2306,6 +2365,7 @@ bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const
return true;
}
// ================================================================================================
void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2328,41 +2388,7 @@ void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) {
profilingEnd(cmd);
}
void VirtualGPU::dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t* packet,
hsa_amd_vendor_packet_header_t header) {
assert(packet->completion_signal.handle != 0);
const uint32_t queueSize = gpu_queue_->size;
const uint32_t queueMask = queueSize - 1;
uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1);
while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
amd::Os::yield();
}
hsa_amd_barrier_value_packet_t* aql_loc = &(reinterpret_cast<hsa_amd_barrier_value_packet_t*>(
gpu_queue_->base_address))[index & queueMask];
*aql_loc = *packet;
unsigned int* headerPtr = reinterpret_cast<unsigned int*>(&header);
__atomic_store_n(reinterpret_cast<uint32_t*>(aql_loc), *headerPtr, __ATOMIC_RELEASE);
auto cache_state = extractAqlBits(header.header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE);
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x ",
"(type=%d, barrier=%d, acquire=%d, release=%d), "
"completion_signal=0x%zx value = 0x%llx mask = 0x%llx cond: %d (GTE: %d EQ: %d NE: %d)",
gpu_queue_, header.header, header.AmdFormat,
extractAqlBits(header.header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(header.header, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
extractAqlBits(header.header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE),
cache_state,
packet->completion_signal, packet->value, packet->mask, packet->cond,
HSA_SIGNAL_CONDITION_GTE, HSA_SIGNAL_CONDITION_EQ, HSA_SIGNAL_CONDITION_NE);
}
// ================================================================================================
void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2380,44 +2406,38 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
if (type == ROCCLR_COMMAND_STREAM_WAIT_VALUE) {
if (GPU_STREAMOPS_CP_WAIT) {
hsa_amd_barrier_value_packet_t aqlPacket;
hsa_amd_vendor_packet_header_t header;
hsa_signal_t signal;
uint16_t header = kBarrierVendorPacketHeader;
Buffer* buff = static_cast<Buffer*>(memory);
header.header = kBarrierVendorPacketHeader;
header.AmdFormat = HSA_AMD_PACKET_TYPE_BARRIER_VALUE;
aqlPacket.signal = buff->getSignal();
aqlPacket.completion_signal = Barriers().ActiveSignal();
hsa_signal_t signal = buff->getSignal();
// mask is always applied on value at signal before performing
// the comparision defiend by 'condition'
switch (flags) {
case ROCCLR_STREAM_WAIT_VALUE_GTE:
aqlPacket.value = value;
aqlPacket.mask = mask;
aqlPacket.cond = HSA_SIGNAL_CONDITION_GTE;
case ROCCLR_STREAM_WAIT_VALUE_GTE: {
dispatchBarrierValuePacket(header, signal, value, mask,
HSA_SIGNAL_CONDITION_GTE, true);
break;
case ROCCLR_STREAM_WAIT_VALUE_EQ:
aqlPacket.value = value;
aqlPacket.mask = mask;
aqlPacket.cond = HSA_SIGNAL_CONDITION_EQ;
}
case ROCCLR_STREAM_WAIT_VALUE_EQ: {
dispatchBarrierValuePacket(header, signal, value, mask,
HSA_SIGNAL_CONDITION_EQ, true);
break;
case ROCCLR_STREAM_WAIT_VALUE_AND:
aqlPacket.value = 0;
aqlPacket.mask = (value & mask);
aqlPacket.cond = HSA_SIGNAL_CONDITION_NE;
}
case ROCCLR_STREAM_WAIT_VALUE_AND: {
dispatchBarrierValuePacket(header, signal, 0, (value & mask),
HSA_SIGNAL_CONDITION_NE, true);
break;
case ROCCLR_STREAM_WAIT_VALUE_NOR:
aqlPacket.value = ~value & mask;
aqlPacket.mask = ~value & mask;
aqlPacket.cond = HSA_SIGNAL_CONDITION_NE;
}
case ROCCLR_STREAM_WAIT_VALUE_NOR: {
uint64_t norValue = ~value & mask;
dispatchBarrierValuePacket(header, signal, norValue, norValue,
HSA_SIGNAL_CONDITION_NE, true);
break;
}
default:
ShouldNotReachHere();
break;
}
dispatchBarrierValuePacket(&aqlPacket, header);
}
// Use a blit kernel to perform the wait operation
else {
@@ -2449,6 +2469,7 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2487,6 +2508,7 @@ void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) {
profilingEnd(cmd);
}
// ================================================================================================
void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
@@ -2517,6 +2539,7 @@ void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) {
profilingEnd(vcmd);
}
// ================================================================================================
static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
// Abort on device exceptions.
@@ -2542,6 +2565,7 @@ static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
}
}
// ================================================================================================
bool VirtualGPU::createSchedulerParam()
{
if (nullptr != schedulerParam_) {
@@ -2599,12 +2623,14 @@ bool VirtualGPU::createSchedulerParam()
return false;
}
// ================================================================================================
uint64_t VirtualGPU::getVQVirtualAddress()
{
Memory* vqMem = dev().getRocMemory(virtualQueue_);
return reinterpret_cast<uint64_t>(vqMem->getDeviceMemory());
}
// ================================================================================================
bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
{
uint MinDeviceQueueSize = 16 * 1024;
projects/clr/rocclr/device/rocm/rocvirtual.hpp
+7 -2
Zobrazit soubor
@@ -424,8 +424,13 @@ class VirtualGPU : public device::VirtualDevice {
hsa_signal_t signal = hsa_signal_t{0});
bool dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion,
bool blocking, const hsa_ven_amd_aqlprofile_1_00_pfn_t* extApi);
void dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t* packet,
hsa_amd_vendor_packet_header_t header);
void dispatchBarrierValuePacket(uint16_t packetHeader,
hsa_signal_t signal = hsa_signal_t{0},
hsa_signal_value_t value = 0,
hsa_signal_value_t mask = 0,
hsa_signal_condition32_t cond = HSA_SIGNAL_CONDITION_EQ,
bool skipTs = false,
hsa_signal_t completionSignal = hsa_signal_t{0});
void initializeDispatchPacket(hsa_kernel_dispatch_packet_t* packet,
amd::NDRangeContainer& sizes);