From e4b73f08966b53ca2f5e1d6b606fd5a643a56519 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 4 Jul 2014 17:54:13 -0400 Subject: [PATCH] P4 to Git Change 1052436 by gandryey@gera-dev-w7 on 2014/07/04 17:46:53 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 --- rocclr/runtime/device/gpu/gpuvirtual.cpp | 81 +++++++++++++++++++++++- rocclr/runtime/utils/flags.hpp | 2 + 2 files changed, 81 insertions(+), 2 deletions(-) diff --git a/rocclr/runtime/device/gpu/gpuvirtual.cpp b/rocclr/runtime/device/gpu/gpuvirtual.cpp index 006ce88b8d..7058b7e276 100644 --- a/rocclr/runtime/device/gpu/gpuvirtual.cpp +++ b/rocclr/runtime/device/gpu/gpuvirtual.cpp @@ -17,6 +17,7 @@ #include "newcore.h" #include "sc-hsa/Interface/SCHSAInterface.h" #include +#include #ifdef _WIN32 #include @@ -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
( @@ -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(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(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(it->second)->index()) { + child = static_cast(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(gpuDefQueue->blitMgr()).runScheduler( diff --git a/rocclr/runtime/utils/flags.hpp b/rocclr/runtime/utils/flags.hpp index db9cc8c14b..8bea06f379 100644 --- a/rocclr/runtime/utils/flags.hpp +++ b/rocclr/runtime/utils/flags.hpp @@ -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, \