diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index 7fbe0d1dffe76..87359ff6c779b 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -17,6 +17,13 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
+// This guard is needed because the libsycl.so can be compiled with C++ <=14
+// while the code requires C++17. This code is not supposed to be used by the
+// libsycl.so so it should not be a problem.
+#if __cplusplus > 201402L
+template <auto &S> struct specialization_id_name_generator {};
+#endif
+
 #ifndef __SYCL_DEVICE_ONLY__
 #define _Bool bool
 #endif
@@ -49,6 +56,14 @@ template <class Name> struct SpecConstantInfo {
   static constexpr const char *getName() { return ""; }
 };
 
+#if __cplusplus >= 201703L
+// Translates SYCL 2020 specialization constant type to its name.
+template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
+  return __builtin_unique_stable_name(
+      specialization_id_name_generator<SpecName>);
+}
+#endif
+
 #ifndef __SYCL_UNNAMED_LAMBDA__
 template <class KernelNameType> struct KernelInfo {
   static constexpr unsigned getNumParams() { return 0; }
diff --git a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp
index 461b1b9b52d2a..a79c496c3a225 100644
--- a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp
+++ b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp
@@ -37,13 +37,12 @@ SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID);
 // are not available.
 template <typename T>
 SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID,
-                                                      void *DefaultValue,
-                                                      void *RTBuffer);
+                                                      const void *DefaultValue,
+                                                      const void *RTBuffer);
 
 template <typename T>
-SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(const char *SymbolicID,
-                                                         void *DefaultValue,
-                                                         void *RTBuffer);
+SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(
+    const char *SymbolicID, const void *DefaultValue, const void *RTBuffer);
 
 // Request a fixed-size allocation in local address space at kernel scope.
 extern "C" SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t *
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 2523e917890e1..dc1a6a2d072dc 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -1048,7 +1048,7 @@ class __SYCL_EXPORT handler {
   }
 
   template <auto &SpecName>
