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

SWDEV-384557 - Fix engine status query

Browse Source 
- Maintain a map of SDMA engine# to stream allocated following a greedy
approach
- Anything past that will query SDMA engine status always and go with a
SDMA or Blit copy path

Change-Id: Ibfaed7f951ab84d80cb0430596a4d11b5aec9202


[ROCm/clr commit: 5865c642d4]
This commit is contained in:
Saleel Kudchadker
2023-04-10 15:19:57 -07:00
committed by Rahul Garg
parent 74405e021d
commit 8ae30e036d
5 changed files with 161 additions and 68 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/device.hpp
+1
View File
@@ -611,6 +611,7 @@ struct Info : public amd::EmbeddedObject {
//! Number of VGPRs per SIMD
uint32_t vgprsPerSimd_;
uint32_t vgprAllocGranularity_;
uint32_t numSDMAengines_; //!< Number of available SDMA engines
};
//! Device settings
projects/clr/rocclr/device/rocm/rocblit.cpp
+75 -56
View File
@@ -32,9 +32,7 @@ DmaBlitManager::DmaBlitManager(VirtualGPU& gpu, Setup setup)
: HostBlitManager(gpu, setup),
MinSizeForPinnedTransfer(dev().settings().pinnedMinXferSize_),
completeOperation_(false),
context_(nullptr),
lastCopyMask_(0),
lastUsedCopyEngine_(HwQueueEngine::Unknown) {}
context_(nullptr) {}
inline void DmaBlitManager::synchronize() const {
if (syncOperation_) {
@@ -118,8 +116,9 @@ bool DmaBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost,
if (pinned != nullptr) {
// Get device memory for this virtual device
Memory* dstMemory = dev().getRocMemory(pinned);
if (!hsaCopy(gpuMem(srcMemory), *dstMemory, srcPin, dst, copySizePin)) {
const KernelBlitManager *kb = dynamic_cast<const KernelBlitManager*>(this);
if (!kb->copyBuffer(gpuMem(srcMemory), *dstMemory, srcPin, dst,
copySizePin)) {
LogWarning("DmaBlitManager::readBuffer failed a pinned copy!");
gpu().addPinnedMem(pinned);
break;
@@ -287,8 +286,9 @@ bool DmaBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory,
if (pinned != nullptr) {
// Get device memory for this virtual device
Memory* srcMemory = dev().getRocMemory(pinned);
if (!hsaCopy(*srcMemory, gpuMem(dstMemory), src, dstPin, copySizePin)) {
const KernelBlitManager *kb = dynamic_cast<const KernelBlitManager*>(this);
if (!kb->copyBuffer(*srcMemory, gpuMem(dstMemory), src, dstPin,
copySizePin)) {
LogWarning("DmaBlitManager::writeBuffer failed a pinned copy!");
gpu().addPinnedMem(pinned);
break;
@@ -677,63 +677,74 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
srcAgent = dstAgent = dev().getBackendDevice();
}
uint32_t copyMask = 0;
uint32_t freeEngineMask = 0;
HwQueueEngine engine = HwQueueEngine::Unknown;
if ((srcAgent.handle == dev().getCpuAgent().handle) &&
(dstAgent.handle != dev().getCpuAgent().handle)) {
engine = HwQueueEngine::SdmaWrite;
copyMask = dev().fetchSDMAMask(this, false);
} else if ((srcAgent.handle != dev().getCpuAgent().handle) &&
(dstAgent.handle == dev().getCpuAgent().handle)) {
engine = HwQueueEngine::SdmaRead;
copyMask = dev().fetchSDMAMask(this, true);
}
auto wait_events = gpu().Barriers().WaitingSignal(engine);
hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
uint32_t freeEngineMask = 0;
uint32_t copyMask = lastCopyMask_;
if (engine != HwQueueEngine::Unknown) {
if (copyMask == Device::kSkipQueryStatus) {
// Do not query engine status or take copy_on_engine path
status = HSA_STATUS_ERROR_OUT_OF_RESOURCES;
}
if ((engine != lastUsedCopyEngine_)) {
// Check SDMA engine status
status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, freemask %x",
status, freeEngineMask);
// Return a mask with the rightmost bit set
copyMask = freeEngineMask - (freeEngineMask & (freeEngineMask - 1));
}
if (copyMask == 0) {
// Check SDMA engine status
status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, free_engine mask %x",
status, freeEngineMask);
// Return a mask with the rightmost bit set
copyMask = freeEngineMask - (freeEngineMask & (freeEngineMask - 1));
}
if (copyMask != 0 && engine != HwQueueEngine::Unknown) {
// Copy on the first available free engine if ROCr returns a valid mask
hsa_amd_sdma_engine_id_t copyEngine = static_cast<hsa_amd_sdma_engine_id_t>(copyMask);
if (copyMask != 0 && status == HSA_STATUS_SUCCESS) {
auto wait_events = gpu().Barriers().WaitingSignal(engine);
hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
// Copy on the first available free engine if ROCr returns a valid mask
hsa_amd_sdma_engine_id_t copyEngine = static_cast<hsa_amd_sdma_engine_id_t>(copyMask);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"HSA Async Copy on copy_engine=%x, dst=0x%zx, src=0x%zx, "
"size=%ld, wait_event=0x%zx, completion_signal=0x%zx", copyEngine,
dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"HSA Async Copy on copy_engine=%x, dst=0x%zx, src=0x%zx, "
"size=%ld, wait_event=0x%zx, completion_signal=0x%zx", copyEngine,
dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent,
size[0], wait_events.size(),
wait_events.data(), active, copyEngine, false);
status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent,
size[0], wait_events.size(),
wait_events.data(), active, copyEngine, false);
if (status != HSA_STATUS_SUCCESS) {
gpu().Barriers().ResetCurrentSignal();
}
}
} else {
// Force copy with BLIT in ROCr. Forcing agents to the GPU device causes ROCr to take
// blit path internally
srcAgent = dstAgent = dev().getBackendDevice();
auto wait_events = gpu().Barriers().WaitingSignal(engine);
hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp());
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"HSA Async Blit Copy dst=0x%zx, src=0x%zx, size=%ld, wait_event=0x%zx, "
"completion_signal=0x%zx",
dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
"HSA Async Blit Copy dst=0x%zx, src=0x%zx, size=%ld, wait_event=0x%zx, "
"completion_signal=0x%zx",
dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent,
size[0], wait_events.size(), wait_events.data(), active);
if (status != HSA_STATUS_SUCCESS) {
gpu().Barriers().ResetCurrentSignal();
}
}
if (status == HSA_STATUS_SUCCESS) {
lastUsedCopyEngine_ = engine;
lastCopyMask_ = copyMask;
gpu().addSystemScope();
} else {
gpu().Barriers().ResetCurrentSignal();
LogPrintfError("HSA copy from host to device failed with code %d", status);
LogPrintfError("HSA copy failed with code %d, falling to Blit copy", status);
}
return (status == HSA_STATUS_SUCCESS);
@@ -855,6 +866,9 @@ KernelBlitManager::~KernelBlitManager() {
kernels_[i]->release();
}
}
dev().resetSDMAMask(this);
if (nullptr != program_) {
program_->release();
}
@@ -2264,9 +2278,26 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
#endif
#endif
if (setup_.disableHwlCopyBuffer_ ||
(!srcMemory.isHostMemDirectAccess() && !dstMemory.isHostMemDirectAccess() &&
!(p2p || asan) && !ipcShared)) {
bool useShaderCopyPath = setup_.disableHwlCopyBuffer_ ||
(!srcMemory.isHostMemDirectAccess() &&
!dstMemory.isHostMemDirectAccess() &&
!(p2p || asan) && !ipcShared);
if (!useShaderCopyPath) {
if (amd::IS_HIP) {
// Update the command type for ROC profiler
if (srcMemory.isHostMemDirectAccess()) {
gpu().SetCopyCommandType(CL_COMMAND_WRITE_BUFFER);
}
if (dstMemory.isHostMemDirectAccess()) {
gpu().SetCopyCommandType(CL_COMMAND_READ_BUFFER);
}
}
result = DmaBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire,
copyMetadata);
}
if (!result) {
uint blitType = BlitCopyBuffer;
size_t dim = 1;
size_t globalWorkOffset[3] = {0, 0, 0};
@@ -2338,18 +2369,6 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
address parameters = captureArguments(kernels_[blitType]);
result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr);
releaseArguments(parameters);
} else {
if (amd::IS_HIP) {
// Update the command type for ROC profiler
if (srcMemory.isHostMemDirectAccess()) {
gpu().SetCopyCommandType(CL_COMMAND_WRITE_BUFFER);
}
if (dstMemory.isHostMemDirectAccess()) {
gpu().SetCopyCommandType(CL_COMMAND_READ_BUFFER);
}
}
result = DmaBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire,
copyMetadata);
}
synchronize();
projects/clr/rocclr/device/rocm/rocblit.hpp
-2
View File
@@ -237,8 +237,6 @@ class DmaBlitManager : public device::HostBlitManager {
const size_t MinSizeForPinnedTransfer;
bool completeOperation_; //!< DMA blit manager must complete operation
amd::Context* context_; //!< A dummy context
mutable uint32_t lastCopyMask_; //!< Last used copy mask
mutable HwQueueEngine lastUsedCopyEngine_; //!< Last used copy engine
private:
//! Disable copy constructor
projects/clr/rocclr/device/rocm/rocdevice.cpp
+75 -10
View File
@@ -240,12 +240,12 @@ Device::~Device() {
hsa_queue_t* queue = qIter->first;
auto& qInfo = qIter->second;
if (qInfo.hostcallBuffer_) {
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "deleting hostcall buffer %p for hardware queue %p",
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hostcall buffer %p for hardware queue %p",
qInfo.hostcallBuffer_, qIter->first);
disableHostcalls(qInfo.hostcallBuffer_);
context().svmFree(qInfo.hostcallBuffer_);
}
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "deleting hardware queue %p with refCount 0", queue);
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0", queue);
qIter = it.erase(qIter);
hsa_queue_destroy(queue);
}
@@ -1183,6 +1183,17 @@ bool Device::populateOCLDeviceConstants() {
return false;
}
if (HSA_STATUS_SUCCESS !=
hsa_agent_get_info(bkendDevice_,
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG),
&info_.numSDMAengines_)) {
return false;
}
for (uint32_t i = 0; i < info_.numSDMAengines_; i++) {
engineAssignMap_[1 << i] = 0;
}
setupCpuAgent();
checkAtomicSupport();
@@ -1200,7 +1211,8 @@ bool Device::populateOCLDeviceConstants() {
hsa_status_t err;
// Can another GPU (agent) have access to the current GPU memory pool (gpuvm_segment_)?
hsa_amd_memory_pool_access_t access;
err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access);
err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_,
HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access);
if (err != HSA_STATUS_SUCCESS) {
continue;
}
@@ -1214,7 +1226,7 @@ bool Device::populateOCLDeviceConstants() {
}
}
/* Keep track of all P2P Agents in a Array including current device handle for IPC */
// Keep track of all P2P Agents in a Array including current device handle for IPC
p2p_agents_list_ = new hsa_agent_t[1 + p2p_agents_.size()];
p2p_agents_list_[0] = getBackendDevice();
for (size_t agent_idx = 0; agent_idx < p2p_agents_.size(); ++agent_idx) {
@@ -1229,6 +1241,20 @@ bool Device::populateOCLDeviceConstants() {
}
assert(group_segment_size > 0);
// Find SDMA read mask
if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getCpuAgent(), getBackendDevice(),
&maxSdmaReadMask)) {
return false;
}
assert(maxSdmaReadMask > 0 && "No SDMA engines available for Read");
// Find SDMA write mask
if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getBackendDevice(), getCpuAgent(),
&maxSdmaWriteMask)) {
return false;
}
assert(maxSdmaWriteMask > 0 && "No SDMA engines available for Write");
info_.localMemSizePerCU_ = group_segment_size;
info_.localMemSize_ = group_segment_size;
@@ -1631,7 +1657,7 @@ bool Device::populateOCLDeviceConstants() {
LogError("HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS query failed.");
}
ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, xnack: %d, direct host access: %d\n",
ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, xnack: %d, direct host access: %d",
info_.hmmSupported_, info_.hmmCpuMemoryAccessible_, info_.hmmDirectHostAccess_);
info_.globalCUMask_ = {};
@@ -2185,7 +2211,7 @@ bool Device::IpcCreate(void* dev_ptr, size_t* mem_size, void* handle, size_t* me
amd::Memory* amd_mem_obj = amd::MemObjMap::FindMemObj(dev_ptr);
if (amd_mem_obj == nullptr) {
DevLogPrintfError("Cannot retrieve amd_mem_obj for dev_ptr: 0x%x \n", dev_ptr);
DevLogPrintfError("Cannot retrieve amd_mem_obj for dev_ptr: 0x%x", dev_ptr);
return false;
}
@@ -2788,7 +2814,7 @@ hsa_queue_t* Device::getQueueFromPool(const uint qIndex) {
for (auto it = queuePool_[qIndex].begin(); it != queuePool_[qIndex].end(); it++) {
if (it->second.refCount == 0) {
it->second.refCount++;
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "selected queue refCount: %p (%d)\n", it->first,
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "selected queue refCount: %p (%d)", it->first,
it->second.refCount);
return it->first;
}
@@ -2930,7 +2956,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue,
for (int i = mask.size() - 1; i >= 0; i--) {
ss << mask[i];
}
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "setting CU mask 0x%s for hardware queue %p",
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Setting CU mask 0x%s for hardware queue %p",
ss.str().c_str(), queue);
hsa_status_t status = hsa_amd_queue_cu_set_mask(queue, mask.size() * 32, mask.data());
@@ -2960,7 +2986,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue,
assert(result.second && "QueueInfo already exists");
auto &qInfo = result.first->second;
qInfo.refCount = 1;
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)\n", result.first->first,
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)", result.first->first,
result.first->second.refCount);
return queue;
}
@@ -2972,7 +2998,7 @@ void Device::releaseQueue(hsa_queue_t* queue, const std::vector<uint32_t>& cuMas
auto &qInfo = qIter->second;
assert(qInfo.refCount > 0);
qInfo.refCount--;
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)\n", qIter->first,
ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)", qIter->first,
qIter->second.refCount);
}
}
@@ -3292,6 +3318,45 @@ void Device::HiddenHeapAlloc(const VirtualGPU& gpu) {
std::call_once(heap_initialized_, HeapAllocZeroOut);
}
// ================================================================================================
uint32_t Device::fetchSDMAMask(const device::BlitManager* handle, bool readEngine) const {
uint32_t engine = 0;
{
amd::ScopedLock lock(vgpusAccess());
for (auto it = engineAssignMap_.rbegin(); it != engineAssignMap_.rend(); ++it) {
// If blitManager handle is in the map return the engine ID else
// add to the map
if (it->second == handle) {
engine = it->first;
break;
} else if (it->second == 0) {
it->second = handle;
engine = it->first;
break;
}
}
}
uint32_t mask = (readEngine ? maxSdmaReadMask : maxSdmaWriteMask) & engine;
if (engine != 0 && mask == 0 ) {
return kSkipQueryStatus;
} else {
return mask;
}
}
// ================================================================================================
void Device::resetSDMAMask(const device::BlitManager* handle) const {
amd::ScopedLock lock(vgpusAccess());
for (auto& it : engineAssignMap_) {
if (it.second == handle) {
it.second = 0;
break;
}
}
}
// ================================================================================================
ProfilingSignal::~ProfilingSignal() {
if (signal_.handle != 0) {
projects/clr/rocclr/device/rocm/rocdevice.hpp
+10
View File
@@ -551,6 +551,9 @@ class Device : public NullDevice {
//! Allocates hidden heap for device memory allocations
void HiddenHeapAlloc(const VirtualGPU& gpu);
uint32_t fetchSDMAMask(const device::BlitManager* handle, bool readEngine = true) const;
void resetSDMAMask(const device::BlitManager* handle) const ;
private:
bool create();
@@ -619,7 +622,14 @@ class Device : public NullDevice {
//! Pool of HSA queues with custom CU masks
std::vector<std::map<hsa_queue_t*, QueueInfo>> queueWithCUMaskPool_;
//! Read and Write mask for device<->host
uint32_t maxSdmaReadMask;
uint32_t maxSdmaWriteMask;
//! Map of SDMA engineId<->stream
mutable std::map<uint32_t, const device::BlitManager*> engineAssignMap_;
public:
constexpr static uint32_t kSkipQueryStatus = 1 << 31;
std::atomic<uint> numOfVgpus_; //!< Virtual gpu unique index
//! enum for keeping the total and available queue priorities