Merge branch 'sycl' into get_native_br

Author: maximdimakov

clang/test/SemaSYCL/intel-fpga-global-const.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
 
 // Test that checks global constant variable (which allows the redeclaration) since
 // IntelFPGAConstVar is one of the subjects listed for [[intel::max_replicates()]] attribute.

sycl/include/CL/sycl/handler.hpp

@@ -731,14 +731,37 @@ class __SYCL_EXPORT handler {
 
     // FIXME Remove the ESIMD check once rounding of execution range works well
     // with ESIMD compilation flow.
+    // Range rounding can be disabled by the user.
+    // Range rounding is not done on the host device.
     // Range rounding is supported only for newer SYCL standards.
-    // Range rounding can also be disabled by the user.
 #if !defined(__SYCL_EXPLICIT_SIMD__) &&                                        \
     !defined(SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING) &&                      \
-    SYCL_LANGUAGE_VERSION >= 202001
-    // The work group size preferred by this device.
-    // A reasonable choice for rounding up the range is 32.
-    constexpr size_t GoodLocalSizeX = 32;
+    !defined(DPCPP_HOST_DEVICE_OPENMP) &&                                      \
+    !defined(DPCPP_HOST_DEVICE_PERF_NATIVE) && SYCL_LANGUAGE_VERSION >= 202001
+    // Range should be a multiple of this for reasonable performance.
+    size_t MinFactorX = 16;
+    // Range should be a multiple of this for improved performance.
+    size_t GoodFactorX = 32;
+    // Range should be at least this to make rounding worthwhile.
+    size_t MinRangeX = 1024;
+
+    // Parse optional parameters of this form:
+    // MinRound:PreferredRound:MinRange
+    char *RoundParams = getenv("SYCL_PARALLEL_FOR_RANGE_ROUNDING_PARAMS");
+    if (RoundParams != nullptr) {
+      std::string Params(RoundParams);
+      size_t Pos = Params.find(':');
+      if (Pos != std::string::npos) {
+        MinFactorX = std::stoi(Params.substr(0, Pos));
+        Params.erase(0, Pos + 1);
+        Pos = Params.find(':');
+        if (Pos != std::string::npos) {
+          GoodFactorX = std::stoi(Params.substr(0, Pos));
+          Params.erase(0, Pos + 1);
+          MinRangeX = std::stoi(Params);
+        }
+      }
+    }
 
     // Disable the rounding-up optimizations under these conditions:
     // 1. The env var SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING is set.
@@ -769,15 +792,16 @@ class __SYCL_EXPORT handler {
         (KI::callsThisItem());
 
     // Perform range rounding if rounding-up is enabled
+    // and there are sufficient work-items to need rounding
     // and the user-specified range is not a multiple of a "good" value.
-    if (!DisableRounding && NumWorkItems[0] % GoodLocalSizeX != 0) {
+    if (!DisableRounding && (NumWorkItems[0] >= MinRangeX) &&
+        (NumWorkItems[0] % MinFactorX != 0)) {
       // It is sufficient to round up just the first dimension.
       // Multiplying the rounded-up value of the first dimension
       // by the values of the remaining dimensions (if any)
       // will yield a rounded-up value for the total range.
       size_t NewValX =
-          ((NumWorkItems[0] + GoodLocalSizeX - 1) / GoodLocalSizeX) *
-          GoodLocalSizeX;
+          ((NumWorkItems[0] + GoodFactorX - 1) / GoodFactorX) * GoodFactorX;
       using NameWT = typename detail::get_kernel_wrapper_name_t<NameT>::name;
       if (getenv("SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE") != nullptr)
         std::cout << "parallel_for range adjusted from " << NumWorkItems[0]

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1802,19 +1802,26 @@ pi_result piDevicePartition(pi_device Device,
 
   PI_ASSERT(Device, PI_INVALID_DEVICE);
 
+  // Check if Device was already partitioned into the same or bigger size
+  // before. If so, we can return immediately without searching the global
+  // device cache. Note that L0 driver always returns the same handles in the
+  // same order for the given number of sub-devices.
+  if (OutDevices && NumDevices <= Device->SubDevices.size()) {
+    for (uint32_t I = 0; I < NumDevices; I++) {
+      OutDevices[I] = Device->SubDevices[I];
+      // reusing the same pi_device needs to increment the reference count
+      piDeviceRetain(OutDevices[I]);
+    }
+    if (OutNumDevices)
+      *OutNumDevices = NumDevices;
+    return PI_SUCCESS;
+  }
+
   // Get the number of subdevices available.
   // TODO: maybe add interface to create the specified # of subdevices.
   uint32_t Count = 0;
   ZE_CALL(zeDeviceGetSubDevices(Device->ZeDevice, &Count, nullptr));
 
-  // Check that the requested/allocated # of sub-devices is the same
-  // as was reported by the above call.
-  // TODO: we may want to support smaller/larger # devices too.
-  if (Count != NumDevices) {
-    zePrint("piDevicePartition: unsupported # of sub-devices requested\n");
-    return PI_INVALID_OPERATION;
-  }
-
   if (OutNumDevices) {
     *OutNumDevices = Count;
   }
@@ -1825,17 +1832,29 @@ pi_result piDevicePartition(pi_device Device,
   }
 
   try {
+    pi_platform Platform = Device->Platform;
     auto ZeSubdevices = new ze_device_handle_t[Count];
     ZE_CALL(zeDeviceGetSubDevices(Device->ZeDevice, &Count, ZeSubdevices));
 
     // Wrap the Level Zero sub-devices into PI sub-devices, and write them out.
     for (uint32_t I = 0; I < Count; ++I) {
-      OutDevices[I] = new _pi_device(ZeSubdevices[I], Device->Platform,
-                                     true /* isSubDevice */);
-      pi_result Result = OutDevices[I]->initialize();
-      if (Result != PI_SUCCESS) {
-        delete[] ZeSubdevices;
-        return Result;
+      pi_device Dev = Platform->getDeviceFromNativeHandle(ZeSubdevices[I]);
+      if (Dev) {
+        OutDevices[I] = Dev;
+        // reusing the same pi_device needs to increment the reference count
+        piDeviceRetain(OutDevices[I]);
+      } else {
+        std::unique_ptr<_pi_device> PiSubDevice(
+            new _pi_device(ZeSubdevices[I], Platform));
+        pi_result Result = PiSubDevice->initialize();
+        if (Result != PI_SUCCESS) {
+          delete[] ZeSubdevices;
+          return Result;
+        }
+        OutDevices[I] = PiSubDevice.get();
+        Platform->PiDevicesCache.push_back(std::move(PiSubDevice));
+        // save pointers to sub-devices for quick retrieval in the future.
+        Device->SubDevices.push_back(Dev);
       }
     }
     delete[] ZeSubdevices;
@@ -1911,29 +1930,25 @@ pi_result piextDeviceCreateWithNativeHandle(pi_native_handle NativeHandle,
   PI_ASSERT(Device, PI_INVALID_DEVICE);
   PI_ASSERT(NativeHandle, PI_INVALID_VALUE);
   PI_ASSERT(Platform, PI_INVALID_PLATFORM);
-
-  std::lock_guard<std::mutex> Lock(Platform->PiDevicesCacheMutex);
-  pi_result Res = populateDeviceCacheIfNeeded(Platform);
-  if (Res != PI_SUCCESS) {
-    return Res;
+  {
+    std::lock_guard<std::mutex> Lock(Platform->PiDevicesCacheMutex);
+    pi_result Res = populateDeviceCacheIfNeeded(Platform);
+    if (Res != PI_SUCCESS) {
+      return Res;
+    }
   }
-
   auto ZeDevice = pi_cast<ze_device_handle_t>(NativeHandle);
 
   // The SYCL spec requires that the set of devices must remain fixed for the
   // duration of the application's execution. We assume that we found all of the
   // Level Zero devices when we initialized the device cache, so the
   // "NativeHandle" must already be in the cache. If it is not, this must not be
   // a valid Level Zero device.
-  for (const std::unique_ptr<_pi_device> &CachedDevice :
-       Platform->PiDevicesCache) {
-    if (CachedDevice->ZeDevice == ZeDevice) {
-      *Device = CachedDevice.get();
-      return PI_SUCCESS;
-    }
-  }
-
-  return PI_INVALID_VALUE;
+  pi_device Dev = Platform->getDeviceFromNativeHandle(ZeDevice);
+  if (Dev == nullptr)
+    return PI_INVALID_VALUE;
+  *Device = Dev;
+  return PI_SUCCESS;
 }
 
 pi_result piContextCreate(const pi_context_properties *Properties,

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -156,6 +156,12 @@ struct _pi_device : _pi_object {
   // Level Zero device handle.
   ze_device_handle_t ZeDevice;
 
+  // Keep the subdevices that are partitioned from this pi_device for reuse
+  // The order of sub-devices in this vector is repeated from the
+  // ze_device_handle_t array that are returned from zeDeviceGetSubDevices()
+  // call, which will always return sub-devices in the fixed same order.
+  std::vector<pi_device> SubDevices;
+
   // PI platform to which this device belongs.
   pi_platform Platform;
 

sycl/source/detail/device_image_impl.hpp

@@ -47,7 +47,7 @@ class device_image_impl {
 
     // Collect kernel names for the image
     pi_device_binary DevBin =
-        const_cast<pi_device_binary>(&BinImage->getRawData());
+        const_cast<pi_device_binary>(&MBinImage->getRawData());
     for (_pi_offload_entry EntriesIt = DevBin->EntriesBegin;
          EntriesIt != DevBin->EntriesEnd; ++EntriesIt) {
 

sycl/test/basic_tests/kernel_bundle/kernel_bundle_api.cpp renamed to sycl/test/on-device/basic_tests/kernel_bundle/kernel_bundle_api.cpp

@@ -13,6 +13,11 @@ class Kernel2Name;
 
 int main() {
   sycl::queue Q;
+
+  // No support for host device so far.
+  if (Q.is_host())
+    return 0;
+
   const sycl::context Ctx = Q.get_context();
   const sycl::device Dev = Q.get_device();