diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp
index 62eb7806e4517..fa1f95f8c5a69 100644
--- a/sycl/include/sycl/nd_item.hpp
+++ b/sycl/include/sycl/nd_item.hpp
@@ -47,7 +47,9 @@ template <int Dimensions = 1> class nd_item {
 public:
   static constexpr int dimensions = Dimensions;
 
-  id<Dimensions> get_global_id() const {
+  nd_item() = delete;
+
+  id<Dimensions> get_global_id() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInGlobalInvocationId<Dimensions, id<Dimensions>>();
 #else
@@ -55,13 +57,13 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
+  size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const noexcept {
     size_t Id = get_global_id()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Id);
     return Id;
   }
 
-  size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const {
+  size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const noexcept {
     size_t LinId = 0;
     id<Dimensions> Index = get_global_id();
     range<Dimensions> Extent = get_global_range();
@@ -78,7 +80,7 @@ template <int Dimensions = 1> class nd_item {
     return LinId;
   }
 
-  id<Dimensions> get_local_id() const {
+  id<Dimensions> get_local_id() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
 #else
@@ -86,13 +88,13 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
+  size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const noexcept {
     size_t Id = get_local_id()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Id);
     return Id;
   }
 
-  size_t get_local_linear_id() const {
+  size_t get_local_linear_id() const noexcept {
     size_t LinId = 0;
     id<Dimensions> Index = get_local_id();
     range<Dimensions> Extent = get_local_range();
@@ -108,7 +110,7 @@ template <int Dimensions = 1> class nd_item {
     return LinId;
   }
 
-  group<Dimensions> get_group() const {
+  group<Dimensions> get_group() const noexcept {
     // TODO: ideally Group object should be stateless and have a contructor with
     // no arguments.
     return detail::Builder::createGroup(get_global_range(), get_local_range(),
@@ -116,15 +118,15 @@ template <int Dimensions = 1> class nd_item {
   }
 
   // Out-of-class definition in sub_group.hpp
-  sub_group get_sub_group() const;
+  sub_group get_sub_group() const noexcept;
 
-  size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
+  size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const noexcept {
     size_t Id = get_group_id()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Id);
     return Id;
   }
 
-  size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const {
+  size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const noexcept {
     size_t LinId = 0;
     id<Dimensions> Index = get_group_id();
     range<Dimensions> Extent = get_group_range();
@@ -140,7 +142,7 @@ template <int Dimensions = 1> class nd_item {
     return LinId;
   }
 
-  range<Dimensions> get_group_range() const {
+  range<Dimensions> get_group_range() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInNumWorkgroups<Dimensions, range<Dimensions>>();
 #else
@@ -148,13 +150,13 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
+  size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const noexcept {
     size_t Range = get_group_range()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Range);
     return Range;
   }
 
-  range<Dimensions> get_global_range() const {
+  range<Dimensions> get_global_range() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInGlobalSize<Dimensions, range<Dimensions>>();
 #else
@@ -162,13 +164,13 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  size_t get_global_range(int Dimension) const {
+  size_t get_global_range(int Dimension) const noexcept {
     size_t Val = get_global_range()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Val);
     return Val;
   }
 
-  range<Dimensions> get_local_range() const {
+  range<Dimensions> get_local_range() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInWorkgroupSize<Dimensions, range<Dimensions>>();
 #else
@@ -176,14 +178,14 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  size_t get_local_range(int Dimension) const {
+  size_t get_local_range(int Dimension) const noexcept {
     size_t Id = get_local_range()[Dimension];
     __SYCL_ASSUME_ID_RANGE(Id);
     return Id;
   }
 
   __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
-  id<Dimensions> get_offset() const {
+  id<Dimensions> get_offset() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>();
 #else
@@ -191,7 +193,7 @@ template <int Dimensions = 1> class nd_item {
 #endif
   }
 
-  nd_range<Dimensions> get_nd_range() const {
+  nd_range<Dimensions> get_nd_range() const noexcept {
     return nd_range<Dimensions>(get_global_range(), get_local_range(),
                                 get_offset());
   }
@@ -248,7 +250,7 @@ template <int Dimensions = 1> class nd_item {
                                               src,
                                           [[maybe_unused]] size_t numElements,
                                           [[maybe_unused]] size_t srcStride)
-      const {
+      const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
         __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -274,7 +276,7 @@ template <int Dimensions = 1> class nd_item {
                                           [[maybe_unused]] local_ptr<dataT> src,
                                           [[maybe_unused]] size_t numElements,
                                           [[maybe_unused]] size_t destStride)
-      const {
+      const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
         __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -299,7 +301,7 @@ template <int Dimensions = 1> class nd_item {
   async_work_group_copy([[maybe_unused]] decorated_local_ptr<DestDataT> dest,
                         [[maybe_unused]] decorated_global_ptr<SrcDataT> src,
                         [[maybe_unused]] size_t numElements,
-                        [[maybe_unused]] size_t srcStride) const {
+                        [[maybe_unused]] size_t srcStride) const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
         __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -324,7 +326,7 @@ template <int Dimensions = 1> class nd_item {
   async_work_group_copy([[maybe_unused]] decorated_global_ptr<DestDataT> dest,
                         [[maybe_unused]] decorated_local_ptr<SrcDataT> src,
                         [[maybe_unused]] size_t numElements,
-                        [[maybe_unused]] size_t destStride) const {
+                        [[maybe_unused]] size_t destStride) const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
         __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -352,7 +354,7 @@ template <int Dimensions = 1> class nd_item {
                                                     access::decorated::legacy>
                                               Src,
                                           size_t NumElements,
-                                          size_t Stride) const {
+                                          size_t Stride) const noexcept {
     static_assert(sizeof(bool) == sizeof(uint8_t),
                   "Async copy to/from bool memory is not supported.");
     auto DestP = multi_ptr<uint8_t, DestS, access::decorated::legacy>(
@@ -378,7 +380,7 @@ template <int Dimensions = 1> class nd_item {
                                                     access::decorated::legacy>
                                               Src,
                                           size_t NumElements,
-                                          size_t Stride) const {
+                                          size_t Stride) const noexcept {
     static_assert(sizeof(bool) == sizeof(uint8_t),
                   "Async copy to/from bool memory is not supported.");
     using VecT = detail::change_base_type_t<T, uint8_t>;
@@ -401,7 +403,7 @@ template <int Dimensions = 1> class nd_item {
                    device_event>
   async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
                         multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
-                        size_t NumElements, size_t Stride) const {
+                        size_t NumElements, size_t Stride) const noexcept {
     static_assert(sizeof(bool) == sizeof(uint8_t),
                   "Async copy to/from bool memory is not supported.");
     using QualSrcT =
@@ -429,7 +431,7 @@ template <int Dimensions = 1> class nd_item {
                    device_event>
   async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
                         multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
-                        size_t NumElements, size_t Stride) const {
+                        size_t NumElements, size_t Stride) const noexcept {
     static_assert(sizeof(bool) == sizeof(uint8_t),
                   "Async copy to/from bool memory is not supported.");
     using VecT = detail::change_base_type_t<DestT, uint8_t>;
@@ -455,7 +457,7 @@ template <int Dimensions = 1> class nd_item {
   __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
   device_event
       async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
-                            size_t numElements) const {
+                            size_t numElements) const noexcept {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 
@@ -468,7 +470,7 @@ template <int Dimensions = 1> class nd_item {
   __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
   device_event
       async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
-                            size_t numElements) const {
+                            size_t numElements) const noexcept {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 
@@ -483,7 +485,7 @@ template <int Dimensions = 1> class nd_item {
       std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
   async_work_group_copy(decorated_local_ptr<DestDataT> dest,
                         decorated_global_ptr<SrcDataT> src,
-                        size_t numElements) const {
+                        size_t numElements) const noexcept {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 
@@ -498,11 +500,12 @@ template <int Dimensions = 1> class nd_item {
       std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
   async_work_group_copy(decorated_global_ptr<DestDataT> dest,
                         decorated_local_ptr<SrcDataT> src,
-                        size_t numElements) const {
+                        size_t numElements) const noexcept {
     return async_work_group_copy(dest, src, numElements, 1);
   }
 
-  template <typename... eventTN> void wait_for(eventTN... events) const {
+  template <typename... eventTN>
+  void wait_for(eventTN... events) const noexcept {
     (events.wait(), ...);
   }
 
@@ -517,16 +520,20 @@ template <int Dimensions = 1> class nd_item {
   nd_item &operator=(const nd_item &rhs) = default;
   nd_item &operator=(nd_item &&rhs) = default;
 
-  bool operator==(const nd_item &) const { return true; }
-  bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
+  friend bool operator==(const nd_item &, const nd_item &) noexcept {
+    return true;
+  }
+
+  friend bool operator!=(const nd_item &lhs, const nd_item &rhs) noexcept {
+    return !(lhs == rhs);
+  }
 
 protected:
   friend class detail::Builder;
-  nd_item() {}
   nd_item(const item<Dimensions, true> &, const item<Dimensions, false> &,
           const group<Dimensions> &) {}
 
-  id<Dimensions> get_group_id() const {
+  id<Dimensions> get_group_id() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initBuiltInWorkgroupId<Dimensions, id<Dimensions>>();
 #else
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
index d25ea31a1b88a..443c7c62c0c12 100644
--- a/sycl/include/sycl/sub_group.hpp
+++ b/sycl/include/sycl/sub_group.hpp
@@ -197,7 +197,8 @@ struct sub_group {
   sub_group() = default;
 };
 
-template <int Dimensions> sub_group nd_item<Dimensions>::get_sub_group() const {
+template <int Dimensions>
+sub_group nd_item<Dimensions>::get_sub_group() const noexcept {
   return sub_group();
 }
 
diff --git a/sycl/test/basic_tests/nd_item_interface.cpp b/sycl/test/basic_tests/nd_item_interface.cpp
new file mode 100644
index 0000000000000..7c0f1385767af
--- /dev/null
+++ b/sycl/test/basic_tests/nd_item_interface.cpp
@@ -0,0 +1,136 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Wno-deprecated-declarations -fsyntax-only %s
+
+//==--------- nd_item_interface.cpp - nd_item interface 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <sycl/multi_ptr.hpp>
+#include <sycl/nd_item.hpp>
+#include <sycl/sub_group.hpp>
+#include <sycl/vector.hpp>
+
+#include <cstddef>
+#include <type_traits>
+#include <utility>
+
+template <typename T, typename = void>
+struct has_member_equal : std::false_type {};
+
+template <typename T>
+struct has_member_equal<
+    T, std::void_t<decltype(std::declval<const T &>().operator==(
+           std::declval<const T &>()))>> : std::true_type {};
+
+template <typename T, typename = void>
+struct has_member_not_equal : std::false_type {};
+
+template <typename T>
+struct has_member_not_equal<
+    T, std::void_t<decltype(std::declval<const T &>().operator!=(
+           std::declval<const T &>()))>> : std::true_type {};
+
+using Item = sycl::nd_item<1>;
+
+static_assert(!std::is_default_constructible_v<Item>);
+
+static_assert(noexcept(std::declval<const Item &>().get_global_id()));
+static_assert(noexcept(std::declval<const Item &>().get_global_id(0)));
+static_assert(noexcept(std::declval<const Item &>().get_global_linear_id()));
+static_assert(noexcept(std::declval<const Item &>().get_local_id()));
+static_assert(noexcept(std::declval<const Item &>().get_local_id(0)));
+static_assert(noexcept(std::declval<const Item &>().get_local_linear_id()));
+static_assert(noexcept(std::declval<const Item &>().get_group()));
+static_assert(noexcept(std::declval<const Item &>().get_sub_group()));
+static_assert(noexcept(std::declval<const Item &>().get_group(0)));
+static_assert(noexcept(std::declval<const Item &>().get_group_linear_id()));
+static_assert(noexcept(std::declval<const Item &>().get_group_range()));
+static_assert(noexcept(std::declval<const Item &>().get_group_range(0)));
+static_assert(noexcept(std::declval<const Item &>().get_global_range()));
+static_assert(noexcept(std::declval<const Item &>().get_global_range(0)));
+static_assert(noexcept(std::declval<const Item &>().get_local_range()));
+static_assert(noexcept(std::declval<const Item &>().get_local_range(0)));
+static_assert(noexcept(std::declval<const Item &>().get_offset()));
+static_assert(noexcept(std::declval<const Item &>().get_nd_range()));
+
+static_assert(!has_member_equal<Item>::value);
+static_assert(!has_member_not_equal<Item>::value);
+static_assert(std::is_same_v<decltype(std::declval<const Item &>() ==
+                                      std::declval<const Item &>()),
+                             bool>);
+static_assert(std::is_same_v<decltype(std::declval<const Item &>() !=
+                                      std::declval<const Item &>()),
+                             bool>);
+static_assert(noexcept(std::declval<const Item &>() ==
+                       std::declval<const Item &>()));
+static_assert(noexcept(std::declval<const Item &>() !=
+                       std::declval<const Item &>()));
+static_assert(noexcept(operator==(std::declval<const Item &>(),
+                                  std::declval<const Item &>())));
+static_assert(noexcept(operator!=(std::declval<const Item &>(),
+                                  std::declval<const Item &>())));
+
+using LegacyLocalIntPtr = sycl::local_ptr<int>;
+using LegacyGlobalIntPtr = sycl::global_ptr<int>;
+using DecoratedLocalIntPtr = sycl::decorated_local_ptr<int>;
+using DecoratedGlobalIntPtr = sycl::decorated_global_ptr<int>;
+using DecoratedLocalConstIntPtr = sycl::decorated_local_ptr<const int>;
+using DecoratedGlobalConstIntPtr = sycl::decorated_global_ptr<const int>;
+
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyLocalIntPtr>(), std::declval<LegacyGlobalIntPtr>(),
+    std::size_t{}, std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyGlobalIntPtr>(), std::declval<LegacyLocalIntPtr>(),
+    std::size_t{}, std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedLocalIntPtr>(),
+    std::declval<DecoratedGlobalConstIntPtr>(), std::size_t{}, std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedGlobalIntPtr>(),
+    std::declval<DecoratedLocalConstIntPtr>(), std::size_t{}, std::size_t{})));
+
+using BoolVector = sycl::vec<bool, 2>;
+using LegacyLocalBoolPtr = sycl::local_ptr<bool>;
+using LegacyGlobalBoolPtr = sycl::global_ptr<bool>;
+using LegacyLocalBoolVectorPtr = sycl::local_ptr<BoolVector>;
+using LegacyGlobalBoolVectorPtr = sycl::global_ptr<BoolVector>;
+using DecoratedLocalBoolPtr = sycl::decorated_local_ptr<bool>;
+using DecoratedGlobalConstBoolPtr = sycl::decorated_global_ptr<const bool>;
+using DecoratedLocalBoolVectorPtr = sycl::decorated_local_ptr<BoolVector>;
+using DecoratedGlobalConstBoolVectorPtr =
+    sycl::decorated_global_ptr<const BoolVector>;
+
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyLocalBoolPtr>(), std::declval<LegacyGlobalBoolPtr>(),
+    std::size_t{}, std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyLocalBoolVectorPtr>(),
+    std::declval<LegacyGlobalBoolVectorPtr>(), std::size_t{}, std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedLocalBoolPtr>(),
+    std::declval<DecoratedGlobalConstBoolPtr>(), std::size_t{},
+    std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedLocalBoolVectorPtr>(),
+    std::declval<DecoratedGlobalConstBoolVectorPtr>(), std::size_t{},
+    std::size_t{})));
+
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyLocalIntPtr>(), std::declval<LegacyGlobalIntPtr>(),
+    std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<LegacyGlobalIntPtr>(), std::declval<LegacyLocalIntPtr>(),
+    std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedLocalIntPtr>(),
+    std::declval<DecoratedGlobalConstIntPtr>(), std::size_t{})));
+static_assert(noexcept(std::declval<const Item &>().async_work_group_copy(
+    std::declval<DecoratedGlobalIntPtr>(),
+    std::declval<DecoratedLocalConstIntPtr>(), std::size_t{})));
+
+static_assert(noexcept(
+    std::declval<const Item &>().wait_for(std::declval<sycl::device_event>())));