Merge branch 'main' into review/yang/local_accessor

Author: AllanZyne

include/ur_api.h

@@ -1546,7 +1546,7 @@ typedef enum ur_device_info_t {
                                                                     ///< shared memory access
     UR_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT = 87,                  ///< [::ur_device_usm_access_capability_flags_t] support USM system wide
                                                                     ///< shared memory access
-    UR_DEVICE_INFO_UUID = 88,                                       ///< [char[]] return device UUID
+    UR_DEVICE_INFO_UUID = 88,                                       ///< [uint8_t[]] return device UUID
     UR_DEVICE_INFO_PCI_ADDRESS = 89,                                ///< [char[]] return device PCI address
     UR_DEVICE_INFO_GPU_EU_COUNT = 90,                               ///< [uint32_t] return Intel GPU EU count
     UR_DEVICE_INFO_GPU_EU_SIMD_WIDTH = 91,                          ///< [uint32_t] return Intel GPU EU SIMD width

include/ur_print.hpp

@@ -3596,8 +3596,17 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr, ur_device_info
     } break;
     case UR_DEVICE_INFO_UUID: {
 
-        const char *tptr = (const char *)ptr;
-        printPtr(os, tptr);
+        const uint8_t *tptr = (const uint8_t *)ptr;
+        os << "{";
+        size_t nelems = size / sizeof(uint8_t);
+        for (size_t i = 0; i < nelems; ++i) {
+            if (i != 0) {
+                os << ", ";
+            }
+
+            os << tptr[i];
+        }
+        os << "}";
     } break;
     case UR_DEVICE_INFO_PCI_ADDRESS: {
 

scripts/core/device.yml

@@ -377,7 +377,7 @@ etors:
     - name: USM_SYSTEM_SHARED_SUPPORT
       desc: "[$x_device_usm_access_capability_flags_t] support USM system wide shared memory access"
     - name: UUID
-      desc: "[char[]] return device UUID"
+      desc: "[uint8_t[]] return device UUID"
     - name: PCI_ADDRESS
       desc: "[char[]] return device PCI address"
     - name: GPU_EU_COUNT

scripts/templates/ldrddi.cpp.mako

@@ -58,6 +58,8 @@ namespace ur_loader
         {
             for( auto& platform : context->platforms )
             {
+                if(platform.initStatus != ${X}_RESULT_SUCCESS)
+                    continue;
                 platform.dditable.${n}.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}( 1, &${obj['params'][1]['name']}[adapterIndex], nullptr );
                 try
                 {

scripts/templates/tools-info.hpp.mako

@@ -62,7 +62,11 @@ inline void printPlatformInfos(${x}_platform_handle_t hPlatform, std::string_vie
 inline void printDeviceInfos(${x}_device_handle_t hDevice, std::string_view prefix = "      ") {
 %for etor in obj['etors']:
     std::cout << prefix;
+%if etor['name'] == '$X_DEVICE_INFO_UUID':
+    printDeviceUUID(hDevice, ${etor['name'].replace('$X', X)});
+%else:
     printDeviceInfo<${etor['desc'][1:etor['desc'].find(' ')-1].replace('$x', x)}>(hDevice, ${etor['name'].replace('$X', X)});
+%endif
 %endfor
 }
 %endif

source/adapters/level_zero/adapter.cpp

@@ -19,6 +19,21 @@ ur_adapter_handle_t_ *GlobalAdapter = new ur_adapter_handle_t_();
 ur_adapter_handle_t_ *GlobalAdapter;
 #endif
 
+class ur_legacy_sink : public logger::Sink {
+public:
+  ur_legacy_sink(std::string logger_name = "", bool skip_prefix = true)
+      : Sink(std::move(logger_name), skip_prefix) {
+    this->ostream = &std::cerr;
+  }
+
+  virtual void print([[maybe_unused]] logger::Level level,
+                     const std::string &msg) override {
+    fprintf(stderr, "%s", msg.c_str());
+  }
+
+  ~ur_legacy_sink() = default;
+};
+
 ur_result_t initPlatforms(PlatformVec &platforms) noexcept try {
   uint32_t ZeDriverCount = 0;
   ZE2UR_CALL(zeDriverGet, (&ZeDriverCount, nullptr));
@@ -44,7 +59,18 @@ ur_result_t initPlatforms(PlatformVec &platforms) noexcept try {
 
 ur_result_t adapterStateInit() { return UR_RESULT_SUCCESS; }
 
-ur_adapter_handle_t_::ur_adapter_handle_t_() {
+ur_adapter_handle_t_::ur_adapter_handle_t_()
+    : logger(logger::get_logger("level_zero")) {
+
+  if (UrL0Debug & UR_L0_DEBUG_BASIC) {
+    logger.setLegacySink(std::make_unique<ur_legacy_sink>());
+  };
+
+  if (UrL0Debug & UR_L0_DEBUG_VALIDATION) {
+    setEnvVar("ZE_ENABLE_VALIDATION_LAYER", "1");
+    setEnvVar("ZE_ENABLE_PARAMETER_VALIDATION", "1");
+  }
+
   PlatformCache.Compute = [](Result<PlatformVec> &result) {
     static std::once_flag ZeCallCountInitialized;
     try {
@@ -68,7 +94,7 @@ ur_adapter_handle_t_::ur_adapter_handle_t_() {
       }
 
       if (getenv("SYCL_ENABLE_PCI") != nullptr) {
-        urPrint(
+        logger::warning(
             "WARNING: SYCL_ENABLE_PCI is deprecated and no longer needed.\n");
       }
 
@@ -91,8 +117,9 @@ ur_adapter_handle_t_::ur_adapter_handle_t_() {
       return;
     }
     if (*GlobalAdapter->ZeResult != ZE_RESULT_SUCCESS) {
-      urPrint("zeInit: Level Zero initialization failure\n");
+      logger::error("zeInit: Level Zero initialization failure\n");
       result = ze2urResult(*GlobalAdapter->ZeResult);
+
       return;
     }
 
@@ -157,10 +184,10 @@ ur_result_t adapterStateTeardown() {
     //               zeMemAllocShared = 0     \--->               zeMemFree = 1
     //
     // clang-format on
-
-    fprintf(stderr, "Check balance of create/destroy calls\n");
-    fprintf(stderr,
-            "----------------------------------------------------------\n");
+    // TODO: use logger to print this messages
+    std::cerr << "Check balance of create/destroy calls\n";
+    std::cerr << "----------------------------------------------------------\n";
+    std::stringstream ss;
     for (const auto &Row : CreateDestroySet) {
       int diff = 0;
       for (auto I = Row.begin(); I != Row.end();) {
@@ -171,23 +198,30 @@ ur_result_t adapterStateTeardown() {
         bool Last = (++I == Row.end());
 
         if (Last) {
-          fprintf(stderr, " \\--->");
+          ss << " \\--->";
           diff -= ZeCount;
         } else {
           diff += ZeCount;
           if (!First) {
-            fprintf(stderr, " | \n");
+            ss << " | ";
+            std::cerr << ss.str() << "\n";
+            ss.str("");
+            ss.clear();
           }
         }
-
-        fprintf(stderr, "%30s = %-5d", ZeName, ZeCount);
+        ss << std::setw(30) << std::right << ZeName;
+        ss << " = ";
+        ss << std::setw(5) << std::left << ZeCount;
       }
 
       if (diff) {
         LeakFound = true;
-        fprintf(stderr, " ---> LEAK = %d", diff);
+        ss << " ---> LEAK = " << diff;
       }
-      fprintf(stderr, "\n");
+
+      std::cerr << ss.str() << '\n';
+      ss.str("");
+      ss.clear();
     }
 
     ZeCallCount->clear();

source/adapters/level_zero/adapter.hpp

@@ -8,6 +8,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "logger/ur_logger.hpp"
 #include <atomic>
 #include <mutex>
 #include <optional>
@@ -16,13 +17,16 @@
 
 using PlatformVec = std::vector<std::unique_ptr<ur_platform_handle_t_>>;
 
+class ur_legacy_sink;
+
 struct ur_adapter_handle_t_ {
   ur_adapter_handle_t_();
   std::atomic<uint32_t> RefCount = 0;
   std::mutex Mutex;
 
   std::optional<ze_result_t> ZeResult;
   ZeCache<Result<PlatformVec>> PlatformCache;
+  logger::Logger &logger;
 };
 
 extern ur_adapter_handle_t_ *GlobalAdapter;

source/adapters/level_zero/command_buffer.cpp

@@ -8,6 +8,7 @@
 //
 //===----------------------------------------------------------------------===//
 #include "command_buffer.hpp"
+#include "logger/ur_logger.hpp"
 #include "ur_level_zero.hpp"
 
 /* L0 Command-buffer Extension Doc see:
@@ -140,16 +141,16 @@ ur_result_t calculateKernelWorkDimensions(
         while (GlobalWorkSize3D[I] % GroupSize[I]) {
           --GroupSize[I];
         }
-        if (GlobalWorkSize3D[I] / GroupSize[I] > UINT32_MAX) {
-          urPrint("calculateKernelWorkDimensions: can't find a WG size "
-                  "suitable for global work size > UINT32_MAX\n");
+        if (GlobalWorkSize[I] / GroupSize[I] > UINT32_MAX) {
+          logger::debug("calculateKernelWorkDimensions: can't find a WG size "
+                        "suitable for global work size > UINT32_MAX");
           return UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE;
         }
         WG[I] = GroupSize[I];
       }
-      urPrint("calculateKernelWorkDimensions: using computed WG size = {%d, "
-              "%d, %d}\n",
-              WG[0], WG[1], WG[2]);
+      logger::debug("calculateKernelWorkDimensions: using computed WG "
+                    "size = {{{}, {}, {}}}",
+                    WG[0], WG[1], WG[2]);
     }
   }
 
@@ -177,30 +178,27 @@ ur_result_t calculateKernelWorkDimensions(
     break;
 
   default:
-    urPrint("calculateKernelWorkDimensions: unsupported work_dim\n");
+    logger::error("calculateKernelWorkDimensions: unsupported work_dim");
     return UR_RESULT_ERROR_INVALID_VALUE;
   }
 
   // Error handling for non-uniform group size case
   if (GlobalWorkSize3D[0] !=
       size_t(ZeThreadGroupDimensions.groupCountX) * WG[0]) {
-    urPrint("calculateKernelWorkDimensions: invalid work_dim. The range "
-            "is not a "
-            "multiple of the group size in the 1st dimension\n");
+    logger::error("calculateKernelWorkDimensions: invalid work_dim. The range "
+                  "is not a multiple of the group size in the 1st dimension");
     return UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE;
   }
   if (GlobalWorkSize3D[1] !=
       size_t(ZeThreadGroupDimensions.groupCountY) * WG[1]) {
-    urPrint("calculateKernelWorkDimensions: invalid work_dim. The range "
-            "is not a "
-            "multiple of the group size in the 2nd dimension\n");
+    logger::error("calculateKernelWorkDimensions: invalid work_dim. The range "
+                  "is not a multiple of the group size in the 2nd dimension");
     return UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE;
   }
   if (GlobalWorkSize3D[2] !=
       size_t(ZeThreadGroupDimensions.groupCountZ) * WG[2]) {
-    urPrint("calculateKernelWorkDimensions: invalid work_dim. The range "
-            "is not a "
-            "multiple of the group size in the 3rd dimension\n");
+    logger::error("calculateKernelWorkDimensions: invalid work_dim. The range "
+                  "is not a multiple of the group size in the 3rd dimension");
     return UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE;
   }
 
@@ -268,9 +266,9 @@ static ur_result_t enqueueCommandBufferMemCopyHelper(
              (CommandBuffer->ZeCommandList, Dst, Src, Size,
               LaunchEvent->ZeEvent, ZeEventList.size(), ZeEventList.data()));
 
-  urPrint("calling zeCommandListAppendMemoryCopy() with"
-          "  ZeEvent %#" PRIxPTR "\n",
-          ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
+  logger::debug("calling zeCommandListAppendMemoryCopy() with"
+                "  ZeEvent {}",
+                ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
 
   return UR_RESULT_SUCCESS;
 }
@@ -335,9 +333,9 @@ static ur_result_t enqueueCommandBufferMemCopyRectHelper(
               DstSlicePitch, Src, &ZeSrcRegion, SrcPitch, SrcSlicePitch,
               LaunchEvent->ZeEvent, ZeEventList.size(), ZeEventList.data()));
 
-  urPrint("calling zeCommandListAppendMemoryCopyRegion() with"
-          "  ZeEvent %#" PRIxPTR "\n",
-          ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
+  logger::debug("calling zeCommandListAppendMemoryCopyRegion() with"
+                "  ZeEvent {}",
+                ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
 
   return UR_RESULT_SUCCESS;
 }
@@ -378,9 +376,9 @@ static ur_result_t enqueueCommandBufferFillHelper(
              (CommandBuffer->ZeCommandList, Ptr, Pattern, PatternSize, Size,
               LaunchEvent->ZeEvent, ZeEventList.size(), ZeEventList.data()));
 
-  urPrint("calling zeCommandListAppendMemoryFill() with"
-          "  ZeEvent %#lx\n",
-          ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
+  logger::debug("calling zeCommandListAppendMemoryFill() with"
+                "  ZeEvent {}",
+                ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
 
   return UR_RESULT_SUCCESS;
 }
@@ -519,7 +517,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendKernelLaunchExp(
   if (GlobalWorkOffset != NULL) {
     if (!CommandBuffer->Context->getPlatform()
              ->ZeDriverGlobalOffsetExtensionFound) {
-      urPrint("No global offset extension found on this driver\n");
+      logger::debug("No global offset extension found on this driver");
       return UR_RESULT_ERROR_INVALID_VALUE;
     }
 
@@ -606,9 +604,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendKernelLaunchExp(
               &ZeThreadGroupDimensions, LaunchEvent->ZeEvent,
               ZeEventList.size(), ZeEventList.data()));
 
-  urPrint("calling zeCommandListAppendLaunchKernel() with"
-          "  ZeEvent %#" PRIxPTR "\n",
-          ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
+  logger::debug("calling zeCommandListAppendLaunchKernel() with"
+                "  ZeEvent {}",
+                ur_cast<std::uintptr_t>(LaunchEvent->ZeEvent));
 
   return UR_RESULT_SUCCESS;
 }
@@ -1068,7 +1066,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferUpdateKernelLaunchExp(
   if (NewGlobalWorkOffset && Dim > 0) {
     if (!CommandBuffer->Context->getPlatform()
              ->ZeDriverGlobalOffsetExtensionFound) {
-      urPrint("No global offset extension found on this driver\n");
+      logger::error("No global offset extension found on this driver");
       return UR_RESULT_ERROR_INVALID_VALUE;
     }
     auto MutableGroupOffestDesc =
@@ -1277,8 +1275,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferUpdateKernelLaunchExp(
       // Zero.
       continue;
     } else {
-      urPrint("urCommandBufferUpdateKernelLaunchExp: unsupported name of "
-              "execution attribute.\n");
+      logger::error("urCommandBufferUpdateKernelLaunchExp: unsupported name of "
+                    "execution attribute.");
       return UR_RESULT_ERROR_INVALID_VALUE;
     }
   }