-  typename std::remove_reference_t<decltype(SpecName)>::type
+  typename std::remove_reference_t<decltype(SpecName)>::value_type
   get_specialization_constant() const {
 
     std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp
index 41ab9c3405899..e8c93a02f079a 100644
--- a/sycl/include/CL/sycl/kernel_bundle.hpp
+++ b/sycl/include/CL/sycl/kernel_bundle.hpp
@@ -168,6 +168,16 @@ class __SYCL_EXPORT kernel_bundle_plain {
   // \returns an iterator to the last device image kernel_bundle contains
   const device_image_plain *end() const;
 
+  bool has_specialization_constant_impl(const char *SpecName) const noexcept;
+
+  void set_specialization_constant_impl(const char *SpecName,
+                                        void *Value) noexcept;
+
+  void get_specialization_constant_impl(const char *SpecName, void *Value) const
+      noexcept;
+
+  bool is_specialization_constant_set(const char *SpecName) const noexcept;
+
   detail::KernelBundleImplPtr impl;
 };
 
@@ -247,9 +257,8 @@ class kernel_bundle : public detail::kernel_bundle_plain {
   /// \returns true if any device image in the kernel_bundle uses specialization
   /// constant whose address is SpecName
   template <auto &SpecName> bool has_specialization_constant() const noexcept {
-    throw sycl::runtime_error(
-        "kernel_bundle::has_specialization_constant is not implemented yet",
-        PI_INVALID_OPERATION);
+    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    return has_specialization_constant_impl(SpecSymName);
   }
 
   /// Sets the value of the specialization constant whose address is SpecName
@@ -259,20 +268,27 @@ class kernel_bundle : public detail::kernel_bundle_plain {
             typename = detail::enable_if_t<_State == bundle_state::input>>
   void set_specialization_constant(
       typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
-    (void)Value;
-    throw sycl::runtime_error(
-        "kernel_bundle::set_specialization_constant is not implemented yet",
-        PI_INVALID_OPERATION);
+    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    set_specialization_constant_impl(SpecSymName, &Value);
   }
 
-  /// The value of the specialization constant whose address is SpecName for
-  /// this kernel bundle.
+  /// \returns the value of the specialization constant whose address is
+  /// SpecName for this kernel bundle.
   template <auto &SpecName>
   typename std::remove_reference_t<decltype(SpecName)>::value_type
   get_specialization_constant() const {
-    throw sycl::runtime_error(
-        "kernel_bundle::get_specialization_constant is not implemented yet",
-        PI_INVALID_OPERATION);
+    const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
+    if (!is_specialization_constant_set(SpecSymName))
+      return SpecName.getDefaultValue();
+
+    using SCType =
+        typename std::remove_reference_t<decltype(SpecName)>::value_type;
+
+    std::array<char *, sizeof(SCType)> RetValue;
+
+    get_specialization_constant_impl(SpecSymName, RetValue.data());
+
+    return *reinterpret_cast<SCType *>(RetValue.data());
   }
 #endif
 
diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp
index 0a3f77381b364..fd5415b802e28 100644
--- a/sycl/include/CL/sycl/kernel_handler.hpp
+++ b/sycl/include/CL/sycl/kernel_handler.hpp
@@ -8,19 +8,13 @@
 
 #pragma once
 
-__SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl {
-namespace detail {
-
-// This guard is needed because the libsycl.so can compiled with C++ <=14
-// while the code requires C++17. This code is not supposed to be used by the
-// libsycl.so so it should not be a problem.
-#if __cplusplus > 201402L
-template <auto &S> struct specialization_id_name_generator {};
-#endif
+#include <CL/sycl/detail/kernel_desc.hpp>
+#include <CL/sycl/exception.hpp>
 
-} // namespace detail
+#include <type_traits>
 
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
 /// Reading the value of a specialization constant
 ///
 /// \ingroup sycl_api
@@ -28,7 +22,8 @@ class kernel_handler {
 public:
 #if __cplusplus > 201402L
   template <auto &S>
-  typename std::remove_reference_t<decltype(S)> get_specialization_constant() {
+  typename std::remove_reference_t<decltype(S)>::value_type
+  get_specialization_constant() {
 #ifdef __SYCL_DEVICE_ONLY__
     return getSpecializationConstantOnDevice<S>();
 #else
@@ -48,16 +43,20 @@ class kernel_handler {
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
-  template <auto &S, typename T = std::remove_reference_t<decltype(S)>,
-            std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
+  template <
+      auto &S,
+      typename T = typename std::remove_reference_t<decltype(S)>::value_type,
+      std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
   T getSpecializationConstantOnDevice() {
     const char *SymbolicID = __builtin_unique_stable_name(
         detail::specialization_id_name_generator<S>);
     return __sycl_getScalar2020SpecConstantValue<T>(
         SymbolicID, &S, MSpecializationConstantsBuffer);
   }
-  template <auto &S, typename T = std::remove_reference_t<decltype(S)>,
-            std::enable_if_t<std::is_compound_v<T>> * = nullptr>
+  template <
+      auto &S,
+      typename T = typename std::remove_reference_t<decltype(S)>::value_type,
+      std::enable_if_t<std::is_compound_v<T>> * = nullptr>
   T getSpecializationConstantOnDevice() {
     const char *SymbolicID = __builtin_unique_stable_name(
         detail::specialization_id_name_generator<S>);
diff --git a/sycl/include/CL/sycl/specialization_id.hpp b/sycl/include/CL/sycl/specialization_id.hpp
index cc4f25d8ebad6..0023507a7702a 100644
--- a/sycl/include/CL/sycl/specialization_id.hpp
+++ b/sycl/include/CL/sycl/specialization_id.hpp
@@ -28,6 +28,9 @@ template <typename T> class specialization_id {
   specialization_id &operator=(specialization_id &&rhs) = delete;
 
 private:
+  template <bundle_state State> friend class kernel_bundle;
+  T getDefaultValue() const noexcept { return MDefaultValue; }
+
   T MDefaultValue;
 };
 
diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
index dbfc545bcea3d..9fa83db6aa150 100644
--- a/sycl/source/detail/device_image_impl.hpp
+++ b/sycl/source/detail/device_image_impl.hpp
@@ -39,7 +39,9 @@ class device_image_impl {
                     std::vector<kernel_id> KernelIDs, RT::PiProgram Program)
       : MBinImage(BinImage), MContext(std::move(Context)),
         MDevices(std::move(Devices)), MState(State), MProgram(Program),
-        MKernelIDs(std::move(KernelIDs)) {}
+        MKernelIDs(std::move(KernelIDs)) {
+    updateSpecConstSymMap();
+  }
 
   bool has_kernel(const kernel_id &KernelIDCand) const noexcept {
     return std::binary_search(MKernelIDs.begin(), MKernelIDs.end(),
@@ -60,7 +62,11 @@ class device_image_impl {
   }
 
   bool has_specialization_constants() const noexcept {
-    return !MSpecConstsBlob.empty();
+    // Lock the mutex to prevent when one thread in the middle of writing a
+    // new value while another thread is reading the value to pass it to
+    // JIT compiler.
+    const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
+    return !MSpecConstSymMap.empty();
   }
 
   bool all_specialization_constant_native() const noexcept {
@@ -72,45 +78,69 @@ class device_image_impl {
   // for this spec const should be.
   struct SpecConstDescT {
     unsigned int ID = 0;
-    unsigned int Offset = 0;
+    unsigned int CompositeOffset = 0;
+    unsigned int Size = 0;
+    unsigned int BlobOffset = 0;
     bool IsSet = false;
   };
 
-  bool has_specialization_constant(unsigned int SpecID) const noexcept {
-    return std::any_of(MSpecConstDescs.begin(), MSpecConstDescs.end(),
-                       [SpecID](const SpecConstDescT &SpecConstDesc) {
-                         return SpecConstDesc.ID == SpecID;
-                       });
-  }
-
-  void set_specialization_constant_raw_value(unsigned int SpecID,
-                                             const void *Value,
-                                             size_t ValueSize) noexcept {
-    for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs)
-      if (SpecConstDesc.ID == SpecID) {
-        // Lock the mutex to prevent when one thread in the middle of writing a
-        // new value while another thread is reading the value to pass it to
-        // JIT compiler.
-        const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
-        std::memcpy(MSpecConstsBlob.data() + SpecConstDesc.Offset, Value,
-                    ValueSize);
-        return;
-      }
+  bool has_specialization_constant(const char *SpecName) const noexcept {
+    // Lock the mutex to prevent when one thread in the middle of writing a
+    // new value while another thread is reading the value to pass it to
+    // JIT compiler.
+    const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
+    return MSpecConstSymMap.count(SpecName) != 0;
   }
 
-  void get_specialization_constant_raw_value(unsigned int SpecID,
-                                             void *ValueRet,
-                                             size_t ValueSize) const noexcept {
-    for (const SpecConstDescT &SpecConstDesc : MSpecConstDescs)
-      if (SpecConstDesc.ID == SpecID) {
-        // Lock the mutex to prevent when one thread in the middle of writing a
-        // new value while another thread is reading the value to pass it to
-        // JIT compiler.
-        const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
-        std::memcpy(ValueRet, MSpecConstsBlob.data() + SpecConstDesc.Offset,
-                    ValueSize);
-        return;
-      }
+  void set_specialization_constant_raw_value(const char *SpecName,
+                                             const void *Value) noexcept {
+    // Lock the mutex to prevent when one thread in the middle of writing a
+    // new value while another thread is reading the value to pass it to
+    // JIT compiler.
+    const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
+
+    if (MSpecConstSymMap.count(std::string{SpecName}) == 0)
+      return;
+
+    std::vector<SpecConstDescT> &Descs =
+        MSpecConstSymMap[std::string{SpecName}];
+    for (SpecConstDescT &Desc : Descs) {
+      Desc.IsSet = true;
+      std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset,
+                  static_cast<const char *>(Value) + Desc.CompositeOffset,
+                  Desc.Size);
+    }
+  }
+
+  void get_specialization_constant_raw_value(const char *SpecName,
+                                             void *ValueRet) const noexcept {
+    assert(is_specialization_constant_set(SpecName));
+    // Lock the mutex to prevent when one thread in the middle of writing a
+    // new value while another thread is reading the value to pass it to
+    // JIT compiler.
+    const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
+
+    // operator[] can't be used here, since it's not marked as const
+    const std::vector<SpecConstDescT> &Descs =
+        MSpecConstSymMap.at(std::string{SpecName});
+    for (const SpecConstDescT &Desc : Descs) {
+
+      std::memcpy(static_cast<char *>(ValueRet) + Desc.CompositeOffset,
+                  MSpecConstsBlob.data() + Desc.BlobOffset, Desc.Size);
+    }
+  }
+
+  bool is_specialization_constant_set(const char *SpecName) const noexcept {
+    // Lock the mutex to prevent when one thread in the middle of writing a
+    // new value while another thread is reading the value to pass it to
+    // JIT compiler.
+    const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
+    if (MSpecConstSymMap.count(std::string{SpecName}) == 0)
+      return false;
+
+    const std::vector<SpecConstDescT> &Descs =
+        MSpecConstSymMap.at(std::string{SpecName});
+    return Descs.front().IsSet;
   }
 
   bundle_state get_state() const noexcept { return MState; }
@@ -137,8 +167,13 @@ class device_image_impl {
     return MSpecConstsBlob;
   }
 
-  std::vector<SpecConstDescT> &get_spec_const_offsets_ref() noexcept {
-    return MSpecConstDescs;
+  const std::map<std::string, std::vector<SpecConstDescT>> &
+  get_spec_const_data_ref() const noexcept {
+    return MSpecConstSymMap;
+  }
+
+  std::mutex &get_spec_const_data_lock() noexcept {
+    return MSpecConstAccessMtx;
   }
 
   ~device_image_impl() {
@@ -150,6 +185,49 @@ class device_image_impl {
   }
 
 private:
+  void updateSpecConstSymMap() {
+    if (MBinImage) {
+      const pi::DeviceBinaryImage::PropertyRange &SCRange =
+          MBinImage->getSpecConstants();
+      using SCItTy = pi::DeviceBinaryImage::PropertyRange::ConstIterator;
+
+      // This variable is used to calculate spec constant value offset in a
+      // flat byte array.
+      unsigned BlobOffset = 0;
+      for (SCItTy SCIt : SCRange) {
+        const char *SCName = (*SCIt)->Name;
+
+        pi::ByteArray Descriptors =
+            pi::DeviceBinaryProperty(*SCIt).asByteArray();
+        assert(Descriptors.size() > 8 && "Unexpected property size");
+
+        // Expected layout is vector of 3-component tuples (flattened into a
+        // vector of scalars), where each tuple consists of: ID of a scalar spec
+        // constant, (which might be a member of the composite); offset, which
+        // is used to calculate location of scalar member within the composite
+        // or zero for scalar spec constants; size of a spec constant
+        constexpr size_t NumElements = 3;
+        assert(((Descriptors.size() - 8) / sizeof(std::uint32_t)) %
+                       NumElements ==
+                   0 &&
+               "unexpected layout of composite spec const descriptors");
+        auto *It = reinterpret_cast<const std::uint32_t *>(&Descriptors[8]);
+        auto *End = reinterpret_cast<const std::uint32_t *>(&Descriptors[0] +
+                                                            Descriptors.size());
+        while (It != End) {
+          // The map is not locked here because updateSpecConstSymMap() is only
+          // supposed to be called from c'tor.
+          MSpecConstSymMap[std::string{SCName}].push_back(
+              SpecConstDescT{/*ID*/ It[0], /*CompositeOffset*/ It[1],
+                             /*Size*/ It[2], BlobOffset});
+          BlobOffset += /*Size*/ It[2];
+          It += NumElements;
+        }
+      }
+      MSpecConstsBlob.resize(BlobOffset);
+    }
+  }
+
   const RTDeviceBinaryImage *MBinImage = nullptr;
   context MContext;
   std::vector<device> MDevices;
@@ -166,8 +244,9 @@ class device_image_impl {
   // Binary blob which can have values of all specialization constants in the
   // image
   std::vector<unsigned char> MSpecConstsBlob;
-  // Contains list of spec ID + their offsets in the MSpecConstsBlob
-  std::vector<SpecConstDescT> MSpecConstDescs;
+  // Contains map of spec const names to their descriptions + offsets in
+  // the MSpecConstsBlob
+  std::map<std::string, std::vector<SpecConstDescT>> MSpecConstSymMap;
 };
 
 } // namespace detail
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index a7a5a93990725..984fd6651189e 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -238,33 +238,39 @@ class kernel_bundle_impl {
                        });
   }
 
-  bool has_specialization_constant(unsigned int SpecID) const noexcept {
+  bool has_specialization_constant(const char *SpecName) const noexcept {
     return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
-                       [SpecID](const device_image_plain &DeviceImage) {
+                       [SpecName](const device_image_plain &DeviceImage) {
                          return getSyclObjImpl(DeviceImage)
-                             ->has_specialization_constant(SpecID);
+                             ->has_specialization_constant(SpecName);
                        });
   }
 
-  void set_specialization_constant_raw_value(unsigned int SpecID,
-                                             const void *Value,
-                                             size_t ValueSize) {
+  void set_specialization_constant_raw_value(const char *SpecName,
+                                             const void *Value) noexcept {
+    // TODO add support for specialization constants, that are missing in any
+    // device image.
     for (const device_image_plain &DeviceImage : MDeviceImages)
       getSyclObjImpl(DeviceImage)
-          ->set_specialization_constant_raw_value(SpecID, Value, ValueSize);
+          ->set_specialization_constant_raw_value(SpecName, Value);
   }
 
-  const void *get_specialization_constant_raw_value(unsigned int SpecID,
-                                                    void *ValueRet,
-                                                    size_t ValueSize) const {
+  void get_specialization_constant_raw_value(const char *SpecName,
+                                             void *ValueRet) const noexcept {
     for (const device_image_plain &DeviceImage : MDeviceImages)
-      if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecID)) {
+      if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) {
         getSyclObjImpl(DeviceImage)
-            ->get_specialization_constant_raw_value(SpecID, ValueRet,
-                                                    ValueSize);
+            ->get_specialization_constant_raw_value(SpecName, ValueRet);
+        return;
       }
+  }
 
-    return nullptr;
+  bool is_specialization_constant_set(const char *SpecName) const noexcept {
+    return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
+                       [SpecName](const device_image_plain &DeviceImage) {
+                         return getSyclObjImpl(DeviceImage)
+                             ->is_specialization_constant_set(SpecName);
+                       });
   }
 
   const device_image_plain *begin() const { return &MDeviceImages.front(); }
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index dd00f357032a2..f04b6cc2aa74e 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1540,17 +1540,22 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage,
     const std::vector<unsigned char> &SpecConstsBlob =
         InputImpl->get_spec_const_blob_ref();
 
-    std::vector<device_image_impl::SpecConstDescT> &SpecConstOffsets =
-        InputImpl->get_spec_const_offsets_ref();
-
-    unsigned int PrevOffset = 0;
-    for (const device_image_impl::SpecConstDescT &SpecIDDesc :
-         SpecConstOffsets) {
-
-      Plugin.call<PiApiKind::piextProgramSetSpecializationConstant>(
-          NativePrg, SpecIDDesc.ID, SpecIDDesc.Offset - PrevOffset,
-          SpecConstsBlob.data() + SpecIDDesc.Offset);
-      PrevOffset = SpecIDDesc.Offset;
+    {
+      std::lock_guard<std::mutex> Lock{InputImpl->get_spec_const_data_lock()};
+      const std::map<std::string,
+                     std::vector<device_image_impl::SpecConstDescT>>
+          &SpecConstData = InputImpl->get_spec_const_data_ref();
+
+      for (const auto &DescPair : SpecConstData) {
+        for (const device_image_impl::SpecConstDescT &SpecIDDesc :
+             DescPair.second) {
+          if (SpecIDDesc.IsSet) {
+            Plugin.call<PiApiKind::piextProgramSetSpecializationConstant>(
+                NativePrg, SpecIDDesc.ID, SpecIDDesc.Size,
+                SpecConstsBlob.data() + SpecIDDesc.BlobOffset);
+          }
+        }
+      }
     }
 
     ProgramPtr ProgramManaged(
diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
index 7f35903792bf7..b5ea314428d38 100644
--- a/sycl/source/kernel_bundle.cpp
+++ b/sycl/source/kernel_bundle.cpp
@@ -83,6 +83,27 @@ bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID,
   return impl->has_kernel(KernelID, Dev);
 }
 
+bool kernel_bundle_plain::has_specialization_constant_impl(
+    const char *SpecName) const noexcept {
+  return impl->has_specialization_constant(SpecName);
+}
+
+void kernel_bundle_plain::set_specialization_constant_impl(
+    const char *SpecName, void *Value) noexcept {
+  impl->set_specialization_constant_raw_value(SpecName, Value);
+}
+
+void kernel_bundle_plain::get_specialization_constant_impl(const char *SpecName,
+                                                           void *Value) const
+    noexcept {
+  impl->get_specialization_constant_raw_value(SpecName, Value);
+}
+
+bool kernel_bundle_plain::is_specialization_constant_set(
+    const char *SpecName) const noexcept {
+  return impl->is_specialization_constant_set(SpecName);
+}
+
 ////////////////////////////
 ///// free functions
 ///////////////////////////
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 4601549761311..a21cab277a49b 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3786,6 +3786,7 @@ _ZN2cl4sycl6detail18stringifyErrorCodeEi
 _ZN2cl4sycl6detail19convertChannelOrderE23_pi_image_channel_order
 _ZN2cl4sycl6detail19convertChannelOrderENS0_19image_channel_orderE
 _ZN2cl4sycl6detail19getImageElementSizeEhNS0_18image_channel_typeE
+_ZN2cl4sycl6detail19kernel_bundle_plain32set_specialization_constant_implEPKcPv
 _ZN2cl4sycl6detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE
 _ZN2cl4sycl6detail20getDeviceFromHandlerERNS0_7handlerE
 _ZN2cl4sycl6detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE
@@ -3976,7 +3977,10 @@ _ZNK2cl4sycl6detail19kernel_bundle_plain11get_backendEv
 _ZNK2cl4sycl6detail19kernel_bundle_plain11get_contextEv
 _ZNK2cl4sycl6detail19kernel_bundle_plain11get_devicesEv
 _ZNK2cl4sycl6detail19kernel_bundle_plain14get_kernel_idsEv
+_ZNK2cl4sycl6detail19kernel_bundle_plain30is_specialization_constant_setEPKc
 _ZNK2cl4sycl6detail19kernel_bundle_plain30native_specialization_constantEv
+_ZNK2cl4sycl6detail19kernel_bundle_plain32get_specialization_constant_implEPKcPv
+_ZNK2cl4sycl6detail19kernel_bundle_plain32has_specialization_constant_implEPKc
 _ZNK2cl4sycl6detail19kernel_bundle_plain33contains_specialization_constantsEv
 _ZNK2cl4sycl6detail19kernel_bundle_plain3endEv
 _ZNK2cl4sycl6detail19kernel_bundle_plain5beginEv
diff --git a/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp
new file mode 100644
index 0000000000000..e73d2944dba66
--- /dev/null
+++ b/sycl/test/on-device/basic_tests/specialization_constants/host_apis.cpp
@@ -0,0 +1,76 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %t.out
+
+// UNSUPPORTED: cuda
+
+#include <sycl/sycl.hpp>
+
+#include <cmath>
+
+class Kernel1Name;
+class Kernel2Name;
+
+struct TestStruct {
+  int a;
+  int b;
+};
+
+const static sycl::specialization_id<int> SpecConst1{42};
+const static sycl::specialization_id<int> SpecConst2{42};
+const static sycl::specialization_id<TestStruct> SpecConst3{TestStruct{42, 42}};
+const static sycl::specialization_id<short> SpecConst4{42};
+
+int main() {
+  sycl::queue Q;
+
+  // No support for host device so far
+  if (Q.is_host())
+    return 0;
+
+  // The code is needed to just have device images in the executable
+  if (0) {
+    Q.submit([](sycl::handler &CGH) { CGH.single_task<Kernel1Name>([] {}); });
+    Q.submit([](sycl::handler &CGH) { CGH.single_task<Kernel2Name>([] {}); });
+  }
+
+  const sycl::context Ctx = Q.get_context();
+  const sycl::device Dev = Q.get_device();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(Ctx, {Dev});
+
+  assert(KernelBundle.contains_specialization_constants() == true);
+  assert(KernelBundle.has_specialization_constant<SpecConst1>() == false);
+  assert(KernelBundle.has_specialization_constant<SpecConst2>() == true);
+  KernelBundle.set_specialization_constant<SpecConst2>(1);
+  {
+    auto ExecBundle = sycl::build(KernelBundle);
+    sycl::buffer<int, 1> Buf{sycl::range{1}};
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.use_kernel_bundle(ExecBundle);
+      auto Acc = Buf.get_access<sycl::access::mode::read_write>(CGH);
+      CGH.single_task<class Kernel3Name>([=](sycl::kernel_handler KH) {
+        Acc[0] = KH.get_specialization_constant<SpecConst2>();
+      });
+    });
+    auto Acc = Buf.get_access<sycl::access::mode::read>();
+    assert(Acc[0] == 1);
+  }
+
+  {
+    sycl::buffer<TestStruct, 1> Buf{sycl::range{1}};
+    Q.submit([&](sycl::handler &CGH) {
+      auto Acc = Buf.get_access<sycl::access::mode::read_write>(CGH);
+      CGH.set_specialization_constant<SpecConst3>(TestStruct{1, 2});
+      const auto SC = CGH.get_specialization_constant<SpecConst4>();
+      assert(SC == 42);
+      CGH.single_task<class Kernel4Name>([=](sycl::kernel_handler KH) {
+        Acc[0] = KH.get_specialization_constant<SpecConst3>();
+      });
+    });
+    auto Acc = Buf.get_access<sycl::access::mode::read>();
+    assert(Acc[0].a == 1 && Acc[0].b == 2);
+  }
+
+  return 0;
+}