[SYCL] Add a unittest for is_compatible

sycl/source/detail/program_manager/program_manager.cpp

@@ -1605,6 +1605,17 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_Ptr2DeviceGlobal.insert({DeviceGlobalPtr, NewEntry.first->second.get()});
 }
 
+void ProgramManager::getRawDeviceImages(
+    const std::vector<kernel_id> &KernelIDs,
+    std::set<RTDeviceBinaryImage *> &BinImages) {
+  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+  for (const kernel_id &KID : KernelIDs) {
+    auto Range = m_KernelIDs2BinImage.equal_range(KID);
+    for (auto It = Range.first, End = Range.second; It != End; ++It)
+      BinImages.insert(It->second);
+  }
+}
+
 std::vector<device_image_plain>
 ProgramManager::getSYCLDeviceImagesWithCompatibleState(
     const context &Ctx, const std::vector<device> &Devs,
@@ -1614,12 +1625,7 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
   // TODO: Can we avoid repacking?
   std::set<RTDeviceBinaryImage *> BinImages;
   if (!KernelIDs.empty()) {
-    std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
-    for (const kernel_id &KID : KernelIDs) {
-      auto Range = m_KernelIDs2BinImage.equal_range(KID);
-      for (auto It = Range.first, End = Range.second; It != End; ++It)
-        BinImages.insert(It->second);
-    }
+    getRawDeviceImages(KernelIDs, BinImages);
   } else {
     std::lock_guard<std::mutex> Guard(Sync::getGlobalLock());
     for (auto &ImagesSets : m_DeviceImages) {
@@ -1628,7 +1634,7 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
         BinImages.insert(ImageUPtr.get());
     }
   }
-  assert(BinImages.size() > 0 && "Expected to find at least on device image");
+  assert(BinImages.size() > 0 && "Expected to find at least one device image");
 
   std::vector<device_image_plain> SYCLDeviceImages;
   for (RTDeviceBinaryImage *BinImage : BinImages) {

sycl/source/detail/program_manager/program_manager.hpp

@@ -255,6 +255,9 @@ class ProgramManager {
 
   bool kernelUsesAssert(OSModuleHandle M, const std::string &KernelName) const;
 
+  void getRawDeviceImages(const std::vector<kernel_id> &KernelIDs,
+                          std::set<RTDeviceBinaryImage *> &BinImages);
+
 private:
   ProgramManager(ProgramManager const &) = delete;
   ProgramManager &operator=(ProgramManager const &) = delete;

sycl/source/kernel_bundle.cpp

@@ -296,20 +296,17 @@ std::vector<kernel_id> get_kernel_ids() {
 }
 
 bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
-  for (const auto &KernelId : KernelIDs) {
-    const detail::RTDeviceBinaryImage &Img =
-        detail::ProgramManager::getInstance().getDeviceImage(
-            detail::OSUtil::ExeModuleHandle, KernelId.get_name(), context(Dev),
-            Dev);
-    const detail::RTDeviceBinaryImage::PropertyRange &ARange =
-        Img.getDeviceRequirements();
-    for (detail::RTDeviceBinaryImage::PropertyRange::ConstIterator It :
-         ARange) {
+  using namespace detail;
+  std::set<RTDeviceBinaryImage *> BinImages;
+  ProgramManager::getInstance().getRawDeviceImages(KernelIDs, BinImages);
+  for (RTDeviceBinaryImage *Img : BinImages) {
+    const RTDeviceBinaryImage::PropertyRange &PropRange =
+        Img->getDeviceRequirements();
+    for (RTDeviceBinaryImage::PropertyRange::ConstIterator It : PropRange) {
       using namespace std::literals;
       if ((*It)->Name != "aspects"sv)
         continue;
-      detail::ByteArray Aspects =
-          detail::DeviceBinaryProperty(*It).asByteArray();
+      ByteArray Aspects = DeviceBinaryProperty(*It).asByteArray();
       // Drop 8 bytes describing the size of the byte array
       Aspects.dropBytes(8);
       while (!Aspects.empty()) {

sycl/unittests/SYCL2020/CMakeLists.txt

@@ -6,5 +6,6 @@ add_sycl_unittest(SYCL2020Tests OBJECT
   KernelBundle.cpp
   KernelID.cpp
   HasExtension.cpp
+  IsCompatible.cpp
 )
 

sycl/unittests/SYCL2020/IsCompatible.cpp

@@ -0,0 +1,177 @@
+#include <sycl/sycl.hpp>
+
+#include <helpers/PiImage.hpp>
+#include <helpers/PiMock.hpp>
+
+#include <gtest/gtest.h>
+
+class TestKernelCPU;
+class TestKernelGPU;
+class TestKernelACC;
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+template <> struct KernelInfo<TestKernelCPU> {
+  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 "TestKernelCPU"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
+};
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+template <> struct KernelInfo<TestKernelGPU> {
+  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 "TestKernelGPU"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
+};
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace detail {
+template <> struct KernelInfo<TestKernelACC> {
+  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 "TestKernelACC"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
+};
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
+
+static sycl::unittest::PiImage
+generateDefaultImage(std::initializer_list<std::string> KernelNames,
+                     const std::vector<sycl::aspect> &Aspects) {
+  using namespace sycl::unittest;
+
+  PiPropertySet PropSet;
+  addAspects(PropSet, Aspects);
+
+  std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data
+
+  PiArray<PiOffloadEntry> Entries = makeEmptyKernels(KernelNames);
+
+  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 Imgs[3] = {
+    generateDefaultImage({"TestKernelCPU"}, {sycl::aspect::cpu}),
+    generateDefaultImage({"TestKernelGPU"}, {sycl::aspect::gpu}),
+    generateDefaultImage({"TestKernelACC"}, {sycl::aspect::accelerator})};
+
+static sycl::unittest::PiImageArray<3> ImgArray{Imgs};
+
+static pi_result redefinedDeviceGetInfoCPU(pi_device device,
+                                           pi_device_info param_name,
+                                           size_t param_value_size,
+                                           void *param_value,
+                                           size_t *param_value_size_ret) {
+  if (param_name == PI_DEVICE_INFO_TYPE) {
+    auto *Result = reinterpret_cast<_pi_device_type *>(param_value);
+    *Result = PI_DEVICE_TYPE_CPU;
+  }
+  return PI_SUCCESS;
+}
+
+// Mock device is "GPU" by default, but we need to redefine it just in case
+// if there are some changes in the future
+static pi_result redefinedDeviceGetInfoGPU(pi_device device,
+                                           pi_device_info param_name,
+                                           size_t param_value_size,
+                                           void *param_value,
+                                           size_t *param_value_size_ret) {
+  if (param_name == PI_DEVICE_INFO_TYPE) {
+    auto *Result = reinterpret_cast<_pi_device_type *>(param_value);
+    *Result = PI_DEVICE_TYPE_GPU;
+  }
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedDeviceGetInfoACC(pi_device device,
+                                           pi_device_info param_name,
+                                           size_t param_value_size,
+                                           void *param_value,
+                                           size_t *param_value_size_ret) {
+  if (param_name == PI_DEVICE_INFO_TYPE) {
+    auto *Result = reinterpret_cast<_pi_device_type *>(param_value);
+    *Result = PI_DEVICE_TYPE_ACC;
+  }
+  return PI_SUCCESS;
+}
+
+TEST(IsCompatible, CPU) {
+  sycl::unittest::PiMock Mock;
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoCPU);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  EXPECT_TRUE(Dev.is_cpu());
+  EXPECT_TRUE(sycl::is_compatible<TestKernelCPU>(Dev));
+  EXPECT_FALSE(sycl::is_compatible<TestKernelGPU>(Dev));
+  EXPECT_FALSE(sycl::is_compatible<TestKernelACC>(Dev));
+}
+
+TEST(IsCompatible, GPU) {
+  sycl::unittest::PiMock Mock;
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoGPU);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  EXPECT_TRUE(Dev.is_gpu());
+  EXPECT_FALSE(sycl::is_compatible<TestKernelCPU>(Dev));
+  EXPECT_TRUE(sycl::is_compatible<TestKernelGPU>(Dev));
+  EXPECT_FALSE(sycl::is_compatible<TestKernelACC>(Dev));
+}
+
+TEST(IsCompatible, ACC) {
+  sycl::unittest::PiMock Mock;
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoACC);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  EXPECT_TRUE(Dev.is_accelerator());
+  EXPECT_FALSE(sycl::is_compatible<TestKernelCPU>(Dev));
+  EXPECT_FALSE(sycl::is_compatible<TestKernelGPU>(Dev));
+  EXPECT_TRUE(sycl::is_compatible<TestKernelACC>(Dev));
+}

sycl/unittests/helpers/PiImage.hpp

@@ -429,6 +429,24 @@ makeKernelParamOptInfo(const std::string &Name, const size_t NumArgs,
   return Prop;
 }
 
+/// Utility function to add aspects to property set.
+inline void addAspects(PiPropertySet &Props,
+                       const std::vector<sycl::aspect> &Aspects) {
+  const size_t BYTES_FOR_SIZE = 8;
+  std::vector<char> ValData(BYTES_FOR_SIZE +
+                            Aspects.size() * sizeof(sycl::aspect));
+  uint64_t ValDataSize = ValData.size();
+  std::uninitialized_copy(&ValDataSize, &ValDataSize + sizeof(uint64_t),
+                          ValData.data());
+  auto *AspectsPtr = reinterpret_cast<const unsigned char *>(&Aspects[0]);
+  std::uninitialized_copy(AspectsPtr, AspectsPtr + Aspects.size(),
+                          ValData.data() + BYTES_FOR_SIZE);
+  PiProperty Prop{"aspects", ValData, PI_PROPERTY_TYPE_BYTE_ARRAY};
+  PiArray<PiProperty> Value{std::move(Prop)};
+  Props.insert(__SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS,
+               std::move(Value));
+}
+
 } // namespace unittest
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl