From 9d30d4922cc6a734dc6930b28ee3cdfa8a9ecf3f Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitry.vodopyanov@intel.com>
Date: Fri, 30 Apr 2021 20:08:57 +0300
Subject: [PATCH 1/9] [SYCL] Add DPC++ RT support for SYCL2020 spec constants'
 default values

This patch introduces the support of default values for SYCL2020
specialization constants in DPC++ runtime.
---
 sycl/include/CL/sycl/detail/pi.h         |  4 ++++
 sycl/include/CL/sycl/detail/pi.hpp       |  4 ++++
 sycl/source/detail/device_image_impl.hpp | 14 ++++++++++++++
 sycl/source/detail/pi.cpp                |  2 ++
 4 files changed, 24 insertions(+)

diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h
index a6cdadf664310..e308f5e8f63e2 100644
--- a/sycl/include/CL/sycl/detail/pi.h
+++ b/sycl/include/CL/sycl/detail/pi.h
@@ -673,6 +673,10 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 /// PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS defined in
 /// PropertySetIO.h
 #define __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP "SYCL/specialization constants"
+/// PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES defined in
+/// PropertySetIO.h
+#define __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP                   \
+  "SYCL/specialization constants default values"
 /// PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK defined in PropertySetIO.h
 #define __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK "SYCL/devicelib req mask"
 /// PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO defined in PropertySetIO.h
diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp
index 03f3915b1dcb6..5f1cf1600d6c8 100644
--- a/sycl/include/CL/sycl/detail/pi.hpp
+++ b/sycl/include/CL/sycl/detail/pi.hpp
@@ -318,6 +318,9 @@ class DeviceBinaryImage {
   /// like:
   /// { ID5, 0, 4 }
   const PropertyRange &getSpecConstants() const { return SpecConstIDMap; }
+  const PropertyRange &getSpecConstantsDefaultValues() const {
+    return SpecConstDefaultValuesMap;
+  }
   const PropertyRange &getDeviceLibReqMask() const { return DeviceLibReqMask; }
   const PropertyRange &getKernelParamOptInfo() const {
     return KernelParamOptInfo;
@@ -331,6 +334,7 @@ class DeviceBinaryImage {
   pi_device_binary Bin;
   pi::PiDeviceBinaryType Format = PI_DEVICE_BINARY_TYPE_NONE;
   DeviceBinaryImage::PropertyRange SpecConstIDMap;
+  DeviceBinaryImage::PropertyRange SpecConstDefaultValuesMap;
   DeviceBinaryImage::PropertyRange DeviceLibReqMask;
   DeviceBinaryImage::PropertyRange KernelParamOptInfo;
 };
diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index 697404d27a5db..b60497520dd5e 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -243,6 +243,20 @@ class device_image_impl {
         }
       }
       MSpecConstsBlob.resize(BlobOffset);
+
+      // set default values for specialization constants
+      const pi::DeviceBinaryImage::PropertyRange &SCDefValRange =
+          MBinImage->getSpecConstantsDefaultValues();
+      for (SCItTy SCIt : SCDefValRange) {
+        const char *SCName = (*SCIt)->Name;
+        pi::ByteArray Descriptors =
+            pi::DeviceBinaryProperty(*SCIt).asByteArray();
+        // TODO: same 8 bytes are the size of this new property?
+        assert(Descriptors.size() > 8 && "Unexpected property size");
+        // TODO: need to simplify it
+        auto Value = reinterpret_cast<const char *>(&Descriptors[8])[0];
+        set_specialization_constant_raw_value(SCName, Value);
+      }
     }
   }
 
diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp
index f5a393d2c0d82..f7e78145f4b36 100644
--- a/sycl/source/detail/pi.cpp
+++ b/sycl/source/detail/pi.cpp
@@ -615,6 +615,8 @@ void DeviceBinaryImage::init(pi_device_binary Bin) {
     Format = getBinaryImageFormat(Bin->BinaryStart, getSize());
 
   SpecConstIDMap.init(Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP);
+  SpecConstDefaultValuesMap.init(
+      Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP);
   DeviceLibReqMask.init(Bin, __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK);
   KernelParamOptInfo.init(Bin, __SYCL_PI_PROPERTY_SET_KERNEL_PARAM_OPT_INFO);
 }

From 41893a978566ce62d9788e660d6a3771751d9e28 Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitry.vodopyanov@intel.com>
Date: Fri, 30 Apr 2021 21:25:34 +0300
Subject: [PATCH 2/9] Fix comp error

---
 sycl/source/detail/device_image_impl.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index b60497520dd5e..f91d57c119820 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -254,7 +254,7 @@ class device_image_impl {
         // TODO: same 8 bytes are the size of this new property?
         assert(Descriptors.size() > 8 && "Unexpected property size");
         // TODO: need to simplify it
-        auto Value = reinterpret_cast<const char *>(&Descriptors[8])[0];
+        const auto Value = reinterpret_cast<const void *>(&Descriptors[8]);
         set_specialization_constant_raw_value(SCName, Value);
       }
     }

From 61b139ac2885fb3f794b7eb35cd222524effd864 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 7 May 2021 12:44:21 +0300
Subject: [PATCH 3/9] Add unit tests

