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

P4 to Git Change 1052436 by gandryey@gera-dev-w7 on 2014/07/04 17:46:53

Browse Source 
ECR #304775 - Device enqueuing
	- Added debug print for the generated child kernels. GPU_PRINT_CHILD_KERNEL=N, where N is the number of child kernels for dump.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.cpp#318 edit
... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#205 edit
This commit is contained in:
foreman
2014-07-04 17:54:13 -04:00
parent 3694ab2ce8
commit e4b73f0896
2 changed files with 81 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/runtime/device/gpu/gpuvirtual.cpp
+79 -2
View File
@@ -17,6 +17,7 @@
#include "newcore.h"
#include "sc-hsa/Interface/SCHSAInterface.h"
#include <fstream>
#include <sstream>
#ifdef _WIN32
#include <d3d10_1.h>
@@ -290,7 +291,9 @@ VirtualGPU::createVirtualQueue(uint deviceQueueSize)
allocSize += amd::alignUp(numSlots, 32) / 32;
virtualQueue_ = new Memory(dev(), allocSize);
if ((virtualQueue_ == NULL) || !virtualQueue_->create(Resource::Local)) {
Resource::MemoryType type = (GPU_PRINT_CHILD_KERNEL == 0) ?
Resource::Local : Resource::Remote;
if ((virtualQueue_ == NULL) || !virtualQueue_->create(type)) {
return false;
}
address ptr = reinterpret_cast<address>(
@@ -325,7 +328,9 @@ VirtualGPU::createVirtualQueue(uint deviceQueueSize)
slots[i].wait_list = argStart + dev().info().maxParameterSize_ + 64;
}
// Upload data back to local memory
virtualQueue_->unmap(this);
if (GPU_PRINT_CHILD_KERNEL == 0) {
virtualQueue_->unmap(this);
}
schedParams_ = new Memory(dev(), 64 * Ki);
if ((schedParams_ == NULL) || !schedParams_->create(Resource::RemoteUSWC)) {
@@ -1718,6 +1723,78 @@ VirtualGPU::submitKernelInternalHSA(
//! \todo Remove flush. We start parent earlier.
flushDMA(MainEngine);
if (GPU_PRINT_CHILD_KERNEL != 0) {
waitForEvent(&gpuEvent);
AmdAqlWrap* wraps = (AmdAqlWrap*)(&((AmdVQueueHeader*)gpuDefQueue->virtualQueue_->data())[1]);
uint p = 0;
for (uint i = 0; i < gpuDefQueue->vqHeader_->aql_slot_num; ++i) {
if (wraps[i].state != 0) {
if (p == GPU_PRINT_CHILD_KERNEL) {
break;
}
p++;
std::stringstream print;
print.flags(std::ios::right | std::ios_base::hex | std::ios_base::uppercase);
print << "Slot#: " << i << "\n";
print << "\tenqueue_flags: " << wraps[i].enqueue_flags << "\n";
print << "\tcommand_id: " << wraps[i].command_id << "\n";
print << "\tchild_counter: " << wraps[i].child_counter << "\n";
print << "\tcompletion: " << wraps[i].completion << "\n";
print << "\tparent_wrap: " << wraps[i].parent_wrap << "\n";
print << "\twait_list: " << wraps[i].wait_list << "\n";
print << "\twait_num: " << wraps[i].wait_num << "\n";
print << "WorkGroupSize[ " << wraps[i].aql.workgroup_size[0] << ", ";
print << wraps[i].aql.workgroup_size[1] << ", ";
print << wraps[i].aql.workgroup_size[2] << "]\n";
print << "GridSize[ " << wraps[i].aql.grid_size[0] << ", ";
print << wraps[i].aql.grid_size[1] << ", ";
print << wraps[i].aql.grid_size[2] << "]\n";
uint64_t* kernels = (uint64_t*)(
const_cast<Memory*>(hsaKernel.prog().kernelTable())->map(this));
uint j;
for (j = 0; j < hsaKernel.prog().kernels().size(); ++j) {
if (kernels[j] == wraps[i].aql.kernel_object_address) {
break;
}
}
const_cast<Memory*>(hsaKernel.prog().kernelTable())->unmap(this);
HSAILKernel* child = NULL;
for (auto it = hsaKernel.prog().kernels().begin();
it != hsaKernel.prog().kernels().end(); ++it) {
if (j == static_cast<HSAILKernel*>(it->second)->index()) {
child = static_cast<HSAILKernel*>(it->second);
}
}
if (child == NULL) {
printf("Error: couldn't find child kernel!\n");
continue;
}
uint offsArg = wraps[i].aql.kernel_arg_address -
gpuDefQueue->virtualQueue_->vmAddress();
address argum = gpuDefQueue->virtualQueue_->data() + offsArg;
print << "Kernel: " << child->name() << "\n";
static const char* Names[HSAILKernel::ExtraArguments] = {
"Offset0: ", "Offset1: ","Offset2: ","PrintfBuf: ", "VqueuePtr: ", "AqlWarap: "};
for (j = 0; j < HSAILKernel::ExtraArguments; ++j) {
print << "\t" << Names[j] << *(size_t*)argum;
print << "\n";
argum += sizeof(size_t);
}
for (j = 0; j < child->numArguments(); ++j) {
print << "\t" << child->argument(j)->name_ << ": ";
for (int s = child->argument(j)->size_ - 1; s >= 0; --s) {
print << (uint32_t)(argum[s]);
}
argum += child->argument(j)->size_;
print << "\n";
}
printf("%s", print.str().c_str());
}
}
}
// Get the global loop start before the scheduler
mcaddr loopStart = gpuDefQueue->virtualQueueDispatcherStart();
static_cast<KernelBlitManager&>(gpuDefQueue->blitMgr()).runScheduler(
rocclr/runtime/utils/flags.hpp
+2
View File
@@ -146,6 +146,8 @@ release(bool, GPU_HSAIL_ENABLE, false, \
"Enable HSAIL on dGPU stack (requires CI+ HW)") \
release(bool, GPU_ASSUME_ALIASES, false, \
"Assume memory aliases in the compilation process") \
release(uint, GPU_PRINT_CHILD_KERNEL, 0, \
"Prints the specified number of the child kernels") \
release(bool, GPU_DIRECT_SRD, true, \
"Use indirect SRD access in HSAIL") \
release(bool, AMD_DEPTH_MSAA_INTEROP, false, \