Add JIT timings

src/CHIPBackend.cc

@@ -1044,7 +1044,12 @@ chipstar::Module *chipstar::Device::getOrCreateModule(const SPVModule &SrcMod) {
 
   logDebug("Compile module {}", static_cast<const void *>(&SrcMod));
 
+  auto start = std::chrono::high_resolution_clock::now();
   auto *Module = compile(SrcMod);
+  auto end = std::chrono::high_resolution_clock::now();
+  auto duration =
+      std::chrono::duration_cast<std::chrono::microseconds>(end - start);
+  logInfo("Module compilation took {} microseconds", duration.count());
   if (!Module) { // Probably a compile error.
     logWarn("Compile module returned NULL, probably error");
     return nullptr;

src/backend/Level0/CHIPBackendLevel0.cc

@@ -22,6 +22,7 @@
 
 #include "CHIPBackendLevel0.hh"
 #include "Utils.hh"
+#include <chrono>
 
 // Auto-generated header that lives in <build-dir>/bitcode.
 #include "rtdevlib-modules.h"
@@ -2356,7 +2357,13 @@ static ze_module_handle_t compileIL(ze_context_handle_t ZeCtx,
 
   ze_module_build_log_handle_t Log;
   ze_module_handle_t Object;
+
+  auto start = std::chrono::high_resolution_clock::now();
   zeStatus = zeModuleCreate(ZeCtx, ZeDev, &ModuleDesc, &Object, &Log);
+  auto end = std::chrono::high_resolution_clock::now();
+  std::chrono::duration<double, std::milli> duration = end - start;
+
+  logTrace("zeModuleCreate took {} ms", duration.count());
 
   if (zeStatus != ZE_RESULT_SUCCESS)
     dumpBuildLog(std::move(Log));
@@ -2435,7 +2442,12 @@ void CHIPModuleLevel0::compile(chipstar::Device *ChipDev) {
                                  0, nullptr, nullptr, nullptr};
 
   auto *ChipCtxLz = static_cast<CHIPContextLevel0 *>(ChipDev->getContext());
+  auto start = std::chrono::high_resolution_clock::now();
   ZeModule_ = compileIL(ChipCtxLz->get(), LzDev->get(), ModuleDesc);
+  auto end = std::chrono::high_resolution_clock::now();
+  std::chrono::duration<double, std::milli> duration = end - start;
+
+  logTrace("Module compilation took {} ms", duration.count());
 
   uint32_t KernelCount = 0;
   zeStatus = zeModuleGetKernelNames(ZeModule_, &KernelCount, nullptr);
@@ -2447,6 +2459,8 @@ void CHIPModuleLevel0::compile(chipstar::Device *ChipDev) {
   CHIPERR_CHECK_LOG_AND_THROW_TABLE(zeModuleGetKernelNames);
   for (auto &Kernel : KernelNames)
     logTrace("Kernel {}", Kernel);
+
+  auto kernelCreationStart = std::chrono::high_resolution_clock::now();
   for (uint32_t i = 0; i < KernelCount; i++) {
     std::string HostFName = KernelNames[i];
     logTrace("Registering kernel {}", HostFName);
@@ -2473,13 +2487,26 @@ void CHIPModuleLevel0::compile(chipstar::Device *ChipDev) {
       //       indirectly. This requires kernel code inspection.
       KernelDesc.flags |= ZE_KERNEL_FLAG_FORCE_RESIDENCY;
 
+    auto kernelStart = std::chrono::high_resolution_clock::now();
     zeStatus = zeKernelCreate(ZeModule_, &KernelDesc, &ZeKernel);
+    auto kernelEnd = std::chrono::high_resolution_clock::now();
+    std::chrono::duration<double, std::milli> kernelDuration =
+        kernelEnd - kernelStart;
+
+    logTrace("zeKernelCreate for kernel {} took {} ms", HostFName,
+             kernelDuration.count());
+
     CHIPERR_CHECK_LOG_AND_THROW_TABLE(zeKernelCreate);
     logTrace("LZ KERNEL CREATION via calling zeKernelCreate {} ", zeStatus);
     CHIPKernelLevel0 *ChipZeKernel =
         new CHIPKernelLevel0(ZeKernel, LzDev, HostFName, FuncInfo, this);
     addKernel(ChipZeKernel);
   }
+  auto kernelCreationEnd = std::chrono::high_resolution_clock::now();
+  std::chrono::duration<double, std::milli> kernelCreationDuration =
+      kernelCreationEnd - kernelCreationStart;
+
+  logTrace("Total kernel creation took {} ms", kernelCreationDuration.count());
 }
 
 void CHIPExecItemLevel0::setupAllArgs() {

src/backend/OpenCL/CHIPBackendOpenCL.cc

@@ -851,12 +851,22 @@ static cl::Program compileIL(cl::Context Ctx, CHIPDeviceOpenCL &ChipDev,
                              const void *IL, size_t Length,
                              const std::string &Options = "") {
   cl_int Err;
+  auto start = std::chrono::high_resolution_clock::now();
   cl::Program Prog(clCreateProgramWithIL(Ctx.get(), IL, Length, &Err));
+  auto end = std::chrono::high_resolution_clock::now();
+  auto duration =
+      std::chrono::duration_cast<std::chrono::microseconds>(end - start);
+  logTrace("clCreateProgramWithIL took {} microseconds", duration.count());
   CHIPERR_CHECK_LOG_AND_THROW_TABLE(clCreateProgramWithIL);
 
   cl_device_id DevId = ChipDev.get()->get();
+  auto Start = std::chrono::high_resolution_clock::now();
   Err = clCompileProgram(Prog.get(), 1, &DevId, Options.c_str(), 0, nullptr,
                          nullptr, nullptr, nullptr);
+  auto End = std::chrono::high_resolution_clock::now();
+  auto Duration =
+      std::chrono::duration_cast<std::chrono::milliseconds>(End - Start);
+  logTrace("clCompileProgram took {} ms", Duration.count());
   if (Err != CL_SUCCESS) {
     dumpProgramLog(ChipDev, Prog);
     CHIPERR_LOG_AND_THROW("Compile step failed.", hipErrorInitializationError);
@@ -918,15 +928,24 @@ void CHIPModuleOpenCL::compile(chipstar::Device *ChipDev) {
   std::vector<cl::Program> ClObjects;
   ClObjects.push_back(ClMainObj);
   appendRuntimeObjects(*ChipCtxOcl->get(), *ChipDevOcl, ClObjects);
+  auto start = std::chrono::high_resolution_clock::now();
   Program_ = cl::linkProgram(ClObjects, nullptr, nullptr, nullptr, &Err);
+  auto end = std::chrono::high_resolution_clock::now();
+  auto duration =
+      std::chrono::duration_cast<std::chrono::microseconds>(end - start);
+  logTrace("cl::linkProgram took {} microseconds", duration.count());
   if (Err != CL_SUCCESS) {
     dumpProgramLog(*ChipDevOcl, Program_);
     CHIPERR_LOG_AND_THROW("Device library link step failed.",
                           hipErrorInitializationError);
   }
 
   std::vector<cl::Kernel> Kernels;
+  start = std::chrono::high_resolution_clock::now();
   Err = Program_.createKernels(&Kernels);
+  end = std::chrono::high_resolution_clock::now();
+  duration = std::chrono::duration_cast<std::chrono::microseconds>(end - start);
+  logTrace("clCreateKernelsInProgram took {} microseconds", duration.count());
   CHIPERR_CHECK_LOG_AND_THROW_TABLE(clCreateKernelsInProgram);
 
   logTrace("Kernels in CHIPModuleOpenCL: {} \n", Kernels.size());