SWDEV-229840 - Improve error messages on ROCCLR Layer.

Change-Id: Iab7d9156cdc206db86385aa05023a0095ed40f92

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;
 }
 

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<std::string>* printfStr) {
       }
 
       if (status != AMD_COMGR_STATUS_SUCCESS) {
+        DevLogPrintfError("Comgr API failed with status: %d \n", status);
         amd::Comgr::destroy_metadata(printfMeta);
         return false;
       }

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<Program*>& 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<const unsigned*>(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<std::string>* 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<std::string> var_names;
 
   if (!getUndefinedVarFromCodeObj(&var_names)) {
+    DevLogError("Cannot get Undefined Var from Code Object \n");
     return false;
   }
 

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;
   }
 

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;
   }

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;
   }
 

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;

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);

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<hsa_amd_ipc_memory_t*>(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;
       }
     }

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;
     }
   }

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;
       }
     }

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

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<amd::Device*>& 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<amd::Device*>& 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<amd::Device*>& 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<amd::Device*>& 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<amd::Device*>& 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<amd::Device*>& 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<amd::Device*>& 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;
 }

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;

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_*/