Merge branch 'sycl' into review/yang/msan_device_global

Author: AllanZyne

include/ur_api.h

@@ -5188,7 +5188,6 @@ typedef struct ur_kernel_arg_pointer_properties_t {
 ///     - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE
 ///         + `NULL == hKernel`
 ///     - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
-///     - ::UR_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
 UR_APIEXPORT ur_result_t UR_APICALL
 urKernelSetArgPointer(
     ur_kernel_handle_t hKernel,                            ///< [in] handle of the kernel object

scripts/core/kernel.yml

@@ -352,7 +352,6 @@ params:
       desc: "[in][optional] Pointer obtained by USM allocation or virtual memory mapping operation. If null then argument value is considered null."
 returns:
     - $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX
-    - $X_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE
 --- #--------------------------------------------------------------------------
 type: struct
 desc: "Properties for for $xKernelSetExecInfo."

scripts/templates/ldrddi.cpp.mako

@@ -273,11 +273,17 @@ namespace ur_loader
 
         %endif
         %endif
-        ## Before we can re-enable the releases we will need ref-counted object_t.
-        ## See unified-runtime github issue #1784
-        ##%if item['release']:
-        ##// release loader handle
-        ##${item['factory']}.release( ${item['name']} );
+        ## Possibly handle release/retain ref counting - there are no ur_exp-image factories
+        %if 'factory' in item and '_exp_image_' not in item['factory']:
+            %if item['release']:
+            // release loader handle
+            context->factories.${item['factory']}.release( ${item['name']} );
+            %endif
+            %if item['retain']:
+            // increment refcount of handle
+            context->factories.${item['factory']}.retain( ${item['name']} );
+            %endif
+        %endif
         %if not item['release'] and not item['retain'] and not '_native_object_' in item['obj'] or th.make_func_name(n, tags, obj) == 'urPlatformCreateWithNativeHandle':
         try
         {

source/adapters/level_zero/kernel.cpp

@@ -813,12 +813,13 @@ ur_result_t urKernelGetGroupInfo(
                                       (ZeKernelDevice, &kernelProperties));
       if (ZeResult || workGroupProperties.maxGroupSize == 0) {
         return ReturnValue(
-            uint64_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
+            size_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
       }
-      return ReturnValue(workGroupProperties.maxGroupSize);
+      // Specification states this returns a size_t.
+      return ReturnValue(size_t{workGroupProperties.maxGroupSize});
     } else {
       return ReturnValue(
-          uint64_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
+          size_t{Device->ZeDeviceComputeProperties->maxTotalGroupSize});
     }
   }
   case UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE: {
@@ -830,12 +831,12 @@ ur_result_t urKernelGetGroupInfo(
     return ReturnValue(WgSize);
   }
   case UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE:
-    return ReturnValue(uint32_t{Kernel->ZeKernelProperties->localMemSize});
+    return ReturnValue(size_t{Kernel->ZeKernelProperties->localMemSize});
   case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: {
     return ReturnValue(size_t{Device->ZeDeviceProperties->physicalEUSimdWidth});
   }
   case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: {
-    return ReturnValue(uint32_t{Kernel->ZeKernelProperties->privateMemSize});
+    return ReturnValue(size_t{Kernel->ZeKernelProperties->privateMemSize});
   }
   case UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE:
   case UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE:

source/adapters/level_zero/v2/kernel.cpp

@@ -512,10 +512,11 @@ ur_result_t urKernelGetGroupInfo(
         ZE_CALL_NOCHECK(zeKernelGetProperties, (zeDevice, &kernelProperties));
     if (zeResult == ZE_RESULT_SUCCESS &&
         workGroupProperties.maxGroupSize != 0) {
-      return returnValue(workGroupProperties.maxGroupSize);
+      // Specification states this returns a size_t.
+      return returnValue(size_t{workGroupProperties.maxGroupSize});
     }
     return returnValue(
-        uint64_t{hDevice->ZeDeviceComputeProperties->maxTotalGroupSize});
+        size_t{hDevice->ZeDeviceComputeProperties->maxTotalGroupSize});
   }
   case UR_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE: {
     auto props = hKernel->getProperties(hDevice);
@@ -527,15 +528,15 @@ ur_result_t urKernelGetGroupInfo(
   }
   case UR_KERNEL_GROUP_INFO_LOCAL_MEM_SIZE: {
     auto props = hKernel->getProperties(hDevice);
-    return returnValue(uint32_t{props.localMemSize});
+    return returnValue(size_t{props.localMemSize});
   }
   case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: {
     return returnValue(
         size_t{hDevice->ZeDeviceProperties->physicalEUSimdWidth});
   }
   case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: {
     auto props = hKernel->getProperties(hDevice);
-    return returnValue(uint32_t{props.privateMemSize});
+    return returnValue(size_t{props.privateMemSize});
   }
   case UR_KERNEL_GROUP_INFO_COMPILE_MAX_WORK_GROUP_SIZE:
   case UR_KERNEL_GROUP_INFO_COMPILE_MAX_LINEAR_WORK_GROUP_SIZE:

source/adapters/native_cpu/enqueue.cpp

@@ -138,12 +138,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
 #else
   bool isLocalSizeOne =
       ndr.LocalSize[0] == 1 && ndr.LocalSize[1] == 1 && ndr.LocalSize[2] == 1;
-  if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads) {
+  if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads &&
+      !hKernel->hasLocalArgs()) {
     // If the local size is one, we make the assumption that we are running a
     // parallel_for over a sycl::range.
-    // Todo: we could add compiler checks and
-    // kernel properties for this (e.g. check that no barriers are called, no
-    // local memory args).
+    // Todo: we could add more compiler checks and
+    // kernel properties for this (e.g. check that no barriers are called).
 
     // Todo: this assumes that dim 0 is the best dimension over which we want to
     // parallelize

source/adapters/native_cpu/kernel.hpp

@@ -142,7 +142,9 @@ struct ur_kernel_handle_t_ : RefCounted {
     _localMemPoolSize = reqSize;
   }
 
-  // To be called before executing a work group
+  bool hasLocalArgs() const { return !_localArgInfo.empty(); }
+
+  // To be called before executing a work group if local args are present
   void handleLocalArgs(size_t numParallelThread, size_t threadId) {
     // For each local argument we have size*numthreads
     size_t offset = 0;

source/common/ur_singleton.hpp

@@ -11,20 +11,26 @@
 #ifndef UR_SINGLETON_H
 #define UR_SINGLETON_H 1
 
+#include <cassert>
 #include <memory>
 #include <mutex>
 #include <unordered_map>
 
 //////////////////////////////////////////////////////////////////////////
 /// a abstract factory for creation of singleton objects
 template <typename singleton_tn, typename key_tn> class singleton_factory_t {
+    struct entry_t {
+        std::unique_ptr<singleton_tn> ptr;
+        size_t ref_count;
+    };
+
   protected:
     using singleton_t = singleton_tn;
     using key_t = typename std::conditional<std::is_pointer<key_tn>::value,
                                             size_t, key_tn>::type;
 
     using ptr_t = std::unique_ptr<singleton_t>;
-    using map_t = std::unordered_map<key_t, ptr_t>;
+    using map_t = std::unordered_map<key_t, entry_t>;
 
     std::mutex mut; ///< lock for thread-safety
     map_t map;      ///< single instance of singleton for each unique key
@@ -60,16 +66,31 @@ template <typename singleton_tn, typename key_tn> class singleton_factory_t {
         if (map.end() == iter) {
             auto ptr =
                 std::make_unique<singleton_t>(std::forward<Ts>(params)...);
-            iter = map.emplace(key, std::move(ptr)).first;
+            iter = map.emplace(key, entry_t{std::move(ptr), 0}).first;
+        } else {
+            iter->second.ref_count++;
         }
-        return iter->second.get();
+        return iter->second.ptr.get();
+    }
+
+    void retain(key_tn key) {
+        std::lock_guard<std::mutex> lk(mut);
+        auto iter = map.find(getKey(key));
+        assert(iter != map.end());
+        iter->second.ref_count++;
     }
 
     //////////////////////////////////////////////////////////////////////////
     /// once the key is no longer valid, release the singleton
     void release(key_tn key) {
         std::lock_guard<std::mutex> lk(mut);
-        map.erase(getKey(key));
+        auto iter = map.find(getKey(key));
+        assert(iter != map.end());
+        if (iter->second.ref_count == 0) {
+            map.erase(iter);
+        } else {
+            iter->second.ref_count--;
+        }
     }
 
     void clear() {