From 236705c62ff4c8d6a6c6e6349707d5bcb20dc77a Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 13 Apr 2020 22:41:44 -0400 Subject: [PATCH] SWDEV-229840 - Improve error messages on ROCCLR Layer. Change-Id: Iab7d9156cdc206db86385aa05023a0095ed40f92 [ROCm/clr commit: 7458bf9964ed8603ae91059e5d5cab23d5957db1] --- projects/clr/rocclr/device/device.cpp | 11 +++++ projects/clr/rocclr/device/devkernel.cpp | 4 ++ projects/clr/rocclr/device/devprogram.cpp | 16 +++++++ .../clr/rocclr/device/rocm/rocappprofile.cpp | 2 + projects/clr/rocclr/device/rocm/rocblit.cpp | 1 + projects/clr/rocclr/device/rocm/rocblit.hpp | 0 .../clr/rocclr/device/rocm/roccounters.cpp | 2 + projects/clr/rocclr/device/rocm/rocdevice.cpp | 14 ++++-- projects/clr/rocclr/device/rocm/rockernel.cpp | 26 ++++++---- projects/clr/rocclr/device/rocm/rocmemory.cpp | 14 ++++-- projects/clr/rocclr/device/rocm/rocprintf.cpp | 3 +- projects/clr/rocclr/platform/command.cpp | 3 ++ projects/clr/rocclr/platform/kernel.cpp | 2 + projects/clr/rocclr/platform/memory.cpp | 47 ++++++++++++++++--- projects/clr/rocclr/platform/sampler.hpp | 1 + projects/clr/rocclr/utils/debug.hpp | 8 ++++ 16 files changed, 128 insertions(+), 26 deletions(-) mode change 100644 => 100755 projects/clr/rocclr/device/device.cpp mode change 100644 => 100755 projects/clr/rocclr/device/devkernel.cpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/rocappprofile.cpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/rocblit.cpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/rocblit.hpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/roccounters.cpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/rocmemory.cpp mode change 100644 => 100755 projects/clr/rocclr/device/rocm/rocprintf.cpp mode change 100644 => 100755 projects/clr/rocclr/platform/command.cpp mode change 100644 => 100755 projects/clr/rocclr/platform/kernel.cpp mode change 100644 => 100755 projects/clr/rocclr/platform/memory.cpp mode change 100644 => 100755 projects/clr/rocclr/platform/sampler.hpp mode change 100644 => 100755 projects/clr/rocclr/utils/debug.hpp diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp old mode 100644 new mode 100755 index cafff5f3ea..9f7771798e --- a/projects/clr/rocclr/device/device.cpp +++ b/projects/clr/rocclr/device/device.cpp @@ -158,6 +158,7 @@ bool Device::BlitProgram::create(amd::Device* device, const char* extraKernels, } if (CL_SUCCESS != program_->build(devices, opt.c_str(), nullptr, nullptr, GPU_DUMP_BLIT_KERNELS)) { + DevLogPrintfError("Build failed for Kernel: %s \n", kernels.c_str()); return false; } @@ -185,6 +186,7 @@ bool Device::init() { // that KFD is not installed. // Ignore the failure and assume KFD is not installed. // abort(); + DevLogError("KFD is not installed \n"); } ret |= roc::NullDevice::init(); } @@ -629,6 +631,7 @@ bool ClBinary::isRecompilable(std::string& llvmBinary, amd::OclElf::oclElfPlatfo to check it here. */ if (llvmBinary.empty()) { + DevLogError("LLVM Binary string is empty \n"); return false; } @@ -648,6 +651,7 @@ bool ClBinary::isRecompilable(std::string& llvmBinary, amd::OclElf::oclElfPlatfo } } + DevLogPrintfError("LLVM_Binary: %s is not recompilable \n", llvmBinary.c_str()); return false; } @@ -716,6 +720,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { } if (!elfOut_->dumpImage(&image, &imageSize)) { + DevLogError("Dump Image failed \n"); return false; } @@ -734,6 +739,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { delete[] image; if (!success) { delete[] outBuf; + DevLogError("Cannot succesfully OCL Encrypt Image"); return false; } image = outBuf; @@ -784,6 +790,7 @@ bool ClBinary::decryptElf(const char* binaryIn, size_t size, char** decryptBin, int outDataSize = 0; if (!amd::oclDecrypt(binaryIn, (int)size, outBuf, outBufSize, &outDataSize)) { delete[] outBuf; + DevLogError("Cannot Decrypt Image \n"); return false; } @@ -858,6 +865,7 @@ bool ClBinary::loadLlvmBinary(std::string& llvmBinary, } } + DevLogPrintfError("Cannot Load LLVM Binary: %s \n", llvmBinary.c_str()); return false; } @@ -872,6 +880,8 @@ bool ClBinary::loadCompileOptions(std::string& compileOptions) const { } return true; } + DevLogPrintfError("Cannot Load Compilation Options: %s \n", + compileOptions.c_str()); return false; } @@ -886,6 +896,7 @@ bool ClBinary::loadLinkOptions(std::string& linkOptions) const { } return true; } + DevLogPrintfError("Cannot Load Link Options: %s \n", linkOptions.c_str()); return false; } diff --git a/projects/clr/rocclr/device/devkernel.cpp b/projects/clr/rocclr/device/devkernel.cpp old mode 100644 new mode 100755 index f690e334ad..c6f29a81b7 --- a/projects/clr/rocclr/device/devkernel.cpp +++ b/projects/clr/rocclr/device/devkernel.cpp @@ -1126,6 +1126,8 @@ static inline cl_kernel_arg_type_qualifier GetOclTypeQualOCL(const aclArgData* a bool Kernel::GetAttrCodePropMetadata() { amd_comgr_metadata_node_t kernelMetaNode; if (!prog().getKernelMetadata(name(), &kernelMetaNode)) { + DevLogPrintfError("Cannot get program kernel metadata for %s \n", + name().c_str()); return false; } @@ -1184,6 +1186,7 @@ bool Kernel::GetAttrCodePropMetadata() { if (status != AMD_COMGR_STATUS_SUCCESS) { + LogError("Comgr Api failed with Status: \n"); return false; } @@ -1265,6 +1268,7 @@ bool Kernel::GetPrintfStr(std::vector* printfStr) { } if (status != AMD_COMGR_STATUS_SUCCESS) { + DevLogPrintfError("Comgr API failed with status: %d \n", status); amd::Comgr::destroy_metadata(printfMeta); return false; } diff --git a/projects/clr/rocclr/device/devprogram.cpp b/projects/clr/rocclr/device/devprogram.cpp index 7c06590d14..63f83304aa 100755 --- a/projects/clr/rocclr/device/devprogram.cpp +++ b/projects/clr/rocclr/device/devprogram.cpp @@ -373,6 +373,8 @@ bool Program::linkLLVMBitcode(const amd_comgr_data_set_t inputs, amd_comgr_language_t langver; setLangAndTargetStr(amdOptions->oVariables->CLStd, &langver, targetIdent); if (langver == AMD_COMGR_LANGUAGE_NONE) { + DevLogPrintfError("Cannot set Langauge version for %s \n", + amdOptions->oVariables->CLStd); return false; } @@ -430,6 +432,8 @@ bool Program::compileToLLVMBitcode(const amd_comgr_data_set_t compileInputs, amd_comgr_language_t langver; setLangAndTargetStr(amdOptions->oVariables->CLStd, &langver, targetIdent); if (langver == AMD_COMGR_LANGUAGE_NONE) { + DevLogPrintfError("Cannot set Langauge version for %s \n", + amdOptions->oVariables->CLStd); return false; } @@ -1370,6 +1374,7 @@ bool Program::linkImplHSAIL(amd::option::Options* options) { // Call the device layer to setup all available kernels on the actual device if (!setKernels(options, binary, binSize)) { + buildLog_ += "Error: Cannot set kernel \n"; return false; } @@ -1413,6 +1418,7 @@ bool Program::initBuild(amd::option::Options* options) { } buildLog_.clear(); if (!initClBinary()) { + DevLogError("Init CL Binary failed \n"); return false; } @@ -1903,7 +1909,11 @@ bool Program::getCompileOptionsAtLinking(const std::vector& inputProgr bool isSPIRVMagicL(const void* Image, size_t Length) { const unsigned SPRVMagicNumber = 0x07230203; if (Image == nullptr || Length < sizeof(unsigned)) + { + DevLogPrintfError("Invalid Argument, Image: 0x%x Length: %u \n", + Image, Length); return false; + } auto Magic = static_cast(Image); return *Magic == SPRVMagicNumber; } @@ -1911,6 +1921,7 @@ bool isSPIRVMagicL(const void* Image, size_t Length) { // ================================================================================================ bool Program::initClBinary(const char* binaryIn, size_t size) { if (!initClBinary()) { + DevLogError("Init CL Binary failed \n"); return false; } @@ -1982,6 +1993,7 @@ bool Program::initClBinary(const char* binaryIn, size_t size) { } else { size_t decryptedSize; if (!clBinary()->decryptElf(binaryIn, size, &decryptedBin, &decryptedSize, &encryptCode)) { + DevLogError("Cannot Decrypt Elf \n"); return false; } if (decryptedBin != nullptr) { @@ -1995,6 +2007,7 @@ bool Program::initClBinary(const char* binaryIn, size_t size) { if (decryptedBin != nullptr) { delete[] decryptedBin; } + DevLogError("Bin is not ELF \n"); return false; } } @@ -2007,6 +2020,7 @@ bool Program::initClBinary(const char* binaryIn, size_t size) { // ================================================================================================ bool Program::setBinary(const char* binaryIn, size_t size) { if (!initClBinary(binaryIn, size)) { + DevLogError("Init CL Binary failed \n"); return false; } @@ -2623,6 +2637,7 @@ const bool Program::getLoweredNames(std::vector* mangledNames) cons /* Itrate thru global vars */ if (!getSymbolsFromCodeObj(mangledNames, AMD_COMGR_SYMBOL_TYPE_OBJECT)) { + DevLogError("Cannot get Symbols from Code Obj \n"); return false; } @@ -2677,6 +2692,7 @@ bool Program::defineUndefinedVars() { std::vector var_names; if (!getUndefinedVarFromCodeObj(&var_names)) { + DevLogError("Cannot get Undefined Var from Code Object \n"); return false; } diff --git a/projects/clr/rocclr/device/rocm/rocappprofile.cpp b/projects/clr/rocclr/device/rocm/rocappprofile.cpp old mode 100644 new mode 100755 index 4abe0619d0..26579aa294 --- a/projects/clr/rocclr/device/rocm/rocappprofile.cpp +++ b/projects/clr/rocclr/device/rocm/rocappprofile.cpp @@ -31,6 +31,8 @@ amd::AppProfile* rocCreateAppProfile() { amd::AppProfile* appProfile = new roc::AppProfile; if ((appProfile == nullptr) || !appProfile->init()) { + DevLogPrintfError("App Profile init failed, appProfile: 0x%x \n", + appProfile); return nullptr; } diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp old mode 100644 new mode 100755 index 42c54b623f..2ce45b5ab4 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -2142,6 +2142,7 @@ amd::Memory* DmaBlitManager::pinHostMemory(const void* hostMem, size_t pinSize, amdMemory = new (*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, pinAllocSize); amdMemory->setVirtualDevice(&gpu()); if ((amdMemory != nullptr) && !amdMemory->create(tmpHost, SysMem)) { + DevLogPrintfError("Buffer create failed, Buffer: 0x%x \n", amdMemory); amdMemory->release(); return nullptr; } diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp old mode 100644 new mode 100755 diff --git a/projects/clr/rocclr/device/rocm/roccounters.cpp b/projects/clr/rocclr/device/rocm/roccounters.cpp old mode 100644 new mode 100755 index 93a8251ee3..b0856018b3 --- a/projects/clr/rocclr/device/rocm/roccounters.cpp +++ b/projects/clr/rocclr/device/rocm/roccounters.cpp @@ -590,6 +590,7 @@ hsa_ext_amd_aql_pm4_packet_t* PerfCounterProfile::createStartPacket() { // set up the profile aql packets for capturing performance counter if (api_.hsa_ven_amd_aqlprofile_start(&profile_, &prePacket_) != HSA_STATUS_SUCCESS) { + DevLogError("Cannot Start AQL Profile \n"); return nullptr; } @@ -604,6 +605,7 @@ hsa_ext_amd_aql_pm4_packet_t* PerfCounterProfile::createStopPacket() { // set up the profile aql packets for post-capturing performance counter // and create the completion signal if (api_.hsa_ven_amd_aqlprofile_stop(&profile_, &postPacket_) != HSA_STATUS_SUCCESS) { + DevLogError("Cannot Stop AQL Profile \n"); return nullptr; } diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 8dc8a6e690..54d5e53b6d 100755 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -995,6 +995,7 @@ bool Sampler::create(const amd::Sampler& owner) { hsa_status_t status = hsa_ext_sampler_create(dev_.getBackendDevice(), &samplerDescriptor, &hsa_sampler); if (HSA_STATUS_SUCCESS != status) { + DevLogPrintfError("Sampler creation failed with status: %d \n", status); return false; } @@ -1707,6 +1708,7 @@ device::Memory* Device::createMemory(amd::Memory& owner) const { if (!result) { delete memory; + DevLogError("Cannot Write Image \n"); return nullptr; } @@ -1740,6 +1742,8 @@ void* Device::deviceLocalAlloc(size_t size, bool atomics) const { const hsa_amd_memory_pool_t& pool = (atomics)? gpu_fine_grained_segment_ : gpuvm_segment_; if (pool.handle == 0 || gpuvm_segment_max_alloc_ == 0) { + DevLogPrintfError("Invalid argument, pool_handle: 0x%x , max_alloc: %u \n", + pool.handle, gpuvm_segment_max_alloc_); return nullptr; } @@ -1790,19 +1794,19 @@ amd::Memory *Device::IpcAttach(const void* handle, size_t mem_size, unsigned int mem_size, (1 + p2p_agents_.size()), p2p_agents_list_, dev_ptr); if (hsa_status != HSA_STATUS_SUCCESS) { - LogError("[OCL] HSA failed to attach IPC memory"); + LogPrintfError("HSA failed to attach IPC memory with status: %d \n", hsa_status); return nullptr; } /* Create an amd Memory object for the pointer */ amd_mem_obj = new (context()) amd::Buffer(context(), flags, mem_size, *dev_ptr); if (amd_mem_obj == nullptr) { - LogError("[OCL] failed to create a mem object!"); + LogError("failed to create a mem object!"); return nullptr; } if (!amd_mem_obj->create(nullptr)) { - LogError("[OCL] failed to create a svm hidden buffer!"); + LogError("failed to create a svm hidden buffer!"); amd_mem_obj->release(); return nullptr; } @@ -1825,7 +1829,7 @@ void Device::IpcDetach (amd::Memory& memory) const { /*Detach the memory from HSA */ hsa_status = hsa_amd_ipc_memory_detach(dev_ptr); if (hsa_status != HSA_STATUS_SUCCESS) { - LogError("[OCL] HSA failed to detach memory !"); + LogPrintfError("HSA failed to detach memory with status: %d \n", hsa_status); return; } @@ -1859,6 +1863,7 @@ void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_ // Find the existing amd::mem object mem = amd::MemObjMap::FindMemObj(svmPtr); if (nullptr == mem) { + DevLogPrintfError("Cannot find svm_ptr: 0x%x \n", svmPtr); return nullptr; } @@ -1920,6 +1925,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue) { uint32_t queue_max_packets = 0; if (HSA_STATUS_SUCCESS != hsa_agent_get_info(_bkendDevice, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) { + DevLogError("Cannot get hsa agent info \n"); return nullptr; } auto queue_size = (queue_max_packets < queue_size_hint) ? queue_max_packets : queue_size_hint; diff --git a/projects/clr/rocclr/device/rocm/rockernel.cpp b/projects/clr/rocclr/device/rocm/rockernel.cpp index d0b2b3cfb9..d373bb1bcd 100755 --- a/projects/clr/rocclr/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/device/rocm/rockernel.cpp @@ -76,7 +76,7 @@ bool LightningKernel::init() { } if (!SetAvailableSgprVgpr(targetIdent)) { - LogPrintfError("Cannot set available SGPR/VGPR for target Ident:%s \n", targetIdent.c_str()); + DevLogPrintfError("Cannot set available SGPR/VGPR for target Ident:%s \n", targetIdent.c_str()); return false; } @@ -88,14 +88,16 @@ bool LightningKernel::init() { symbolName().c_str(), &agent, &symbol); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogPrintfError("Cannot Get Symbol : %s \n",symbolName().c_str()); + DevLogPrintfError("Cannot Get Symbol : %s, failed with hsa_status: %d \n", + symbolName().c_str(), hsaStatus); return false; } hsaStatus = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernelCodeHandle_); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogPrintfError(" Cannot Get Symbol Info: %s \n ", symbolName().c_str()); + DevLogPrintfError(" Cannot Get Symbol Info: %s, failed with hsa_status: %d \n ", + symbolName().c_str(), hsaStatus); return false; } @@ -113,7 +115,8 @@ bool LightningKernel::init() { RuntimeHandle().c_str(), &agent, &kernelSymbol); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogPrintfError("Cannot get Kernel Symbol by name: %s \n", RuntimeHandle().c_str()); + DevLogPrintfError("Cannot get Kernel Symbol by name: %s, failed with hsa_status: %d \n", + RuntimeHandle().c_str(), hsaStatus); return false; } @@ -121,7 +124,8 @@ bool LightningKernel::init() { HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogError("[ROC][Kernel] Cannot get Kernel Symbol Info \n"); + DevLogPrintfError("[ROC][Kernel] Cannot get Kernel Symbol Info, failed with hsa_status: %d \n", + hsaStatus); return false; } @@ -129,7 +133,8 @@ bool LightningKernel::init() { HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &variable_address); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogError("[ROC][Kernel] Cannot get Kernel Address \n"); + DevLogPrintfError("[ROC][Kernel] Cannot get Kernel Address, failed with hsa_status: %d \n", + hsaStatus); return false; } @@ -142,7 +147,8 @@ bool LightningKernel::init() { &runtime_handle, variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { - LogError("[ROC][Kernel] HSA Memory copy failed \n"); + DevLogPrintfError("[ROC][Kernel] HSA Memory copy failed, failed with hsa_status: %d \n", + hsaStatus); return false; } } @@ -150,7 +156,8 @@ bool LightningKernel::init() { uint32_t wavefront_size = 0; if (hsa_agent_get_info(program()->hsaDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size) != HSA_STATUS_SUCCESS) { - LogError("[ROC][Kernel] Cannot get Wavefront Size \n"); + DevLogPrintfError("[ROC][Kernel] Cannot get Wavefront Size, failed with hsa_status: %d \n", + hsaStatus); return false; } assert(wavefront_size > 0); @@ -229,7 +236,8 @@ bool HSAILKernel::init() { uint32_t wavefront_size = 0; if (HSA_STATUS_SUCCESS != hsa_agent_get_info(program()->hsaDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size)) { - LogPrintfError("Could not get Wave Info Size: %d \n", errorCode); + DevLogPrintfError("Could not get Wave Info Size: %d, failed with hsa_status: %d \n", + errorCode, hsaStatus); return false; } assert(wavefront_size > 0); diff --git a/projects/clr/rocclr/device/rocm/rocmemory.cpp b/projects/clr/rocclr/device/rocm/rocmemory.cpp old mode 100644 new mode 100755 index 477ddecf98..5d28b9f866 --- a/projects/clr/rocclr/device/rocm/rocmemory.cpp +++ b/projects/clr/rocclr/device/rocm/rocmemory.cpp @@ -126,6 +126,8 @@ void* Memory::allocMapTarget(const amd::Coord3D& origin, const amd::Coord3D& reg if (indirectMapCount_ == 1) { if (!allocateMapMemory(owner()->getSize())) { decIndMapCount(); + DevLogPrintfError("Cannot allocate Map memory for size: %u \n", + owner()->getSize()); return nullptr; } } else { @@ -180,6 +182,7 @@ void* Memory::cpuMap(device::VirtualDevice& vDev, uint flags, uint startLayer, u if (!isHostMemDirectAccess() && !IsPersistentDirectMap()) { if (!vDev.blitMgr().readBuffer(*this, mapTarget, amd::Coord3D(0), amd::Coord3D(size()), true)) { decIndMapCount(); + DevLogError("Cannot read buffer \n"); return nullptr; } } @@ -209,7 +212,7 @@ void Memory::IpcCreate(size_t offset, size_t* mem_size, void* handle) const { reinterpret_cast(handle)); if (hsa_status != HSA_STATUS_SUCCESS) { - LogError("[OCL] Failed to create memory for IPC"); + LogPrintfError("Failed to create memory for IPC, failed with hsa_status: %d \n", hsa_status); return; } } @@ -863,6 +866,7 @@ bool Buffer::create() { hsa_status_t status = hsa_amd_memory_lock_to_pool(owner()->getHostMem(), owner()->getSize(), nullptr, 0, pool, 0, &deviceMemory_); if (status != HSA_STATUS_SUCCESS) { + DevLogPrintfError("Failed to lock memory to pool, failed with hsa_status: %d \n", status); deviceMemory_ = nullptr; } } else { @@ -1076,7 +1080,7 @@ bool Image::create() { permission_, &deviceImageInfo_); if (status != HSA_STATUS_SUCCESS) { - LogError("[OCL] Fail to allocate image memory"); + LogPrintfError("[OCL] Fail to allocate image memory, failed with hsa_status: %d \n", status); return false; } @@ -1110,7 +1114,7 @@ bool Image::create() { permission_, &hsaImageObject_); if (status != HSA_STATUS_SUCCESS) { - LogError("[OCL] Fail to allocate image memory"); + LogPrintfError("[OCL] Fail to allocate image memory, failed with hsa_status: %d \n", status); return false; } @@ -1170,7 +1174,7 @@ bool Image::createView(const Memory& parent) { } if (status != HSA_STATUS_SUCCESS) { - LogError("[OCL] Fail to allocate image memory"); + LogPrintfError("[OCL] Fail to allocate image memory with status: %d \n", status); return false; } @@ -1208,7 +1212,7 @@ void* Image::allocMapTarget(const amd::Coord3D& origin, const amd::Coord3D& regi } else { // Did the map resource allocation fail? if (mapMemory_ == nullptr) { - LogError("Could not map target resource"); + DevLogError("Could not map target resource"); return nullptr; } } diff --git a/projects/clr/rocclr/device/rocm/rocprintf.cpp b/projects/clr/rocclr/device/rocm/rocprintf.cpp old mode 100644 new mode 100755 index d3a277138e..b9e2b3dfec --- a/projects/clr/rocclr/device/rocm/rocprintf.cpp +++ b/projects/clr/rocclr/device/rocm/rocprintf.cpp @@ -398,7 +398,8 @@ bool PrintfDbg::init(bool printfEnabled) { // into the corresponding location in the debug buffer hsa_status_t err = hsa_memory_copy(dbgBuffer_, sysMem, 2 * sizeof(uint32_t)); if (err != HSA_STATUS_SUCCESS) { - LogError("\n Can't copy offset and bytes available data to dgbBuffer_!"); + LogPrintfError("\n Can't copy offset and bytes available data to dgbBuffer_," + "failed with status: %d \n!", err); return false; } } diff --git a/projects/clr/rocclr/platform/command.cpp b/projects/clr/rocclr/platform/command.cpp old mode 100644 new mode 100755 index 3d72f1eae6..db0cce08bf --- a/projects/clr/rocclr/platform/command.cpp +++ b/projects/clr/rocclr/platform/command.cpp @@ -571,6 +571,7 @@ bool TransferBufferFileCommand::validateMemory() { staging_[i] = new (memory_->getContext()) Buffer(memory_->getContext(), StagingBufferMemType, StagingBufferSize); if (NULL == staging_[i] || !staging_[i]->create(nullptr)) { + DevLogPrintfError("Staging Create failed, Staging[%d]: 0x%x", i, staging_[i]); return false; } device::Memory* mem = staging_[i]->getDeviceMemory(queue()->device()); @@ -623,6 +624,8 @@ bool CopyMemoryP2PCommand::validateMemory() { for (uint d = 0; d < devices[0]->GlbCtx().devices().size(); ++d) { device::Memory* mem = devices[0]->P2PStage()->getDeviceMemory(*devices[0]->GlbCtx().devices()[d]); if (nullptr == mem) { + DevLogPrintfError("Cannot get P2P stage Device Memory for device: 0x%x \n", + devices[0]->GlbCtx().devices()[d]); return false; } } diff --git a/projects/clr/rocclr/platform/kernel.cpp b/projects/clr/rocclr/platform/kernel.cpp old mode 100644 new mode 100755 index 9d9a187073..2712b35d3c --- a/projects/clr/rocclr/platform/kernel.cpp +++ b/projects/clr/rocclr/platform/kernel.cpp @@ -58,6 +58,7 @@ bool KernelParameters::check() { for (size_t i = 0; i < signature_.numParameters(); ++i) { if (!test(i)) { + DevLogPrintfError("Kernel Parameter test failed for idx: %d \n", i); return false; } } @@ -236,6 +237,7 @@ address KernelParameters::capture(const Device& device, uint64_t lclMemSize, int bool KernelParameters::boundToSvmPointer(const Device& device, const_address capturedParameter, size_t index) const { if (!device.info().svmCapabilities_) { + DevLogPrintfError("The device: 0x%x does not have SVM Capabilities \n", &device); return false; } //! Information about which arguments are SVM pointers is stored after diff --git a/projects/clr/rocclr/platform/memory.cpp b/projects/clr/rocclr/platform/memory.cpp old mode 100644 new mode 100755 index 2ee2bb1986..27b0362367 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -217,6 +217,7 @@ bool Memory::allocHostMemory(void* initFrom, bool allocHostMem, bool forceCopy) // Allocate host memory buffer if needed else if (allocHostMem && !isInterop()) { if (!hostMemRef_.allocateMemory(size_, context_())) { + DevLogError("Cannot allocate Host Memory Buffer \n"); return false; } @@ -263,6 +264,7 @@ bool Memory::create(void* initFrom, bool sysMemAlloc, bool skipAlloc, bool force } // Allocate host memory if requested else if (!allocHostMemory(initFrom, forceAllocHostMem)) { + DevLogError("Cannot allocate Host Memory \n"); return false; } @@ -571,6 +573,8 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ switch (type) { case CL_MEM_OBJECT_IMAGE3D: if ((width == 0) || (height == 0) || (depth < 1)) { + DevLogPrintfError("Invalid Dimenstions, width: %u height: %u depth: %u \n", + width, height, depth); return false; } for (const auto& dev : devices) { @@ -582,6 +586,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ break; case CL_MEM_OBJECT_IMAGE2D_ARRAY: if (arraySize == 0) { + DevLogError("Array is empty \n"); return false; } for (const auto& dev : devices) { @@ -591,11 +596,13 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ } } if (!sizePass) { + DevLogPrintfError("Cannot allocate image of size: %u \n", arraySize); return false; } // Fall through... case CL_MEM_OBJECT_IMAGE2D: if ((width == 0) || (height == 0)) { + DevLogPrintfError("Invalid dimensions width: %u height: %u \n", width, height); return false; } for (const auto dev : devices) { @@ -606,6 +613,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: if (arraySize == 0) { + DevLogError("Array size cannot be empty \n"); return false; } @@ -616,11 +624,13 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ } } if (!sizePass) { + DevLogPrintfError("Cannot allocate image of size: %u \n", arraySize); return false; } // Fall through... case CL_MEM_OBJECT_IMAGE1D: if (width == 0) { + DevLogError("Invalid dimension \n"); return false; } for (const auto& dev : devices) { @@ -631,6 +641,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ break; case CL_MEM_OBJECT_IMAGE1D_BUFFER: if (width == 0) { + DevLogError("Invalid dimension \n"); return false; } for (const auto& dev : devices) { @@ -643,6 +654,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ break; } + DevLogError("Dimension Validation failed \n"); return false; } @@ -745,8 +757,11 @@ bool Image::Format::isValid() const { case CL_FLOAT: break; - default: + default: { + DevLogPrintfError("Invalid Image format: %u \n", + image_channel_data_type); return false; + } } switch (image_channel_order) { @@ -768,8 +783,11 @@ bool Image::Format::isValid() const { case CL_FLOAT: break; - default: + default: { + DevLogPrintfError("Invalid Luminance: %u \n", + image_channel_data_type); return false; + } } break; @@ -780,8 +798,11 @@ bool Image::Format::isValid() const { case CL_UNORM_INT_101010: break; - default: + default: { + DevLogPrintfError("Invalid RGB: %u \n", + image_channel_data_type); return false; + } } break; @@ -794,8 +815,11 @@ bool Image::Format::isValid() const { case CL_UNSIGNED_INT8: break; - default: + default: { + DevLogPrintfError("Invalid BGRA/ARGB: %u \n", + image_channel_data_type); return false; + } } break; @@ -806,8 +830,11 @@ bool Image::Format::isValid() const { switch (image_channel_data_type) { case CL_UNORM_INT8: break; - default: + default: { + DevLogPrintfError("Invalid sBGRA: %u \n", + image_channel_data_type); return false; + } } break; @@ -816,13 +843,19 @@ bool Image::Format::isValid() const { case CL_UNORM_INT16: case CL_FLOAT: break; - default: + default: { + DevLogPrintfError("Invalid CL Depth: %u \n", + image_channel_data_type); return false; + } } break; - default: + default: { + DevLogPrintfError("Invalid image_channel_order: %u \n", + image_channel_order); return false; + } } return true; } diff --git a/projects/clr/rocclr/platform/sampler.hpp b/projects/clr/rocclr/platform/sampler.hpp old mode 100644 new mode 100755 index ed4ade3cdb..1cb969334c --- a/projects/clr/rocclr/platform/sampler.hpp +++ b/projects/clr/rocclr/platform/sampler.hpp @@ -120,6 +120,7 @@ class Sampler : public RuntimeObject { device::Sampler* sampler = NULL; Device* dev = context_.devices()[i]; if (!dev->createSampler(*this, &sampler)) { + DevLogPrintfError("Sampler creation failed for device: 0x%x \n", dev); return false; } deviceSamplers_[dev] = sampler; diff --git a/projects/clr/rocclr/utils/debug.hpp b/projects/clr/rocclr/utils/debug.hpp old mode 100644 new mode 100755 index 2cc0d3b85f..68d85636c9 --- a/projects/clr/rocclr/utils/debug.hpp +++ b/projects/clr/rocclr/utils/debug.hpp @@ -215,4 +215,12 @@ inline void warning(const char* msg) { amd::report_warning(msg); } #define LogPrintfWarning(format, ...) ClPrint(amd::LOG_WARNING, amd::LOG_ALWAYS, format, __VA_ARGS__) #define LogPrintfInfo(format, ...) ClPrint(amd::LOG_INFO, amd::LOG_ALWAYS, format, __VA_ARGS__) +#ifdef DEV_LOG_ENABLE + #define DevLogPrintfError(format, ...) LogPrintfError(format, __VA_ARGS__) + #define DevLogError(msg) LogError(msg) +#else + #define DevLogPrintfError(format, ...) + #define DevLogError(msg) +#endif + #endif /*DEBUG_HPP_*/