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..e06ae106e65e7 100644
--- a/sycl/include/CL/sycl/detail/pi.hpp
+++ b/sycl/include/CL/sycl/detail/pi.hpp
@@ -318,6 +318,14 @@ class DeviceBinaryImage {
   /// like:
   /// { ID5, 0, 4 }
   const PropertyRange &getSpecConstants() const { return SpecConstIDMap; }
+  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; }
   const PropertyRange &getKernelParamOptInfo() const {
     return KernelParamOptInfo;
diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index 697404d27a5db..49e6bf34556ff 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,12 +243,20 @@ 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);
+
+      if (HasDefaultValues) {
+        pi::ByteArray DefValDescriptors =
+            pi::DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray();
+        std::uninitialized_copy(&DefValDescriptors[8],
+                                &DefValDescriptors[8] + MSpecConstsBlob.size(),
+                                MSpecConstsBlob.data());
+      }
     }
   }
 
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..6306f9961b28f
--- /dev/null
+++ b/sycl/unittests/helpers/PiImage.hpp
@@ -0,0 +1,383 @@
+//==------------- 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;
+
+  /// 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) {
+    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;
+
+  /// 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));
+    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());
+
+  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));
+    // 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,
+                              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;
+}
+
+/// 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));
+
+  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));
+}
+
+/// Utility function to generate offload entries for kernels without arguments.
+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..655547f83d31d
--- /dev/null
+++ b/sycl/unittests/spec_constants/DefaultValues.cpp
@@ -0,0 +1,273 @@
+//==---- 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.
+  }
+
+  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);
+
+  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.
+  }
+
+  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);
+
+  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);
+}