From 740a06d56798c0a30dc3fc579f9b8728ca0d4ff2 Mon Sep 17 00:00:00 2001 From: Karthik Jayaprakash <54370791+kjayapra-amd@users.noreply.github.com> Date: Tue, 25 Nov 2025 19:25:32 -0500 Subject: [PATCH] SWDEV-559267 - Use CLPrint to DevLogPrintf with Log Level - detail debug. (#1160) --- projects/clr/hipamd/CMakeLists.txt | 4 -- projects/clr/hipamd/src/hip_code_object.cpp | 3 +- projects/clr/hipamd/src/hip_context.cpp | 2 +- projects/clr/hipamd/src/hip_memory.cpp | 17 +++++--- projects/clr/hipamd/src/hip_stream.cpp | 2 +- projects/clr/rocclr/device/device.cpp | 39 +++++++++++-------- projects/clr/rocclr/device/devkernel.cpp | 6 ++- projects/clr/rocclr/device/devprogram.cpp | 15 +++---- projects/clr/rocclr/device/pal/paldevice.cpp | 7 ++-- .../clr/rocclr/device/rocm/rocappprofile.cpp | 3 +- projects/clr/rocclr/device/rocm/rocblit.cpp | 3 +- .../clr/rocclr/device/rocm/roccounters.cpp | 4 +- projects/clr/rocclr/device/rocm/rocdevice.cpp | 38 +++++++++++------- projects/clr/rocclr/device/rocm/rockernel.cpp | 23 +++++++---- projects/clr/rocclr/device/rocm/rocmemory.cpp | 7 ++-- projects/clr/rocclr/platform/command.cpp | 5 ++- projects/clr/rocclr/platform/kernel.cpp | 6 ++- projects/clr/rocclr/platform/memory.cpp | 35 +++++++++-------- projects/clr/rocclr/platform/sampler.hpp | 3 +- projects/clr/rocclr/utils/debug.hpp | 7 ---- 20 files changed, 130 insertions(+), 99 deletions(-) diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 044f602450..6d2e03637c 100755 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -301,10 +301,6 @@ if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) set(CMAKE_INSTALL_PREFIX ${HIP_DEFAULT_INSTALL_PREFIX} CACHE PATH "Installation path for HIP" FORCE) endif() -if(DEV_LOG_ENABLE MATCHES "yes") - add_definitions(-DDEV_LOG_ENABLE) -endif() - # Set default install path as "${ROCM_PATH}", can override the path from cmake build. set(CPACK_INSTALL_PREFIX ${HIP_DEFAULT_INSTALL_PREFIX} CACHE PATH "Package Installation path for HIP") diff --git a/projects/clr/hipamd/src/hip_code_object.cpp b/projects/clr/hipamd/src/hip_code_object.cpp index 8c2eef23f2..321f0975ea 100644 --- a/projects/clr/hipamd/src/hip_code_object.cpp +++ b/projects/clr/hipamd/src/hip_code_object.cpp @@ -421,7 +421,8 @@ hipError_t StatCO::registerStatFunction(const void* hostFunction, Function* func amd::ScopedLock lock(sclock_); if (functions_.find(hostFunction) != functions_.end()) { - DevLogPrintfError("hostFunctionPtr: 0x%x already exists", hostFunction); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_API, + "hostFunctionPtr: 0x%x already exists", hostFunction); delete func; } else { functions_.insert(std::make_pair(hostFunction, func)); diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index ca51002293..c672cdfc66 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -282,7 +282,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { } tls.ctxt_stack_.pop(); } else { - DevLogError("Context Stack empty"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_API, "Context Stack empty"); HIP_RETURN(hipErrorInvalidContext); } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 8fce273719..35d9c7517a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1057,7 +1057,7 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder, const cl_channe if (!amd::Image::validateDimensions(devices, imageType, imageWidth, imageHeight, imageDepth, imageArraySize)) { - DevLogError("Image does not have valid dimensions"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Image does not have valid dimensions"); status = hipErrorInvalidValue; return nullptr; } @@ -3431,12 +3431,14 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) { amd_mem_obj = amd::MemObjMap::FindMemObj(dev_ptr); if (amd_mem_obj == nullptr) { - DevLogPrintfError("Memory object for the ptr: 0x%x cannot be null \n", dev_ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Memory object for the ptr: 0x%x cannot be null \n", dev_ptr); HIP_RETURN(hipErrorInvalidValue); } if (!amd_mem_obj->ipcShared()) { - DevLogPrintfError("Memory object for the ptr: 0x%x is not ipcShared \n", dev_ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Memory object for the ptr: 0x%x is not ipcShared \n", dev_ptr); HIP_RETURN(hipErrorInvalidValue); } @@ -3499,7 +3501,8 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void } // getDeviceMemory can fail, hence validate the sanity of the mem obtained if (nullptr == devMem) { - DevLogPrintfError("getDeviceMemory for ptr failed : %p", ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "getDeviceMemory for ptr failed : %p", ptr); HIP_RETURN(hipErrorMemoryAllocation); } @@ -3605,7 +3608,8 @@ hipError_t ihipPointerGetAttributes(void* data, hipPointer_attribute attribute, // getDeviceMemory can fail, hence validate the sanity of the mem obtained if (nullptr == devMem) { - DevLogPrintfError("getDeviceMemory for ptr failed : %p", ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "getDeviceMemory for ptr failed : %p", ptr); return hipErrorMemoryAllocation; } *reinterpret_cast(data) = @@ -3712,7 +3716,8 @@ hipError_t ihipPointerGetAttributes(void* data, hipPointer_attribute attribute, // getDeviceMemory can fail, hence validate the sanity of the mem obtained if (nullptr == devMem) { - DevLogPrintfError("getDeviceMemory for ptr failed : %p", ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "getDeviceMemory for ptr failed : %p", ptr); return hipErrorMemoryAllocation; } *reinterpret_cast(data) = diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 3c4b5e8429..10e6e81273 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -240,7 +240,7 @@ hipStream_t stream_per_thread::get() { hipError_t status = ihipStreamCreate(&m_streams[currDev], hipStreamDefault, hip::Stream::Priority::Normal); if (status != hipSuccess) { - DevLogError("Stream creation failed"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, "Stream creation failed"); } } return m_streams[currDev]; diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp index a1585bb870..45758c1e62 100644 --- a/projects/clr/rocclr/device/device.cpp +++ b/projects/clr/rocclr/device/device.cpp @@ -362,7 +362,7 @@ void MemObjMap::AddMemObj(const void* k, amd::Memory* v) { std::unique_lock lock(AllocatedLock_); auto rval = MemObjMap_.insert({reinterpret_cast(k), v}); if (!rval.second) { - DevLogPrintfError("Memobj map already has an entry for ptr: 0x%x", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Memobj map already has an entry for ptr: 0x%x", reinterpret_cast(k)); } } @@ -435,8 +435,9 @@ void MemObjMap::AddVirtualMemObj(const void* k, amd::Memory* v) { std::unique_lock lock(AllocatedLock_); auto rval = VirtualMemObjMap_.insert({reinterpret_cast(k), v}); if (!rval.second) { - DevLogPrintfError("Virtual Memobj map already has an entry for ptr: 0x%x", - reinterpret_cast(k)); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Virtual Memobj map already has an entry for ptr: 0x%x", + reinterpret_cast(k)); } } @@ -469,7 +470,7 @@ void MemObjMap::AddIpcHandleMemObj(const IpcMemHandle& k, amd::Memory* v) { std::unique_lock lock(AllocatedLock_); auto rval = IpcHandleMemObjMap_.insert({k, v}); if (!rval.second) { - DevLogPrintfError( + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Error adding entry for Memobj 0x%x in IpcHandle map. The handle already exists.", v); } } @@ -629,7 +630,8 @@ bool Device::BlitProgram::create(amd::Device* device, const std::string& extraKe // Create a program with all blit kernels program_ = new Program(*context_, kernels.c_str(), Program::OpenCL_C); if (program_ == nullptr) { - DevLogPrintfError("Program creation for Kernel: %s failed\n", kernels.c_str()); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Program creation for Kernel: %s failed\n", kernels.c_str()); return false; } @@ -651,11 +653,13 @@ bool Device::BlitProgram::create(amd::Device* device, const std::string& extraKe #endif if ((retval = program_->build(devices, opt.c_str(), nullptr, nullptr, GPU_DUMP_BLIT_KERNELS)) != CL_SUCCESS) { - DevLogPrintfError("Build failed for Kernel: %s with error code %d\n", kernels.c_str(), retval); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Build failed for Kernel: %s with error code %d\n", kernels.c_str(), retval); return false; } if (!program_->load()) { - DevLogPrintfError("Could not load the kernels: %s \n", kernels.c_str()); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Could not load the kernels: %s \n", kernels.c_str()); return false; } @@ -684,7 +688,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_INIT, "KFD is not installed \n"); // Disable direct dispatch if ROC initialization wasn't successful AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? false : AMD_DIRECT_DISPATCH; } @@ -1028,7 +1032,8 @@ char* Device::getExtensionString() { bool Device::IpcCreate(void* dev_ptr, size_t* mem_size, char* handle, size_t* mem_offset) const { 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", dev_ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Cannot retrieve amd_mem_obj for dev_ptr: 0x%x", dev_ptr); return false; } @@ -1045,8 +1050,8 @@ bool Device::IpcCreate(void* dev_ptr, size_t* mem_size, char* handle, size_t* me // Check if the dev_ptr is lesser than original dev_ptr if (orig_dev_ptr > dev_ptr) { // If this happens, then revisit FindMemObj logic - DevLogPrintfError("Original dev_ptr: 0x%x cannot be greater than dev_ptr: 0x%x", orig_dev_ptr, - dev_ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Original dev_ptr: 0x%x cannot be greater than dev_ptr: 0x%x", orig_dev_ptr, dev_ptr); return false; } @@ -1138,7 +1143,8 @@ bool Device::GetHandleForAddressRange(void* dev_ptr, size_t size, void* handle) // make sure the memory is allocated. 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", dev_ptr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Cannot retrieve amd_mem_obj for dev_ptr: 0x%x", dev_ptr); return false; } @@ -1337,7 +1343,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { } if (!elfOut_->dumpImage(&image, &imageSize)) { - DevLogError("Dump Image failed \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Dump Image failed \n"); return false; } @@ -1360,7 +1366,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { delete[] image; if (!success) { delete[] outBuf; - DevLogError("Cannot succesfully OCL Encrypt Image"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Cannot succesfully OCL Encrypt Image"); return false; } image = outBuf; @@ -1418,7 +1424,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Cannot Decrypt Image \n"); return false; } @@ -1487,7 +1493,8 @@ bool ClBinary::loadLlvmBinary(std::string& llvmBinary, } } - DevLogPrintfError("Cannot Load LLVM Binary: %s \n", llvmBinary.c_str()); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Cannot Load LLVM Binary: %s \n", llvmBinary.c_str()); return false; } diff --git a/projects/clr/rocclr/device/devkernel.cpp b/projects/clr/rocclr/device/devkernel.cpp index ea9925e648..2a6794b393 100644 --- a/projects/clr/rocclr/device/devkernel.cpp +++ b/projects/clr/rocclr/device/devkernel.cpp @@ -755,7 +755,8 @@ void Kernel::FindLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkSize, 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()); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Cannot get program kernel metadata for %s \n", name().c_str()); return false; } @@ -841,7 +842,8 @@ bool Kernel::GetPrintfStr(std::vector* printfStr) { } if (status != AMD_COMGR_STATUS_SUCCESS) { - DevLogPrintfError("Comgr API failed with status: %d \n", status); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_COMGR, + "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 5589cc1078..170515a832 100644 --- a/projects/clr/rocclr/device/devprogram.cpp +++ b/projects/clr/rocclr/device/devprogram.cpp @@ -221,7 +221,8 @@ static amd_comgr_language_t getCOMGRLanguage(bool isHIP, const amd::option::Opti } } - DevLogPrintfError("Cannot set Language version for %s \n", amdOptions.oVariables->CLStd); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_COMGR, + "Cannot set Language version for %s \n", amdOptions.oVariables->CLStd); return AMD_COMGR_LANGUAGE_NONE; } @@ -969,7 +970,7 @@ bool Program::initBuild(amd::option::Options* options) { } buildLog_.clear(); if (!initClBinary()) { - DevLogError("Init CL Binary failed \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Init CL Binary failed \n"); return false; } @@ -1521,7 +1522,7 @@ bool Program::getCompileOptionsAtLinking(const std::vector& inputProgr bool Program::initClBinary(const char* binaryIn, size_t size, amd::Os::FileDesc fdesc, size_t foffset, std::string uri) { if (!initClBinary()) { - DevLogError("Init CL Binary failed \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Init CL Binary failed \n"); return false; } @@ -1537,7 +1538,7 @@ bool Program::initClBinary(const char* binaryIn, size_t size, amd::Os::FileDesc size_t decryptedSize; if (!clBinary()->decryptElf(binaryIn, size, &decryptedBin, &decryptedSize, &encryptCode)) { - DevLogError("Cannot Decrypt Elf \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Bin is not ELF \n"); return false; } if (decryptedBin != nullptr) { @@ -1549,7 +1550,7 @@ bool Program::initClBinary(const char* binaryIn, size_t size, amd::Os::FileDesc if (!isElf(bin)) { // Invalid binary. delete[] decryptedBin; - DevLogError("Bin is not ELF \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Bin is not ELF \n"); return false; } @@ -1572,7 +1573,7 @@ void Program::addKernel(Kernel* k) { bool Program::setBinary(const char* binaryIn, size_t size, const device::Program* same_dev_prog, amd::Os::FileDesc fdesc, size_t foffset, std::string uri) { if (!initClBinary(binaryIn, size, fdesc, foffset, uri)) { - DevLogError("Init CL Binary failed \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Init CL Binary failed \n"); return false; } @@ -2091,7 +2092,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_COMGR, "Cannot get Symbols from Code Obj \n"); return false; } diff --git a/projects/clr/rocclr/device/pal/paldevice.cpp b/projects/clr/rocclr/device/pal/paldevice.cpp index 6de934fb48..ab476b6a9b 100644 --- a/projects/clr/rocclr/device/pal/paldevice.cpp +++ b/projects/clr/rocclr/device/pal/paldevice.cpp @@ -2737,18 +2737,19 @@ bool Device::createBlitProgram() { std::string opt = "-cl-internal-kernel "; if (auto retval = asm_program->build(devices, opt.c_str(), nullptr, nullptr, false) != CL_SUCCESS) { - DevLogPrintfError("Build failed for trap handler with error code: %d\n", retval); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Build failed for trap handler with error code: %d\n", retval); asm_program->release(); } else { if (asm_program->load()) { trap_handler_ = asm_program; } else { - DevLogError("Could not load the trap handler \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Could not load the trap handler \n"); asm_program->release(); } } } else { - DevLogError("Trap handler creation failed\n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "Trap handler creation failed\n"); } blitProgram_ = new BlitProgram(context_); diff --git a/projects/clr/rocclr/device/rocm/rocappprofile.cpp b/projects/clr/rocclr/device/rocm/rocappprofile.cpp index 8e19f4fe49..c4b5dd46dd 100644 --- a/projects/clr/rocclr/device/rocm/rocappprofile.cpp +++ b/projects/clr/rocclr/device/rocm/rocappprofile.cpp @@ -29,7 +29,8 @@ amd::AppProfile* rocCreateAppProfile() { amd::AppProfile* appProfile = new amd::roc::AppProfile; if ((appProfile == nullptr) || !appProfile->init()) { - DevLogPrintfError("App Profile init failed, appProfile: 0x%x \n", appProfile); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_INIT, + "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 index 51d2672db8..b4c0db7c16 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -2682,7 +2682,8 @@ 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); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Buffer create failed, Buffer: 0x%x \n", amdMemory); amdMemory->release(); return nullptr; } diff --git a/projects/clr/rocclr/device/rocm/roccounters.cpp b/projects/clr/rocclr/device/rocm/roccounters.cpp index 46a9a1bd7c..57aa3a9bf4 100644 --- a/projects/clr/rocclr/device/rocm/roccounters.cpp +++ b/projects/clr/rocclr/device/rocm/roccounters.cpp @@ -579,7 +579,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_AQL, "Cannot Start AQL Profile \n"); return nullptr; } @@ -593,7 +593,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_AQL, "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 5718a7d6d6..38abb065c8 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -935,7 +935,8 @@ bool Sampler::create(const amd::Sampler& owner) { Hsa::sampler_create(dev_.getBackendDevice(), &samplerDescriptor, &hsa_sampler); if (HSA_STATUS_SUCCESS != status) { - DevLogPrintfError("Sampler creation failed with status: %d \n", status); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, + "Sampler creation failed with status: %d \n", status); return false; } @@ -1943,7 +1944,7 @@ device::Memory* Device::createMemory(amd::Memory& owner) const { if (!result) { delete memory; - DevLogError("Cannot Write Image \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Cannot Write Image \n"); return nullptr; } @@ -2075,7 +2076,8 @@ void* Device::hostLock(void* hostMem, size_t size, const MemorySegment memSegmen " deviceMemory = %p, memSegment = %d", pool, size, hostMem, deviceMemory, static_cast(memSegment)); if (status != HSA_STATUS_SUCCESS) { - DevLogPrintfError("Failed to lock memory to pool, failed with hsa_status: %d \n", status); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_LOCK, + "Failed to lock memory to pool, failed with hsa_status: %d \n", status); deviceMemory = nullptr; } return deviceMemory; @@ -2164,8 +2166,9 @@ void* Device::deviceLocalAlloc(size_t size, const AllocationFlags& flags) const : 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_); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Invalid argument, pool_handle: 0x%x , max_alloc: %u \n", + pool.handle, gpuvm_segment_max_alloc_); return nullptr; } @@ -2240,7 +2243,7 @@ void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_ if (flags & CL_MEM_USE_HOST_PTR) { svmPtrUsed = svmPtr; } else { - DevLogPrintfError("Cannot find svm_ptr: 0x%x \n", svmPtr); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Cannot find svm_ptr: 0x%x \n", svmPtr); return nullptr; } } @@ -2911,7 +2914,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, "Cannot get hsa agent info \n"); return nullptr; } auto queue_size = (queue_max_packets < queue_size_hint) ? queue_max_packets : queue_size_hint; @@ -2933,7 +2936,8 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, if (!coop_queue && (cuMask.size() == 0) && (queuePool_[qIndex].size() > 0)) { return getQueueFromPool(qIndex); } - DevLogError("Device::acquireQueue: hsa_queue_create failed!"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, + "Device::acquireQueue: hsa_queue_create failed!"); return nullptr; } } @@ -2942,7 +2946,8 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, if (queue_priority != HSA_AMD_QUEUE_PRIORITY_NORMAL) { hsa_status_t st = Hsa::queue_set_priority(queue, queue_priority); if (st != HSA_STATUS_SUCCESS) { - DevLogError("Device::acquireQueue: hsa_amd_queue_set_priority failed!"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, + "Device::acquireQueue: hsa_amd_queue_set_priority failed!"); Hsa::queue_destroy(queue); return nullptr; } @@ -3017,7 +3022,8 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, hsa_status_t status = Hsa::queue_cu_set_mask(queue, final_mask.size() * 32, final_mask.data()); if (status != HSA_STATUS_SUCCESS) { - DevLogError("Device::acquireQueue: hsa_amd_queue_cu_set_mask failed!"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_QUEUE, + "Device::acquireQueue: hsa_amd_queue_cu_set_mask failed!"); Hsa::queue_destroy(queue); return nullptr; } @@ -3169,7 +3175,8 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, bkendDevice_, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, &hops); if (hsa_status != HSA_STATUS_SUCCESS) { - DevLogPrintfError("Cannot get hops info, hsa failed with status: %d", hsa_status); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Cannot get hops info, hsa failed with status: %d", hsa_status); return false; } @@ -3203,7 +3210,8 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, break; } default: { - DevLogPrintfError("Invalid LinkAttribute: %d ", link_attr.first); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Invalid LinkAttribute: %d ", link_attr.first); return false; } } @@ -3217,7 +3225,8 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, bkendDevice_, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, link_info.data()); if (hsa_status != HSA_STATUS_SUCCESS) { - DevLogPrintfError("Cannot retrieve link info, hsa failed with status: %d", hsa_status); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Cannot retrieve link info, hsa failed with status: %d", hsa_status); return false; } @@ -3254,7 +3263,8 @@ bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool, break; } default: { - DevLogPrintfError("Invalid LinkAttribute: %d ", link_attr.first); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Invalid LinkAttribute: %d ", link_attr.first); return false; } } diff --git a/projects/clr/rocclr/device/rocm/rockernel.cpp b/projects/clr/rocclr/device/rocm/rockernel.cpp index c32a3ba18a..bb0a52d6cd 100644 --- a/projects/clr/rocclr/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/device/rocm/rockernel.cpp @@ -47,7 +47,8 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_get_symbol_by_name(program()->hsaExecutable(), symbolName().c_str(), &agent, &symbol); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError("Cannot Get Symbol : %s, failed with hsa_status: %d \n", symbolName().c_str(), + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Cannot Get Symbol : %s, failed with hsa_status: %d \n", symbolName().c_str(), hsaStatus); return false; } @@ -55,7 +56,8 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernelCodeHandle_); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError(" Cannot Get Symbol Info: %s, failed with hsa_status: %d \n ", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + " Cannot Get Symbol Info: %s, failed with hsa_status: %d \n ", symbolName().c_str(), hsaStatus); return false; } @@ -63,7 +65,8 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_symbol_get_info( symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK, &kernelHasDynamicCallStack_); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError(" Cannot Get Dynamic callstack info, failed with hsa_status: %d \n ", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + " Cannot Get Dynamic callstack info, failed with hsa_status: %d \n ", hsaStatus); return false; } @@ -81,7 +84,8 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_get_symbol_by_name(program()->hsaExecutable(), RuntimeHandle().c_str(), &agent, &kernelSymbol); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError("Cannot get Kernel Symbol by name: %s, failed with hsa_status: %d \n", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Cannot get Kernel Symbol by name: %s, failed with hsa_status: %d \n", RuntimeHandle().c_str(), hsaStatus); return false; } @@ -89,7 +93,7 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_symbol_get_info( kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError( + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, "[ROC][Kernel] Cannot get Kernel Symbol Info, failed with hsa_status: %d \n", hsaStatus); return false; } @@ -97,7 +101,8 @@ bool Kernel::postLoad() { hsaStatus = Hsa::executable_symbol_get_info( kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &variable_address); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError("[ROC][Kernel] Cannot get Kernel Address, failed with hsa_status: %d \n", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "[ROC][Kernel] Cannot get Kernel Address, failed with hsa_status: %d \n", hsaStatus); return false; } @@ -108,7 +113,8 @@ bool Kernel::postLoad() { Hsa::memory_copy(reinterpret_cast(variable_address), &runtime_handle, variable_size); if (hsaStatus != HSA_STATUS_SUCCESS) { - DevLogPrintfError("[ROC][Kernel] HSA Memory copy failed, failed with hsa_status: %d \n", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "[ROC][Kernel] HSA Memory copy failed, failed with hsa_status: %d \n", hsaStatus); return false; } @@ -121,7 +127,8 @@ bool Kernel::postLoad() { if (wavefront_size == 0 && Hsa::agent_get_info(program()->rocDevice().getBackendDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size) != HSA_STATUS_SUCCESS) { - DevLogPrintfError("[ROC][Kernel] Cannot get Wavefront Size, failed with hsa_status: %d \n", + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "[ROC][Kernel] Cannot get Wavefront Size, failed with hsa_status: %d \n", hsaStatus); return false; } diff --git a/projects/clr/rocclr/device/rocm/rocmemory.cpp b/projects/clr/rocclr/device/rocm/rocmemory.cpp index 6ce4f820d6..2b4e59b28e 100644 --- a/projects/clr/rocclr/device/rocm/rocmemory.cpp +++ b/projects/clr/rocclr/device/rocm/rocmemory.cpp @@ -122,7 +122,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", owner()->getSize()); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "Cannot allocate Map memory for size: %u", owner()->getSize()); return nullptr; } } else { @@ -180,7 +181,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Cannot read buffer"); return nullptr; } } @@ -1477,7 +1478,7 @@ void* Image::allocMapTarget(const amd::Coord3D& origin, const amd::Coord3D& regi } else { // Did the map resource allocation fail? if (mapMemory_ == nullptr) { - DevLogError("Could not map target resource"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Could not map target resource"); return nullptr; } } diff --git a/projects/clr/rocclr/platform/command.cpp b/projects/clr/rocclr/platform/command.cpp index 6a268ac77b..9aae99b435 100644 --- a/projects/clr/rocclr/platform/command.cpp +++ b/projects/clr/rocclr/platform/command.cpp @@ -793,8 +793,9 @@ bool CopyMemoryP2PCommand::validateMemory() { 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]); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, + "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 index 2a2332bfed..41c38ed3d4 100644 --- a/projects/clr/rocclr/platform/kernel.cpp +++ b/projects/clr/rocclr/platform/kernel.cpp @@ -58,7 +58,8 @@ 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); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "Kernel Parameter test failed for idx: %d \n", i); return false; } } @@ -327,7 +328,8 @@ address KernelParameters::capture(device::VirtualDevice& vDev, uint64_t lclMemSi 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); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_KERN, + "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 index 3dc21895b8..adf41798cf 100644 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -243,7 +243,7 @@ bool Memory::allocHostMemory(void* initFrom, bool allocHostMem, bool forceCopy) // @note: SVM host memory allocation should be done in the device backend else if (allocHostMem && !isInterop() && !(getMemFlags() & CL_MEM_SVM_FINE_GRAIN_BUFFER)) { if (!hostMemRef_.allocateMemory(size_, context_())) { - DevLogError("Cannot allocate Host Memory Buffer \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Cannot allocate Host Memory Buffer \n"); return false; } @@ -291,7 +291,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_MEM, "Cannot allocate Host Memory \n"); return false; } @@ -654,7 +654,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, + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, + "Invalid Dimenstions, width: %u height: %u depth: %u \n", width, height, depth); return false; } @@ -667,7 +668,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Array is empty \n"); return false; } for (const auto& dev : devices) { @@ -677,7 +678,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ } } if (!sizePass) { - DevLogPrintfError("Cannot allocate image of size: %u \n", arraySize); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Cannot allocate image of size: %u \n", arraySize); return false; } // Fall through... @@ -690,7 +691,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"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Array size cannot be empty \n"); return false; } @@ -701,13 +702,13 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ } } if (!sizePass) { - DevLogPrintfError("Cannot allocate image of size: %u \n", arraySize); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Cannot allocate image of size: %u \n", arraySize); return false; } // Fall through... case CL_MEM_OBJECT_IMAGE1D: if (width == 0) { - DevLogError("Invalid dimension \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid dimension \n"); return false; } for (const auto& dev : devices) { @@ -727,7 +728,7 @@ bool Image::validateDimensions(const std::vector& devices, cl_mem_ break; } - DevLogError("Dimension Validation failed \n"); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Dimension Validation failed \n"); return false; } @@ -830,7 +831,7 @@ bool Image::Format::isValid() const { break; default: { - DevLogPrintfError("Invalid Image format: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid Image format: %u \n", image_channel_data_type); return false; } } @@ -855,7 +856,7 @@ bool Image::Format::isValid() const { break; default: { - DevLogPrintfError("Invalid Luminance: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid Luminance: %u \n", image_channel_data_type); return false; } } @@ -869,7 +870,7 @@ bool Image::Format::isValid() const { break; default: { - DevLogPrintfError("Invalid RGB: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid RGB: %u \n", image_channel_data_type); return false; } } @@ -885,7 +886,7 @@ bool Image::Format::isValid() const { break; default: { - DevLogPrintfError("Invalid BGRA/ARGB: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid BGRA/ARGB: %u \n", image_channel_data_type); return false; } } @@ -899,7 +900,7 @@ bool Image::Format::isValid() const { case CL_UNORM_INT8: break; default: { - DevLogPrintfError("Invalid sBGRA: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid sBGRA: %u \n", image_channel_data_type); return false; } } @@ -911,14 +912,14 @@ bool Image::Format::isValid() const { case CL_FLOAT: break; default: { - DevLogPrintfError("Invalid CL Depth: %u \n", image_channel_data_type); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid CL Depth: %u \n", image_channel_data_type); return false; } } break; default: { - DevLogPrintfError("Invalid image_channel_order: %u \n", image_channel_order); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, "Invalid image_channel_order: %u \n", image_channel_order); return false; } } @@ -1575,4 +1576,4 @@ void IpcBuffer::initDeviceMemory() { memset(deviceMemories_, 0, NumDevicesWithP2P() * sizeof(DeviceMemory)); } -} // namespace amd +} // namespace amd \ No newline at end of file diff --git a/projects/clr/rocclr/platform/sampler.hpp b/projects/clr/rocclr/platform/sampler.hpp index 12fa77639e..3efd245f60 100644 --- a/projects/clr/rocclr/platform/sampler.hpp +++ b/projects/clr/rocclr/platform/sampler.hpp @@ -107,7 +107,8 @@ 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); + ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_RESOURCE, + "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 index 3959e9d767..960722f5a8 100644 --- a/projects/clr/rocclr/utils/debug.hpp +++ b/projects/clr/rocclr/utils/debug.hpp @@ -258,12 +258,5 @@ inline void warning(const char* msg) { amd::report_warning(msg); } ClPrint(amd::LOG_WARNING, amd::LOG_ALWAYS, format, ##__VA_ARGS__) #define LogPrintfInfo(format, ...) ClPrint(amd::LOG_INFO, amd::LOG_ALWAYS, format, ##__VA_ARGS__) -#if (defined(DEBUG) || defined(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_*/