diff --git a/projects/clr/hipamd/src/hip_clang.cpp b/projects/clr/hipamd/src/hip_clang.cpp
index 15a96d298a..6067edd289 100644
--- a/projects/clr/hipamd/src/hip_clang.cpp
+++ b/projects/clr/hipamd/src/hip_clang.cpp
@@ -86,6 +86,7 @@ __hipRegisterFatBinary(const void* data)
 
     std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
       desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
+    tprintf(DB_FB, "Found bundle for %s\n", target.c_str());
 
     for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
       hsa_agent_t agent = g_allAgents[deviceId + 1];
@@ -110,10 +111,25 @@ __hipRegisterFatBinary(const void* data)
 
       if (module->executable.handle) {
         modules->at(deviceId) = module;
+        tprintf(DB_FB, "Loaded code object for %s\n", name);
+      } else {
+        fprintf(stderr, "Failed to load code object for %s\n", name);
+        abort();
       }
     }
   }
 
+  for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
+    hsa_agent_t agent = g_allAgents[deviceId + 1];
+
+    char name[64] = {};
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
+    if (!(*modules)[deviceId]) {
+      fprintf(stderr, "No device code bundle for %s\n", name);
+      abort();
+    }
+  }
+
   tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules);
   return modules;
 }
@@ -132,13 +148,18 @@ extern "C" void __hipRegisterFunction(
   dim3*        gridDim,
   int*         wSize)
 {
+  HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName);
   std::vector<hipFunction_t> functions{g_deviceCnt};
 
+  assert(modules && modules->size() >= g_deviceCnt);
   for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
     hipFunction_t function;
     if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) {
       functions[deviceId] = function;
     }
+    else {
+      tprintf(DB_FB, "missing kernel %s for device %d\n", deviceName, deviceId);
+    }
   }
 
   g_functions.insert(std::make_pair(hostFunction, std::move(functions)));
@@ -180,6 +201,7 @@ hipError_t hipSetupArgument(
   size_t size,
   size_t offset)
 {
+  HIP_INIT_API(arg, size, offset);
   auto ctx = ihipGetTlsDefaultCtx();
   LockedAccessor_CtxCrit_t crit(ctx->criticalData());
   auto& arguments = crit->_execStack.top()._arguments;
@@ -194,6 +216,7 @@ hipError_t hipSetupArgument(
 
 hipError_t hipLaunchByPtr(const void *hostFunction)
 {
+  HIP_INIT_API(hostFunction);
   ihipExec_t exec;
   {
     auto ctx = ihipGetTlsDefaultCtx();
@@ -213,20 +236,26 @@ hipError_t hipLaunchByPtr(const void *hostFunction)
     deviceId = 0;
   }
 
+  hipError_t e = hipSuccess;
   decltype(g_functions)::iterator it;
-  if ((it = g_functions.find(hostFunction)) == g_functions.end())
-    return hipErrorUnknown;
+  if ((it = g_functions.find(hostFunction)) == g_functions.end()) {
+    e = hipErrorUnknown;
+    fprintf(stderr, "kernel %p not found!\n", hostFunction);
+    abort();
+  } else {
+    size_t size = exec._arguments.size();
+    void *extra[] = {
+        HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
+        HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+        HIP_LAUNCH_PARAM_END
+      };
 
-  size_t size = exec._arguments.size();
-  void *extra[] = {
-      HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
-      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-      HIP_LAUNCH_PARAM_END
-    };
+    e = hipModuleLaunchKernel(it->second[deviceId],
+      exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
+      exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
+      exec._sharedMem, exec._hStream, nullptr, extra);
+  }
 
-  return hipModuleLaunchKernel(it->second[deviceId],
-    exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
-    exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
-    exec._sharedMem, exec._hStream, nullptr, extra);
+  return ihipLogStatus(e);
 }