---
 sycl/include/CL/sycl/detail/pi.hpp            |   8 +-
 sycl/source/CMakeLists.txt                    |   3 +
 sycl/source/detail/device_image_impl.hpp      |  26 +-
 sycl/source/detail/pi.cpp                     |   2 -
 sycl/unittests/CMakeLists.txt                 |   1 +
 sycl/unittests/helpers/PiImage.hpp            | 357 ++++++++++++++++++
 sycl/unittests/spec_constants/CMakeLists.txt  |   8 +
 .../spec_constants/DefaultValues.cpp          | 263 +++++++++++++
 8 files changed, 651 insertions(+), 17 deletions(-)
 create mode 100644 sycl/unittests/helpers/PiImage.hpp
 create mode 100644 sycl/unittests/spec_constants/CMakeLists.txt
 create mode 100644 sycl/unittests/spec_constants/DefaultValues.cpp

diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp
index 5f1cf1600d6c8..e06ae106e65e7 100644
--- a/sycl/include/CL/sycl/detail/pi.hpp
+++ b/sycl/include/CL/sycl/detail/pi.hpp
@@ -318,7 +318,12 @@ class DeviceBinaryImage {
   /// like:
   /// { ID5, 0, 4 }
   const PropertyRange &getSpecConstants() const { return SpecConstIDMap; }
-  const PropertyRange &getSpecConstantsDefaultValues() const {
+  const PropertyRange getSpecConstantsDefaultValues() const {
+    // We can't have this variable as a class member, since it would break
+    // the ABI backwards compatibility.
+    DeviceBinaryImage::PropertyRange SpecConstDefaultValuesMap;
+    SpecConstDefaultValuesMap.init(
+        Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP);
     return SpecConstDefaultValuesMap;
   }
   const PropertyRange &getDeviceLibReqMask() const { return DeviceLibReqMask; }
@@ -334,7 +339,6 @@ class DeviceBinaryImage {
   pi_device_binary Bin;
   pi::PiDeviceBinaryType Format = PI_DEVICE_BINARY_TYPE_NONE;
   DeviceBinaryImage::PropertyRange SpecConstIDMap;
-  DeviceBinaryImage::PropertyRange SpecConstDefaultValuesMap;
   DeviceBinaryImage::PropertyRange DeviceLibReqMask;
   DeviceBinaryImage::PropertyRange KernelParamOptInfo;
 };
diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
index 0c6535f21c6de..d392fe5246e38 100644
--- a/sycl/source/CMakeLists.txt
+++ b/sycl/source/CMakeLists.txt
@@ -52,6 +52,9 @@ function(add_sycl_rt_library LIB_NAME)
   else()
     target_compile_options(${LIB_OBJ_NAME} PUBLIC
                            -fvisibility=hidden -fvisibility-inlines-hidden)
+
+    target_compile_options(${LIB_OBJ_NAME} PUBLIC
+                           -ggdb -O0)
     set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt")
     target_link_libraries(
       ${LIB_NAME} PRIVATE "-Wl,--version-script=${linker_script}")
diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index f91d57c119820..8aa7910282c0f 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -205,6 +205,12 @@ class device_image_impl {
           MBinImage->getSpecConstants();
       using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator;
 
+      // get default values for specialization constants
+      const pi::DeviceBinaryImage::PropertyRange &SCDefValRange =
+          MBinImage->getSpecConstantsDefaultValues();
+
+      bool HasDefaultValues = SCDefValRange.begin() != SCDefValRange.end();
+
       // This variable is used to calculate spec constant value offset in a
       // flat byte array.
       unsigned BlobOffset = 0;
@@ -237,25 +243,19 @@ class device_image_impl {
           // supposed to be called from c'tor.
           MSpecConstSymMap[std::string{SCName}].push_back(
               SpecConstDescT{/*ID*/ It[0], /*CompositeOffset*/ It[1],
-                             /*Size*/ It[2], BlobOffset});
+                             /*Size*/ It[2], BlobOffset, HasDefaultValues});
           BlobOffset += /*Size*/ It[2];
           It += NumElements;
         }
       }
       MSpecConstsBlob.resize(BlobOffset);
 
-      // set default values for specialization constants
-      const pi::DeviceBinaryImage::PropertyRange &SCDefValRange =
-          MBinImage->getSpecConstantsDefaultValues();
-      for (SCItTy SCIt : SCDefValRange) {
-        const char *SCName = (*SCIt)->Name;
-        pi::ByteArray Descriptors =
-            pi::DeviceBinaryProperty(*SCIt).asByteArray();
-        // TODO: same 8 bytes are the size of this new property?
-        assert(Descriptors.size() > 8 && "Unexpected property size");
-        // TODO: need to simplify it
-        const auto Value = reinterpret_cast<const void *>(&Descriptors[8]);
-        set_specialization_constant_raw_value(SCName, Value);
+      if (HasDefaultValues) {
+        pi::ByteArray DefValDescriptors =
+            pi::DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray();
+        std::uninitialized_copy(&DefValDescriptors[0],
+                                &DefValDescriptors[0] + MSpecConstsBlob.size(),
+                                MSpecConstsBlob.data());
       }
     }
   }
diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp
index f7e78145f4b36..f5a393d2c0d82 100644
--- a/sycl/source/detail/pi.cpp
+++ b/sycl/source/detail/pi.cpp
@@ -615,8 +615,6 @@ void DeviceBinaryImage::init(pi_device_binary Bin) {
     Format = getBinaryImageFormat(Bin->BinaryStart, getSize());
 
   SpecConstIDMap.init(Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP);
-  SpecConstDefaultValuesMap.init(
-      Bin, __SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP);
   DeviceLibReqMask.init(Bin, __SYCL_PI_PROPERTY_SET_DEVICELIB_REQ_MASK);
   KernelParamOptInfo.init(Bin, __SYCL_PI_PROPERTY_SET_KERNEL_PARAM_OPT_INFO);
 }
diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt
index f2789b4ffb6a1..5da1955d9e5ff 100644
--- a/sycl/unittests/CMakeLists.txt
+++ b/sycl/unittests/CMakeLists.txt
@@ -15,4 +15,5 @@ add_subdirectory(pi)
 add_subdirectory(kernel-and-program)
 add_subdirectory(queue)
 add_subdirectory(scheduler)
+add_subdirectory(spec_constants)
 add_subdirectory(thread_safety)
diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp
new file mode 100644
index 0000000000000..6495a96ac8eb1
--- /dev/null
+++ b/sycl/unittests/helpers/PiImage.hpp
@@ -0,0 +1,357 @@
+//==------------- PiImage.hpp --- PI mock image unit testing library -------==//
+//
+// 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 <CL/sycl.hpp>
+#include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/pi.hpp>
+#include <detail/platform_impl.hpp>
+#include <detail/program_manager/program_manager.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace unittest {
+/// Convinience wrapper around _pi_device_binary_property_struct.
+class PiProperty {
+public:
+  using NativeType = _pi_device_binary_property_struct;
+  PiProperty(const std::string &Name, std::vector<char> Data, uint32_t Type)
+      : MName(Name), MData(std::move(Data)), MType(Type) {
+    updateNativeType();
+  }
+  NativeType convertToNativeType() const { return MNative; }
+
+  PiProperty(const PiProperty &Src) {
+    MName = Src.MName;
+    MData = Src.MData;
+    MType = Src.MType;
+    updateNativeType();
+  }
+  PiProperty &operator=(const PiProperty &Src) {
+    MName = Src.MName;
+    MData = Src.MData;
+    MType = Src.MType;
+    updateNativeType();
+    return *this;
+  }
+
+private:
+  void updateNativeType() {
+    MNative = NativeType{const_cast<char *>(MName.c_str()),
+                         const_cast<char *>(MData.data()), MType, MData.size()};
+  }
+  std::string MName;
+  std::vector<char> MData;
+  uint32_t MType;
+  NativeType MNative;
+};
+
+/// Convinience wrapper for _pi_offload_entry_struct.
+class PiOffloadEntry {
+public:
+  using NativeType = _pi_offload_entry_struct;
+  PiOffloadEntry(const std::string &Name, std::vector<char> Data, int32_t Flags)
+      : MName(Name), MData(std::move(Data)), MFlags(Flags) {
+    updateNativeType();
+  }
+
+  PiOffloadEntry(const PiOffloadEntry &Src) {
+    MName = Src.MName;
+    MData = Src.MData;
+    MFlags = Src.MFlags;
+    updateNativeType();
+  }
+  PiOffloadEntry &operator=(const PiOffloadEntry &Src) {
+    MName = Src.MName;
+    MData = Src.MData;
+    MFlags = Src.MFlags;
+    updateNativeType();
+    return *this;
+  }
+
+  NativeType convertToNativeType() const { return MNative; }
+
+private:
+  void updateNativeType() {
+    MNative = NativeType{
+        const_cast<char *>(MData.data()), MName.data(), MData.size(), MFlags,
+        0 // Reserved
+    };
+  }
+  std::string MName;
+  std::vector<char> MData;
+  int32_t MFlags;
+  NativeType MNative;
+};
+
+/// Generic array of PI entries.
+template <typename T> class PiArray {
+public:
+  explicit PiArray(std::vector<T> Entries) : MMockEntries(std::move(Entries)) {
+    std::transform(MMockEntries.begins(), MMockEntries.end(),
+                   std::back_inserter(MEntries),
+                   [](const T &Entry) { return Entry.convertToNativeType(); });
+  }
+
+  PiArray(std::initializer_list<T> Entries) : MMockEntries(std::move(Entries)) {
+    std::transform(MMockEntries.begin(), MMockEntries.end(),
+                   std::back_inserter(MEntries),
+                   [](const T &Entry) { return Entry.convertToNativeType(); });
+  }
+
+  PiArray() = default;
+
+  void push_back(const T &Entry) {
+    MMockEntries.push_back(Entry);
+    MEntries.push_back(MMockEntries.back().convertToNativeType());
+  }
+
+  typename T::NativeType *begin() { return &*MEntries.begin(); }
+  typename T::NativeType *end() { return &*MEntries.end(); }
+
+private:
+  std::vector<T> MMockEntries;
+  std::vector<typename T::NativeType> MEntries;
+};
+
+/// Convenience wrapper for pi_device_binary_property_set.
+class PiPropertySet {
+public:
+  PiPropertySet() = default;
+
+  void insert(const std::string &Name, PiArray<PiProperty> Props) {
+    MNames.push_back(Name);
+    MMockProperties.push_back(std::move(Props));
+    MProperties.push_back(_pi_device_binary_property_set_struct{
+        MNames.back().data(), MMockProperties.back().begin(),
+        MMockProperties.back().end()});
+  }
+
+  _pi_device_binary_property_set_struct *begin() {
+    if (MProperties.empty())
+      return nullptr;
+    return &*MProperties.begin();
+  }
+
+  _pi_device_binary_property_set_struct *end() {
+    if (MProperties.empty())
+      return nullptr;
+    return &*MProperties.end();
+  }
+
+private:
+  std::vector<std::string> MNames;
+  std::vector<PiArray<PiProperty>> MMockProperties;
+  std::vector<_pi_device_binary_property_set_struct> MProperties;
+};
+
+/// Convenience wrapper around PI internal structures, that manages PI binary
+/// image data lifecycle.
+class PiImage {
+public:
+  /// Constructs an arbitrary device image.
+  PiImage(uint16_t Version, uint8_t Kind, uint8_t Format,
+          const std::string &DeviceTargetSpec,
+          const std::string &CompileOptions, const std::string &LinkOptions,
+          std::vector<char> Manifest, std::vector<unsigned char> Binary,
+          PiArray<PiOffloadEntry> OffloadEntries, PiPropertySet PropertySet)
+      : MDeviceTargetSpec(DeviceTargetSpec), MCompileOptions(CompileOptions),
+        MLinkOptions(LinkOptions), MManifest(std::move(Manifest)),
+        MBinary(std::move(Binary)), MOffloadEntries(std::move(OffloadEntries)),
+        MPropertySet(std::move(PropertySet)) {
+    auto [ManifestStart,
+          ManifestEnd] = [this]() -> std::pair<const char *, const char *> {
+      if (!MManifest.empty())
+        return {&*MManifest.cbegin(), &*MManifest.cend()};
+      return {nullptr, nullptr};
+    }();
+    MBinaryDesc = pi_device_binary_struct{
+        Version,
+        Kind,
+        Format,
+        MDeviceTargetSpec.c_str(),
+        MCompileOptions.c_str(),
+        MLinkOptions.c_str(),
+        ManifestStart,
+        ManifestEnd,
+        &*MBinary.begin(),
+        &*MBinary.end(),
+        MOffloadEntries.begin(),
+        MOffloadEntries.end(),
+        MPropertySet.begin(),
+        MPropertySet.end(),
+    };
+  }
+
+  /// Constructs a SYCL device image of the latest version.
+  PiImage(uint8_t Format, const std::string &DeviceTargetSpec,
+          const std::string &CompileOptions, const std::string &LinkOptions,
+          std::vector<unsigned char> Binary,
+          PiArray<PiOffloadEntry> OffloadEntries, PiPropertySet PropertySet)
+      : PiImage(PI_DEVICE_BINARY_VERSION, PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL,
+                Format, DeviceTargetSpec, CompileOptions, LinkOptions, {},
+                std::move(Binary), std::move(OffloadEntries),
+                std::move(PropertySet)) {}
+
+  pi_device_binary_struct convertToNativeType() const { return MBinaryDesc; }
+
+private:
+  std::string MDeviceTargetSpec;
+  std::string MCompileOptions;
+  std::string MLinkOptions;
+  std::vector<char> MManifest;
+  std::vector<unsigned char> MBinary;
+  PiArray<PiOffloadEntry> MOffloadEntries;
+  PiPropertySet MPropertySet;
+  pi_device_binary_struct MBinaryDesc;
+};
+
+/// Convenience wrapper around pi_device_binaries_struct, that manages mock
+/// device images' lifecycle.
+class PiImageArray {
+public:
+  /// Constructs an array of device images from a single image and registers
+  /// it with SYCL runtime.
+  PiImageArray(PiImage Image) {
+    MImages.push_back(std::move(Image));
+    convertImages();
+    MAllBinaries = pi_device_binaries_struct{
+        PI_DEVICE_BINARIES_VERSION,
+        1, // num binaries
+        MNativeImages.data(),
+        nullptr, // not used, for compatibility with OpenMP
+        nullptr  // not used, for compatibility with OpenMP
+    };
+    __sycl_register_lib(&MAllBinaries);
+  }
+
+  /// Constructs an array of device images and registers it with SYCL runtime.
+  PiImageArray(std::vector<PiImage> Images) : MImages(std::move(Images)) {
+    convertImages();
+    MAllBinaries = pi_device_binaries_struct{
+        PI_DEVICE_BINARIES_VERSION,
+        static_cast<uint16_t>(MNativeImages.size()), // num binaries
+        MNativeImages.data(),
+        nullptr, // not used, for compatibility with OpenMP
+        nullptr  // not used, for compatibility with OpenMP
+    };
+    __sycl_register_lib(&MAllBinaries);
+  }
+
+  ~PiImageArray() { __sycl_unregister_lib(&MAllBinaries); }
+
+private:
+  void convertImages() {
+    std::transform(
+        MImages.begin(), MImages.end(), std::back_inserter(MNativeImages),
+        [](const PiImage &Img) { return Img.convertToNativeType(); });
+  }
+  std::vector<PiImage> MImages;
+  std::vector<pi_device_binary_struct> MNativeImages;
+  pi_device_binaries_struct MAllBinaries;
+};
+
+template <typename Func, uint32_t Idx = 0, typename... Ts>
+std::enable_if_t<Idx == sizeof...(Ts)> iterate_tuple(Func &F,
+                                                     std::tuple<Ts...> &Tuple) {
+  return;
+}
+template <typename Func, uint32_t Idx = 0, typename... Ts>
+    std::enable_if_t <
+    Idx<sizeof...(Ts)> iterate_tuple(Func &F, std::tuple<Ts...> &Tuple) {
+  const auto &Value = std::get<Idx>(Tuple);
+  const char *Begin = reinterpret_cast<const char *>(&Value);
+  const char *End = Begin + sizeof(Value);
+  F(Idx, Begin, End);
+
+  iterate_tuple<Func, Idx + 1, Ts...>(F, Tuple);
+  return;
+}
+
+/// Utility function to create a single spec constant property.
+///
+/// \param ValData is a reference to blob array, that stores default values.
+/// \param Name is a spec constant name.
+/// \param IDs is a list of spec IDs.
+/// \param Offsets is a list of offsets inside composite spec constant.
+/// \param DefaultValues is a tuple of default values for composite spec const.
+template <typename... T>
+PiProperty makeSpecConstant(std::vector<char> &ValData, const std::string &Name,
+                            std::initializer_list<uint32_t> IDs,
+                            std::initializer_list<uint32_t> Offsets,
+                            std::tuple<T...> DefaultValues) {
+  const size_t PropByteArraySize = sizeof...(T) * sizeof(uint32_t) * 3;
+  std::vector<char> DescData;
+  DescData.resize(8 + PropByteArraySize);
+  std::uninitialized_copy(&PropByteArraySize, &PropByteArraySize + 8,
+                          DescData.data());
+
+  size_t PrevSize = ValData.size();
+  // Resize raw data blob to current size + offset of the last element + size of
+  // the last element.
+  ValData.resize(
+      PrevSize + *std::prev(Offsets.end()) +
+      sizeof(typename std::tuple_element<sizeof...(T) - 1,
+                                         decltype(DefaultValues)>::type));
+
+  auto FillData = [PrevOffset = 0, PrevSize, &ValData, &IDs, &Offsets,
+                   &DescData](uint32_t Idx, const char *Begin,
+                              const char *End) mutable {
+    const size_t Offset = 8 + Idx * sizeof(uint32_t) * 3;
+
+    uint32_t ValSize = std::distance(Begin, End);
+    const char *IDsBegin =
+        reinterpret_cast<const char *>(&*std::next(IDs.begin(), Idx));
+    const char *OffsetBegin =
+        reinterpret_cast<const char *>(&*std::next(Offsets.begin(), Idx));
+    const char *ValSizeBegin = reinterpret_cast<const char *>(&ValSize);
+
+    std::uninitialized_copy(IDsBegin, IDsBegin + sizeof(uint32_t),
+                            DescData.data() + Offset);
+    std::uninitialized_copy(OffsetBegin, OffsetBegin + sizeof(uint32_t),
+                            DescData.data() + Offset + sizeof(uint32_t));
+    std::uninitialized_copy(ValSizeBegin, ValSizeBegin + sizeof(uint32_t),
+                            DescData.data() + Offset + 2 * sizeof(uint32_t));
+    std::uninitialized_copy(Begin, End, ValData.data() + PrevSize + PrevOffset);
+    PrevOffset += *std::next(Offsets.begin(), Idx);
+  };
+
+  iterate_tuple(FillData, DefaultValues);
+
+  PiProperty Prop{Name, DescData, PI_PROPERTY_TYPE_BYTE_ARRAY};
+
+  return Prop;
+}
+
+void addSpecConstants(PiArray<PiProperty> SpecConstants,
+                      std::vector<char> ValData, PiPropertySet &Props) {
+  Props.insert(__SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP, std::move(SpecConstants));
+
+  PiProperty Prop{"all", std::move(ValData), PI_PROPERTY_TYPE_BYTE_ARRAY};
+
+  PiArray<PiProperty> DefaultValues{std::move(Prop)};
+
+  Props.insert(__SYCL_PI_PROPERTY_SET_SPEC_CONST_DEFAULT_VALUES_MAP,
+               std::move(DefaultValues));
+}
+
+PiArray<PiOffloadEntry>
+makeEmptyKernels(std::initializer_list<std::string> KernelNames) {
+  PiArray<PiOffloadEntry> Entries;
+
+  for (const auto &Name : KernelNames) {
+    PiOffloadEntry E{Name, {}, 0};
+    Entries.push_back(std::move(E));
+  }
+  return Entries;
+}
+
+} // namespace unittest
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/unittests/spec_constants/CMakeLists.txt b/sycl/unittests/spec_constants/CMakeLists.txt
new file mode 100644
index 0000000000000..cff537cd2963a
--- /dev/null
+++ b/sycl/unittests/spec_constants/CMakeLists.txt
@@ -0,0 +1,8 @@
+set(CMAKE_CXX_EXTENSIONS OFF)
+
+# Enable exception handling for these unit tests
+set(LLVM_REQUIRES_EH 1)
+add_sycl_unittest(SpecConstantsTests OBJECT
+  DefaultValues.cpp
+)
+
diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp
new file mode 100644
index 0000000000000..5ce5c279bf69d
--- /dev/null
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -0,0 +1,263 @@
+//==---- DefaultValues.cpp --- Spec constants default values unit test -----==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#define SYCL2020_DISABLE_DEPRECATION_WARNINGS
+
+#include <CL/sycl.hpp>
+#include <helpers/PiImage.hpp>
+#include <helpers/PiMock.hpp>
+
+#include <gtest/gtest.h>
+
+class TestKernel;
+class TestKernel2;
+const static sycl::specialization_id<int> SpecConst1{42};
+
+__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; }
+};
+
+template <> const char *get_spec_constant_symbolic_ID<SpecConst1>() {
+  return "SC1";
+}
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
+static pi_result redefinedProgramCreate(pi_context, const void *, size_t,
+                                        pi_program *) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramBuild(
+    pi_program prog, pi_uint32, const pi_device *, const char *,
+    void (*pfn_notify)(pi_program program, void *user_data), void *user_data) {
+  if (pfn_notify) {
+    pfn_notify(prog, user_data);
+  }
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramCompile(pi_program, pi_uint32,
+                                         const pi_device *, const char *,
+                                         pi_uint32, const pi_program *,
+                                         const char **,
+                                         void (*)(pi_program, void *), void *) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramLink(pi_context, pi_uint32, const pi_device *,
+                                      const char *, pi_uint32,
+                                      const pi_program *,
+                                      void (*)(pi_program, void *), void *,
+                                      pi_program *) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramGetInfo(pi_program program,
+                                         pi_program_info param_name,
+                                         size_t param_value_size,
+                                         void *param_value,
+                                         size_t *param_value_size_ret) {
+  if (param_name == PI_PROGRAM_INFO_NUM_DEVICES) {
+    auto value = reinterpret_cast<unsigned int *>(param_value);
+    *value = 1;
+  }
+
+  if (param_name == PI_PROGRAM_INFO_BINARY_SIZES) {
+    auto value = reinterpret_cast<size_t *>(param_value);
+    value[0] = 1;
+  }
+
+  if (param_name == PI_PROGRAM_INFO_BINARIES) {
+    auto value = reinterpret_cast<unsigned char *>(param_value);
+    value[0] = 1;
+  }
+
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramRetain(pi_program program) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedProgramRelease(pi_program program) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedKernelCreate(pi_program program,
+                                       const char *kernel_name,
+                                       pi_kernel *ret_kernel) {
+  *ret_kernel = reinterpret_cast<pi_kernel>(new int[1]);
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedKernelRetain(pi_kernel kernel) { return PI_SUCCESS; }
+
+static pi_result redefinedKernelRelease(pi_kernel kernel) {
+  delete[] reinterpret_cast<int *>(kernel);
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedKernelGetInfo(pi_kernel kernel,
+                                        pi_kernel_info param_name,
+                                        size_t param_value_size,
+                                        void *param_value,
+                                        size_t *param_value_size_ret) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedKernelSetExecInfo(pi_kernel kernel,
+                                            pi_kernel_exec_info value_name,
+                                            size_t param_value_size,
+                                            const void *param_value) {
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedEventsWait(pi_uint32 num_events,
+                                     const pi_event *event_list) {
+  return PI_SUCCESS;
+}
+
+int SpecConstVal0 = 0;
+int SpecConstVal1 = 0;
+
+static pi_result
+redefinedProgramSetSpecializationConstant(pi_program prog, pi_uint32 spec_id,
+                                          size_t spec_size,
+                                          const void *spec_value) {
+  if (spec_id == 0)
+    SpecConstVal0 = *static_cast<const int *>(spec_value);
+  if (spec_id == 1)
+    SpecConstVal1 = *static_cast<const int *>(spec_value);
+
+  return PI_SUCCESS;
+}
+
+static pi_result redefinedEnqueueKernelLaunch(pi_queue, pi_kernel, pi_uint32,
+                                              const size_t *, const size_t *,
+                                              const size_t *, pi_uint32,
+                                              const pi_event *, pi_event *) {
+  return PI_SUCCESS;
+}
+
+static void setupDefaultMockAPIs(sycl::unittest::PiMock &Mock) {
+  using namespace sycl::detail;
+  Mock.redefine<PiApiKind::piProgramCreate>(redefinedProgramCreate);
+  Mock.redefine<PiApiKind::piProgramCompile>(redefinedProgramCompile);
+  Mock.redefine<PiApiKind::piProgramLink>(redefinedProgramLink);
+  Mock.redefine<PiApiKind::piProgramBuild>(redefinedProgramBuild);
+  Mock.redefine<PiApiKind::piProgramGetInfo>(redefinedProgramGetInfo);
+  Mock.redefine<PiApiKind::piProgramRetain>(redefinedProgramRetain);
+  Mock.redefine<PiApiKind::piProgramRelease>(redefinedProgramRelease);
+  Mock.redefine<PiApiKind::piKernelCreate>(redefinedKernelCreate);
+  Mock.redefine<PiApiKind::piKernelRetain>(redefinedKernelRetain);
+  Mock.redefine<PiApiKind::piKernelRelease>(redefinedKernelRelease);
+  Mock.redefine<PiApiKind::piKernelGetInfo>(redefinedKernelGetInfo);
+  Mock.redefine<PiApiKind::piKernelSetExecInfo>(redefinedKernelSetExecInfo);
+  Mock.redefine<PiApiKind::piextProgramSetSpecializationConstant>(
+      redefinedProgramSetSpecializationConstant);
+  Mock.redefine<PiApiKind::piEventsWait>(redefinedEventsWait);
+  Mock.redefine<PiApiKind::piEnqueueKernelLaunch>(redefinedEnqueueKernelLaunch);
+}
+
+static sycl::unittest::PiImage generateDefaultImage() {
+  using namespace sycl::unittest;
+
+  std::vector<char> SpecConstData;
+  PiProperty SC1 = makeSpecConstant<int>(SpecConstData, "SC1", {0}, {0}, {42});
+  PiProperty SC2 = makeSpecConstant<int>(SpecConstData, "SC2", {1}, {0}, {8});
+
+  PiPropertySet PropSet;
+  addSpecConstants({SC1, SC2}, std::move(SpecConstData), 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;
+}
+
+sycl::unittest::PiImage Img = generateDefaultImage();
+sycl::unittest::PiImageArray ImgArray{Img};
+
+TEST(DefaultValues, DefaultValuesAreSet) {
+  sycl::platform Plt{sycl::default_selector()};
+  if (Plt.is_host()) {
+    std::cerr << "Test is not supported on host, skipping\n";
+    return; // test is not supported on host.
+  }
+
+  sycl::unittest::PiMock Mock{Plt};
+  setupDefaultMockAPIs(Mock);
+
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  sycl::queue Queue{Dev};
+
+  const sycl::context Ctx = Queue.get_context();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
+  auto ExecBundle = sycl::build(KernelBundle);
+  Queue.submit([&](sycl::handler &CGH) {
+    CGH.use_kernel_bundle(ExecBundle);
+    CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+  });
+
+  EXPECT_EQ(SpecConstVal0, 42);
+  EXPECT_EQ(SpecConstVal1, 8);
+}
+
+TEST(DefaultValues, DefaultValuesAreOverriden) {
+  sycl::platform Plt{sycl::default_selector()};
+  if (Plt.is_host()) {
+    std::cerr << "Test is not supported on host, skipping\n";
+    return; // test is not supported on host.
+  }
+
+  sycl::unittest::PiMock Mock{Plt};
+  setupDefaultMockAPIs(Mock);
+
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  sycl::queue Queue{Dev};
+
+  const sycl::context Ctx = Queue.get_context();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
+  KernelBundle.set_specialization_constant<SpecConst1>(80);
+  auto ExecBundle = sycl::build(KernelBundle);
+  Queue.submit([&](sycl::handler &CGH) {
+    CGH.use_kernel_bundle(ExecBundle);
+    CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+  });
+
+  EXPECT_EQ(SpecConstVal0, 80);
+  EXPECT_EQ(SpecConstVal1, 8);
+}

From 9f60eb71e5f2c08a9dfcec8596b610f252edb21e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 7 May 2021 13:03:59 +0300
Subject: [PATCH 4/9] slight improvements

---
 sycl/source/CMakeLists.txt         |  3 ---
 sycl/unittests/helpers/PiImage.hpp | 17 +++++++++++++++++
 2 files changed, 17 insertions(+), 3 deletions(-)

diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
index d392fe5246e38..0c6535f21c6de 100644
--- a/sycl/source/CMakeLists.txt
+++ b/sycl/source/CMakeLists.txt
@@ -52,9 +52,6 @@ function(add_sycl_rt_library LIB_NAME)
   else()
     target_compile_options(${LIB_OBJ_NAME} PUBLIC
                            -fvisibility=hidden -fvisibility-inlines-hidden)
-
-    target_compile_options(${LIB_OBJ_NAME} PUBLIC
-                           -ggdb -O0)
     set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt")
     target_link_libraries(
       ${LIB_NAME} PRIVATE "-Wl,--version-script=${linker_script}")
diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp
index 6495a96ac8eb1..fa9d39cd91389 100644
--- a/sycl/unittests/helpers/PiImage.hpp
+++ b/sycl/unittests/helpers/PiImage.hpp
@@ -21,10 +21,17 @@ namespace unittest {
 class PiProperty {
 public:
   using NativeType = _pi_device_binary_property_struct;
+
+  /// Constructs a PI property.
+  ///
+  /// \param Name is a property name.
+  /// \param Data is a vector of raw property value bytes.
+  /// \param Type is one of pi_property_type values.
   PiProperty(const std::string &Name, std::vector<char> Data, uint32_t Type)
       : MName(Name), MData(std::move(Data)), MType(Type) {
     updateNativeType();
   }
+
   NativeType convertToNativeType() const { return MNative; }
 
   PiProperty(const PiProperty &Src) {
@@ -33,6 +40,7 @@ class PiProperty {
     MType = Src.MType;
     updateNativeType();
   }
+
   PiProperty &operator=(const PiProperty &Src) {
     MName = Src.MName;
     MData = Src.MData;
@@ -56,6 +64,7 @@ class PiProperty {
 class PiOffloadEntry {
 public:
   using NativeType = _pi_offload_entry_struct;
+
   PiOffloadEntry(const std::string &Name, std::vector<char> Data, int32_t Flags)
       : MName(Name), MData(std::move(Data)), MFlags(Flags) {
     updateNativeType();
@@ -125,6 +134,10 @@ class PiPropertySet {
 public:
   PiPropertySet() = default;
 
+  /// Adds a new array of properties to the set.
+  ///
+  /// \param Name is a property array name. See pi.h for list of known names.
+  /// \param Props is an array of property values.
   void insert(const std::string &Name, PiArray<PiProperty> Props) {
     MNames.push_back(Name);
     MMockProperties.push_back(std::move(Props));
@@ -329,6 +342,9 @@ PiProperty makeSpecConstant(std::vector<char> &ValData, const std::string &Name,
   return Prop;
 }
 
+/// Utility function to add specialization constants to property set.
+///
+/// This function overrides the default spec constant values.
 void addSpecConstants(PiArray<PiProperty> SpecConstants,
                       std::vector<char> ValData, PiPropertySet &Props) {
   Props.insert(__SYCL_PI_PROPERTY_SET_SPEC_CONST_MAP, std::move(SpecConstants));
@@ -341,6 +357,7 @@ void addSpecConstants(PiArray<PiProperty> SpecConstants,
                std::move(DefaultValues));
 }
 
+/// Utility function to generate offload entries for kernels without arguments.
 PiArray<PiOffloadEntry>
 makeEmptyKernels(std::initializer_list<std::string> KernelNames) {
   PiArray<PiOffloadEntry> Entries;

From 37d619e1ed8c3f6caa9e87a4076c79c8c087aa2a Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Fri, 7 May 2021 14:49:25 +0300
Subject: [PATCH 5/9] fix array size

---
 sycl/source/detail/device_image_impl.hpp |  4 ++--
 sycl/unittests/helpers/PiImage.hpp       | 21 +++++++++++++++------
 2 files changed, 17 insertions(+), 8 deletions(-)

diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index 8aa7910282c0f..49e6bf34556ff 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -253,8 +253,8 @@ class device_image_impl {
       if (HasDefaultValues) {
         pi::ByteArray DefValDescriptors =
             pi::DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray();
-        std::uninitialized_copy(&DefValDescriptors[0],
-                                &DefValDescriptors[0] + MSpecConstsBlob.size(),
+        std::uninitialized_copy(&DefValDescriptors[8],
+                                &DefValDescriptors[8] + MSpecConstsBlob.size(),
                                 MSpecConstsBlob.data());
       }
     }
diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp
index fa9d39cd91389..6306f9961b28f 100644
--- a/sycl/unittests/helpers/PiImage.hpp
+++ b/sycl/unittests/helpers/PiImage.hpp
@@ -305,13 +305,22 @@ PiProperty makeSpecConstant(std::vector<char> &ValData, const std::string &Name,
   std::uninitialized_copy(&PropByteArraySize, &PropByteArraySize + 8,
                           DescData.data());
 
+  if (ValData.empty())
+    ValData.resize(8); // Reserve first 8 bytes for array size.
   size_t PrevSize = ValData.size();
-  // Resize raw data blob to current size + offset of the last element + size of
-  // the last element.
-  ValData.resize(
-      PrevSize + *std::prev(Offsets.end()) +
-      sizeof(typename std::tuple_element<sizeof...(T) - 1,
-                                         decltype(DefaultValues)>::type));
+
+  {
+    // Resize raw data blob to current size + offset of the last element + size
+    // of the last element.
+    ValData.resize(
+        PrevSize + *std::prev(Offsets.end()) +
+        sizeof(typename std::tuple_element<sizeof...(T) - 1,
+                                           decltype(DefaultValues)>::type));
+    // Update raw data array size
+    uint64_t NewValSize = ValData.size();
+    std::uninitialized_copy(&NewValSize, &NewValSize + sizeof(uint64_t),
+                            ValData.data());
+  }
 
   auto FillData = [PrevOffset = 0, PrevSize, &ValData, &IDs, &Offsets,
                    &DescData](uint32_t Idx, const char *Begin,

From fee456d3b297c04ab124d60bfb8974b56c3d098f Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 11 May 2021 13:20:51 +0300
Subject: [PATCH 6/9] disable test for CUDA

---
 sycl/unittests/spec_constants/DefaultValues.cpp | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp
index 5ce5c279bf69d..30d099af69e63 100644
--- a/sycl/unittests/spec_constants/DefaultValues.cpp
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -212,6 +212,11 @@ TEST(DefaultValues, DefaultValuesAreSet) {
     return; // test is not supported on host.
   }
 
+  if (Plt.get_backen() == sycl::backend::cuda) {
+    std::cerr << "Test is not supported on CUDA platform, skipping\n";
+    return;
+  }
+
   sycl::unittest::PiMock Mock{Plt};
   setupDefaultMockAPIs(Mock);
 

From 894a99221b3b323a48a68abee6e7342fc76cedf0 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexbatashev@outlook.com>
Date: Tue, 11 May 2021 13:51:33 +0300
Subject: [PATCH 7/9] Update sycl/unittests/spec_constants/DefaultValues.cpp

Co-authored-by: Dmitry Vodopyanov <dmitryvodop@gmail.com>
---
 sycl/unittests/spec_constants/DefaultValues.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp
index 30d099af69e63..c536940006b7a 100644
--- a/sycl/unittests/spec_constants/DefaultValues.cpp
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -212,7 +212,7 @@ TEST(DefaultValues, DefaultValuesAreSet) {
     return; // test is not supported on host.
   }
 
-  if (Plt.get_backen() == sycl::backend::cuda) {
+  if (Plt.get_backend() == sycl::backend::cuda) {
     std::cerr << "Test is not supported on CUDA platform, skipping\n";
     return;
   }

From f1da5a7e4cd036b45f2950d57a79ef3d32c71651 Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitryvodop@gmail.com>
Date: Tue, 11 May 2021 16:19:26 +0300
Subject: [PATCH 8/9] Update DefaultValues.cpp

---
 sycl/unittests/spec_constants/DefaultValues.cpp | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp
index c536940006b7a..14fdc46b7dbbe 100644
--- a/sycl/unittests/spec_constants/DefaultValues.cpp
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -244,6 +244,11 @@ TEST(DefaultValues, DefaultValuesAreOverriden) {
     std::cerr << "Test is not supported on host, skipping\n";
     return; // test is not supported on host.
   }
+  
+  if (Plt.get_backend() == sycl::backend::cuda) {
+    std::cerr << "Test is not supported on CUDA platform, skipping\n";
+    return;
+  }
 
   sycl::unittest::PiMock Mock{Plt};
   setupDefaultMockAPIs(Mock);

From f7a75c82457dafa9bdf0b375726b218a478250bd Mon Sep 17 00:00:00 2001
From: Dmitry Vodopyanov <dmitryvodop@gmail.com>
Date: Tue, 11 May 2021 16:22:48 +0300
Subject: [PATCH 9/9] Fix clang-format

---
 sycl/unittests/spec_constants/DefaultValues.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/unittests/spec_constants/DefaultValues.cpp b/sycl/unittests/spec_constants/DefaultValues.cpp
index 14fdc46b7dbbe..655547f83d31d 100644
--- a/sycl/unittests/spec_constants/DefaultValues.cpp
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -244,7 +244,7 @@ TEST(DefaultValues, DefaultValuesAreOverriden) {
     std::cerr << "Test is not supported on host, skipping\n";
     return; // test is not supported on host.
   }
-  
+
   if (Plt.get_backend() == sycl::backend::cuda) {
     std::cerr << "Test is not supported on CUDA platform, skipping\n";
     return;