[SYCL] Do not build device code for sub-devices

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -2124,8 +2124,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
     }
   }
   case PI_DEVICE_INFO_PARENT_DEVICE:
-    // TODO: all Level Zero devices are parent ?
-    return ReturnValue(pi_device{0});
+    return ReturnValue(Device->RootDevice);
   case PI_DEVICE_INFO_PLATFORM:
     return ReturnValue(Device->Platform);
   case PI_DEVICE_INFO_VENDOR_ID:

sycl/source/detail/device_impl.cpp

@@ -53,13 +53,11 @@ device_impl::device_impl(pi_native_handle InteropDeviceHandle,
   Plugin.call<PiApiKind::piDeviceGetInfo>(
       MDevice, PI_DEVICE_INFO_TYPE, sizeof(RT::PiDeviceType), &MType, nullptr);
 
-  RT::PiDevice parent = nullptr;
   // TODO catch an exception and put it to list of asynchronous exceptions
   Plugin.call<PiApiKind::piDeviceGetInfo>(MDevice, PI_DEVICE_INFO_PARENT_DEVICE,
-                                          sizeof(RT::PiDevice), &parent,
+                                          sizeof(RT::PiDevice), &MRootDevice,
                                           nullptr);
 
-  MIsRootDevice = (nullptr == parent);
   if (!InteroperabilityConstructor) {
     // TODO catch an exception and put it to list of asynchronous exceptions
     // Interoperability Constructor already calls DeviceRetain in

sycl/source/detail/device_impl.hpp

@@ -225,12 +225,14 @@ class device_impl {
 
   bool isAssertFailSupported() const;
 
+  bool isRootDevice() const { return MRootDevice == nullptr; }
+
 private:
   explicit device_impl(pi_native_handle InteropDevice, RT::PiDevice Device,
                        PlatformImplPtr Platform, const plugin &Plugin);
   RT::PiDevice MDevice = 0;
   RT::PiDeviceType MType;
-  bool MIsRootDevice = false;
+  RT::PiDevice MRootDevice = nullptr;
   bool MIsHostDevice;
   PlatformImplPtr MPlatform;
   bool MIsAssertFailSupported = false;

sycl/source/detail/persistent_device_code_cache.cpp

@@ -78,10 +78,13 @@ void PersistentDeviceCodeCache::putItemToDisc(
     const SerializedObj &SpecConsts, const std::string &BuildOptionsString,
     const RT::PiProgram &NativePrg) {
 
+  if (!isImageCached(Img))
+    return;
+
   std::string DirName =
       getCacheItemPath(Device, Img, SpecConsts, BuildOptionsString);
 
-  if (!isImageCached(Img) || DirName.empty())
+  if (DirName.empty())
     return;
 
   auto Plugin = detail::getSyclObjImpl(Device)->getPlugin();
@@ -137,10 +140,13 @@ std::vector<std::vector<char>> PersistentDeviceCodeCache::getItemFromDisc(
     const device &Device, const RTDeviceBinaryImage &Img,
     const SerializedObj &SpecConsts, const std::string &BuildOptionsString) {
 
+  if (!isImageCached(Img))
+    return {};
+
   std::string Path =
       getCacheItemPath(Device, Img, SpecConsts, BuildOptionsString);
 
-  if (!isImageCached(Img) || Path.empty() || !OSUtil::isPathPresent(Path))
+  if (Path.empty() || !OSUtil::isPathPresent(Path))
     return {};
 
   int i = 0;

sycl/source/detail/program_manager/program_manager.cpp

@@ -470,10 +470,16 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(
   if (Prg)
     Prg->stableSerializeSpecConstRegistry(SpecConsts);
 
-  auto BuildF = [this, &M, &KSId, &ContextImpl, &DeviceImpl, Prg, &CompileOpts,
+  // Use root device image to avoid building for the same architecture.
+  DeviceImplPtr RootDev = DeviceImpl;
+  while (!RootDev->isRootDevice())
+    RootDev = detail::getSyclObjImpl(
+        RootDev->get_info<info::device::parent_device>());
+
+  auto BuildF = [this, &M, &KSId, &ContextImpl, &RootDev, Prg, &CompileOpts,
                  &LinkOpts, &JITCompilationIsRequired, SpecConsts] {
     auto Context = createSyclObjFromImpl<context>(ContextImpl);
-    auto Device = createSyclObjFromImpl<device>(DeviceImpl);
+    auto Device = createSyclObjFromImpl<device>(RootDev);
 
     const RTDeviceBinaryImage &Img =
         getDeviceImage(M, KSId, Context, Device, JITCompilationIsRequired);
@@ -523,7 +529,7 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(
     return BuiltProgram.release();
   };
 
-  const RT::PiDevice PiDevice = DeviceImpl->getHandleRef();
+  const RT::PiDevice PiDevice = RootDev->getHandleRef();
 
   auto BuildResult = getOrBuild<PiProgramT, compile_program_error>(
       Cache,
@@ -560,16 +566,21 @@ ProgramManager::getOrCreateKernel(OSModuleHandle M,
     Prg->stableSerializeSpecConstRegistry(SpecConsts);
   }
   applyOptionsFromEnvironment(CompileOpts, LinkOpts);
-  const RT::PiDevice PiDevice = DeviceImpl->getHandleRef();
+
+  // Use root device image to avoid building for the same architecture.
+  DeviceImplPtr D = DeviceImpl;
+  while (!D->isRootDevice())
+    D = detail::getSyclObjImpl(D->get_info<info::device::parent_device>());
+
+  const RT::PiDevice PiDevice = D->getHandleRef();
 
   auto key = std::make_tuple(std::move(SpecConsts), M, PiDevice,
                              CompileOpts + LinkOpts, KernelName);
   auto ret_tuple = Cache.tryToGetKernelFast(key);
   if (std::get<0>(ret_tuple))
     return ret_tuple;
 
-  RT::PiProgram Program =
-      getBuiltPIProgram(M, ContextImpl, DeviceImpl, KernelName, Prg);
+  RT::PiProgram Program = getBuiltPIProgram(M, ContextImpl, D, KernelName, Prg);
 
   auto AcquireF = [](KernelProgramCache &Cache) {
     return Cache.acquireKernelsPerProgramCache();
@@ -830,8 +841,13 @@ ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId,
   for (unsigned I = 0; I < Imgs.size(); I++)
     RawImgs[I] = const_cast<pi_device_binary>(&Imgs[I]->getRawData());
 
+  // Use root device image to avoid building for the same architecture.
+  device RootDevice = Device;
+  while (!getSyclObjImpl(RootDevice)->isRootDevice())
+    RootDevice = Device.get_info<info::device::parent_device>();
+
   Ctx->getPlugin().call<PiApiKind::piextDeviceSelectBinary>(
-      getSyclObjImpl(Device)->getHandleRef(), RawImgs.data(),
+      getSyclObjImpl(RootDevice)->getHandleRef(), RawImgs.data(),
       (cl_uint)RawImgs.size(), &ImgInd);
 
   if (JITCompilationIsRequired) {

sycl/source/detail/program_manager/program_manager.hpp

@@ -107,8 +107,8 @@ class ProgramManager {
                        SerializedObj SpecConsts);
   /// Builds or retrieves from cache a program defining the kernel with given
   /// name.
-  /// \param M idenfies the OS module the kernel comes from (multiple OS modules
-  ///          may have kernels with the same name)
+  /// \param M identifies the OS module the kernel comes from (multiple OS
+  ///        modules may have kernels with the same name)
   /// \param Context the context to build the program with
   /// \param Device the device for which the program is built
   /// \param KernelName the kernel's name
@@ -152,7 +152,7 @@ class ProgramManager {
   /// \param NativePrg the native program, target for spec constant setting; if
   ///        not null then overrides the native program in Prg
   /// \param Img A source of the information about which constants need
-  ///        setting and symboling->integer spec constnant ID mapping. If not
+  ///        setting and symboling->integer spec constant ID mapping. If not
   ///        null, overrides native program->binary image binding maintained by
   ///        the program manager.
   void flushSpecConstants(const program_impl &Prg,

sycl/unittests/helpers/CommonRedefinitions.hpp

@@ -6,7 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <CL/sycl.hpp>
 #include <helpers/PiImage.hpp>
 #include <helpers/PiMock.hpp>
 

sycl/unittests/helpers/PiImage.hpp

@@ -8,7 +8,6 @@
 
 #pragma once
 
-#include <CL/sycl.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/pi.hpp>
 #include <detail/platform_impl.hpp>

sycl/unittests/helpers/PiMock.hpp

@@ -27,9 +27,11 @@
 
 #pragma once
 
-#include <CL/sycl.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/pi.hpp>
+#include <CL/sycl/device_selector.hpp>
+#include <CL/sycl/platform.hpp>
+#include <CL/sycl/queue.hpp>
 #include <detail/platform_impl.hpp>
 
 #include <functional>

sycl/unittests/helpers/TestKernel.hpp

@@ -0,0 +1,55 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "PiImage.hpp"
+
+class TestKernel;
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+template <> struct KernelInfo<TestKernel> {
+  static constexpr unsigned getNumParams() { return 0; }
+  static const kernel_param_desc_t &getParamDesc(int) {
+    static kernel_param_desc_t Dummy;
+    return Dummy;
+  }
+  static constexpr const char *getName() { return "TestKernel"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+};
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
+static sycl::unittest::PiImage generateDefaultImage() {
+  using namespace sycl::unittest;
+
+  PiPropertySet PropSet;
+
+  std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data
+
+  PiArray<PiOffloadEntry> Entries = makeEmptyKernels({"TestKernel"});
+
+  PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV,            // Format
+              __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec
+              "",                                     // Compile options
+              "",                                     // Link options
+              std::move(Bin),
+              std::move(Entries),
+              std::move(PropSet)};
+
+  return Img;
+}
+
+static sycl::unittest::PiImage Img = generateDefaultImage();
+static sycl::unittest::PiImageArray<1> ImgArray{&Img};

sycl/unittests/kernel-and-program/Cache.cpp

@@ -117,6 +117,7 @@ static pi_result redefinedKernelCreate(pi_program program,
                                        pi_kernel *ret_kernel) {
   return PI_SUCCESS;
 }
+
 static pi_result redefinedKernelRelease(pi_kernel kernel) { return PI_SUCCESS; }
 
 class KernelAndProgramCacheTest : public ::testing::Test {

sycl/unittests/pi/PiMock.cpp

@@ -6,10 +6,12 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <gtest/gtest.h>
 #include <helpers/PiMock.hpp>
+
 #include <detail/queue_impl.hpp>
 
+#include <gtest/gtest.h>
+
 using namespace cl::sycl;
 
 pi_result piProgramBuildRedefine(pi_program, pi_uint32, const pi_device *,

sycl/unittests/pi/TestGetPlatforms.hpp

@@ -4,7 +4,8 @@
 
 #pragma once
 
-#include <CL/sycl.hpp>
+#include <CL/sycl/platform.hpp>
+
 #include <algorithm>
 #include <functional>
 #include <vector>

sycl/unittests/program_manager/CMakeLists.txt

@@ -3,5 +3,6 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 add_sycl_unittest(ProgramManagerTests OBJECT
   EliminatedArgMask.cpp
   itt_annotations.cpp
+  SubDevices.cpp
 )