diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp
index 97ea1d7b3e864..400fccf3fb7d3 100644
--- a/sycl/include/sycl/multi_ptr.hpp
+++ b/sycl/include/sycl/multi_ptr.hpp
@@ -825,7 +825,11 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
       : multi_ptr(Accessor.get_pointer()) {}
 
   // Only if Space == local_space || generic_space
-  template <int dimensions>
+  template <
+      int dimensions, access::address_space _Space = Space,
+      typename = typename std::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
   multi_ptr(local_accessor<ElementType, dimensions> Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
diff --git a/sycl/test-e2e/multi_ptr/prefetch.cpp b/sycl/test-e2e/multi_ptr/prefetch.cpp
new file mode 100644
index 0000000000000..221e10ca77a56
--- /dev/null
+++ b/sycl/test-e2e/multi_ptr/prefetch.cpp
@@ -0,0 +1,271 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+//===------------------------- prefetch.cpp
+//--------------------------------===//
+//
+// 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 <cassert>
+#include <cmath>
+#include <iostream>
+#include <sycl/multi_ptr.hpp>
+#include <sycl/queue.hpp>
+#include <vector>
+
+using namespace sycl;
+
+template <access::decorated IsDecorated> class PrefetchKernel;
+
+template <access::decorated IsDecorated> void testPrefetchWithDecoration() {
+  constexpr size_t Size = 1024;
+  std::vector<int> HostData(Size);
+  for (size_t i = 0; i < Size; ++i) {
+    HostData[i] = static_cast<int>(i);
+  }
+
+  queue Q;
+  buffer<int, 1> Buf(HostData.data(), range<1>(Size));
+
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access<access::mode::read_write>(CGH);
+
+    CGH.parallel_for<PrefetchKernel<IsDecorated>>(
+        range<1>(Size), [=](id<1> Index) {
+          auto Ptr = Acc.template get_multi_ptr<IsDecorated>();
+
+          // Test prefetch with different element counts
+          if (Index[0] == 0) {
+            Ptr.prefetch(1);
+            Ptr.prefetch(16);
+            Ptr.prefetch(64);
+            Ptr.prefetch(256);
+          }
+
+          // Test prefetch at different offsets
+          if (Index[0] < Size - 100) {
+            auto OffsetPtr = Ptr + Index[0];
+            OffsetPtr.prefetch(10);
+          }
+
+          // Actual computation to ensure prefetch is useful
+          int Sum = 0;
+          for (size_t i = 0; i < 10 && Index[0] + i < Size; ++i) {
+            Sum += Ptr[Index[0] + i];
+          }
+          Acc[Index] = Sum;
+        });
+  });
+
+  Q.wait();
+
+  // Verify results
+  auto HostAcc = Buf.get_host_access();
+  for (size_t i = 0; i < Size; ++i) {
+    int Expected = 0;
+    for (size_t j = 0; j < 10 && i + j < Size; ++j) {
+      Expected += static_cast<int>(i + j);
+    }
+    assert(HostAcc[i] == Expected && "Prefetch test failed");
+  }
+}
+
+void testPrefetchWithGlobalPointer() {
+  constexpr size_t Size = 512;
+  std::vector<float> HostData(Size);
+  for (size_t i = 0; i < Size; ++i) {
+    HostData[i] = static_cast<float>(i) * 0.5f;
+  }
+
+  queue Q;
+  buffer<float, 1> Buf(HostData.data(), range<1>(Size));
+
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access<access::mode::read_write>(CGH);
+
+    CGH.parallel_for<class GlobalPrefetchKernel>(
+        range<1>(Size), [=](id<1> Index) {
+          using global_ptr =
+              multi_ptr<float, access::address_space::global_space,
+                        access::decorated::yes>;
+
+          global_ptr Ptr =
+              address_space_cast<access::address_space::global_space,
+                                 access::decorated::yes>(&Acc[0]);
+
+          // Prefetch future data
+          if (Index[0] < Size - 50) {
+            auto FuturePtr = Ptr + Index[0] + 10;
+            FuturePtr.prefetch(20);
+          }
+
+          // Process data
+          float Sum = 0.0f;
+          for (size_t i = 0; i < 5 && Index[0] + i < Size; ++i) {
+            Sum += Ptr[Index[0] + i];
+          }
+          Acc[Index] = Sum;
+        });
+  });
+
+  Q.wait();
+
+  // Verify results
+  auto HostAcc = Buf.get_host_access();
+  for (size_t i = 0; i < Size; ++i) {
+    float Expected = 0.0f;
+    for (size_t j = 0; j < 5 && i + j < Size; ++j) {
+      Expected += static_cast<float>(i + j) * 0.5f;
+    }
+    assert(std::abs(HostAcc[i] - Expected) < 0.001f &&
+           "Global prefetch test failed");
+  }
+}
+
+void testPrefetchWithLargeData() {
+  constexpr size_t Size = 4096;
+  std::vector<double> HostData(Size);
+  for (size_t i = 0; i < Size; ++i) {
+    HostData[i] = static_cast<double>(i);
+  }
+
+  queue Q;
+  buffer<double, 1> Buf(HostData.data(), range<1>(Size));
+
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access<access::mode::read>(CGH);
+
+    CGH.parallel_for<class LargePrefetchKernel>(
+        range<1>(Size / 8), [=](id<1> Index) {
+          auto Ptr = Acc.template get_multi_ptr<access::decorated::yes>();
+          size_t BaseIndex = Index[0] * 8;
+
+          // Prefetch a chunk of data
+          auto ChunkPtr = Ptr + BaseIndex;
+          ChunkPtr.prefetch(64);
+
+          // Process the prefetched chunk
+          double Sum = 0.0;
+          for (size_t i = 0; i < 8; ++i) {
+            Sum += ChunkPtr[i];
+          }
+        });
+  });
+
+  Q.wait();
+
+  auto HostAcc = Buf.get_host_access();
+  for (size_t i = 0; i < Size; ++i) {
+    assert(HostAcc[i] == static_cast<double>(i) &&
+           "Large prefetch test failed");
+  }
+}
+
+void testPrefetchAtBoundaries() {
+  constexpr size_t Size = 256;
+  std::vector<int> HostData(Size, 42);
+
+  queue Q;
+  buffer<int, 1> Buf(HostData.data(), range<1>(Size));
+
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access<access::mode::read>(CGH);
+
+    CGH.single_task<class BoundaryPrefetchKernel>([=]() {
+      auto Ptr = Acc.template get_multi_ptr<access::decorated::no>();
+
+      // Prefetch at start
+      Ptr.prefetch(1);
+      Ptr.prefetch(10);
+
+      // Prefetch at end (careful not to go beyond buffer)
+      auto EndPtr = Ptr + (Size - 10);
+      EndPtr.prefetch(10);
+
+      // Prefetch zero elements (edge case)
+      Ptr.prefetch(0);
+    });
+  });
+
+  Q.wait();
+
+  auto HostAcc = Buf.get_host_access();
+  for (size_t i = 0; i < Size; ++i) {
+    assert(HostAcc[i] == 42 && "Boundary prefetch test failed");
+  }
+}
+
+void testPrefetchWithStructs() {
+  struct TestStruct {
+    int A;
+    float B;
+    double C;
+  };
+
+  constexpr size_t Size = 128;
+  std::vector<TestStruct> HostData(Size);
+  for (size_t i = 0; i < Size; ++i) {
+    HostData[i] = {static_cast<int>(i), static_cast<float>(i) * 1.5f,
+                   static_cast<double>(i) * 2.5};
+  }
+
+  queue Q;
+  buffer<TestStruct, 1> Buf(HostData.data(), range<1>(Size));
+
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.get_access<access::mode::read>(CGH);
+
+    CGH.parallel_for<class StructPrefetchKernel>(
+        range<1>(Size / 2), [=](id<1> Index) {
+          auto Ptr = Acc.template get_multi_ptr<access::decorated::yes>();
+          size_t BaseIdx = Index[0] * 2;
+
+          // Prefetch structures
+          auto StructPtr = Ptr + BaseIdx;
+          StructPtr.prefetch(8);
+
+          // Access the data
+          int Sum = 0;
+          for (size_t i = 0; i < 2 && BaseIdx + i < Size; ++i) {
+            Sum += StructPtr[i].A;
+          }
+        });
+  });
+
+  Q.wait();
+
+  auto HostAcc = Buf.get_host_access();
+  for (size_t i = 0; i < Size; ++i) {
+    assert(HostAcc[i].A == static_cast<int>(i) &&
+           "Struct prefetch integer field test failed");
+    assert(std::abs(HostAcc[i].B - static_cast<float>(i) * 1.5f) < 0.001f &&
+           "Struct prefetch float field test failed");
+    assert(std::abs(HostAcc[i].C - static_cast<double>(i) * 2.5) < 0.001 &&
+           "Struct prefetch double field test failed");
+  }
+}
+
+int main() {
+  // Test prefetch with decorated pointers
+  testPrefetchWithDecoration<access::decorated::yes>();
+  testPrefetchWithDecoration<access::decorated::no>();
+
+  // Test prefetch with explicit global pointers
+  testPrefetchWithGlobalPointer();
+
+  // Test prefetch with large datasets
+  testPrefetchWithLargeData();
+
+  // Test prefetch at buffer boundaries
+  testPrefetchAtBoundaries();
+
+  // Test prefetch with complex data structures
+  testPrefetchWithStructs();
+
+  std::cout << "All prefetch tests passed!" << std::endl;
+  return 0;
+}
diff --git a/sycl/test/multi_ptr/restricted_conversions.cpp b/sycl/test/multi_ptr/restricted_conversions.cpp
new file mode 100644
index 0000000000000..ee1e7391736c4
--- /dev/null
+++ b/sycl/test/multi_ptr/restricted_conversions.cpp
@@ -0,0 +1,87 @@
+#include <sycl/detail/core.hpp>
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note
+
+template <typename T, sycl::access::address_space Space,
+          sycl::access::decorated Decorated>
+using multi_ptr_t = sycl::multi_ptr<T, Space, Decorated>;
+
+using legacy_ptr = multi_ptr_t<int, sycl::access::address_space::private_space,
+                               sycl::access::decorated::legacy>;
+using non_legacy_ptr =
+    multi_ptr_t<int, sycl::access::address_space::private_space,
+                sycl::access::decorated::no>;
+using private_ptr = multi_ptr_t<int, sycl::access::address_space::private_space,
+                                sycl::access::decorated::no>;
+using const_void_ptr =
+    multi_ptr_t<const void, sycl::access::address_space::private_space,
+                sycl::access::decorated::no>;
+using void_ptr = multi_ptr_t<void, sycl::access::address_space::private_space,
+                             sycl::access::decorated::no>;
+using global_ptr = multi_ptr_t<int, sycl::access::address_space::global_space,
+                               sycl::access::decorated::no>;
+using local_ptr = multi_ptr_t<int, sycl::access::address_space::local_space,
+                              sycl::access::decorated::no>;
+// expected-warning@+2 {{constant_space' is deprecated: sycl::access::address_space::constant_space is deprecated since SYCL 2020}}
+using constant_ptr =
+    multi_ptr_t<const int, sycl::access::address_space::constant_space,
+                sycl::access::decorated::no>;
+using generic_ptr =
+    multi_ptr_t<const int, sycl::access::address_space::generic_space,
+                sycl::access::decorated::no>;
+
+legacy_ptr leg_ptr{nullptr};
+non_legacy_ptr nonleg_ptr{nullptr};
+
+// expected-error@+1 {{no matching constructor for initialization of}}
+non_legacy_ptr nonleg_ptr1{leg_ptr};
+// expected-error@+1 {{no viable conversion from}}
+non_legacy_ptr nonleg_ptr2 = leg_ptr;
+// TODO: is constructor legal?
+// expected-warning@+1 {{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead}}
+legacy_ptr leg_ptr1{nonleg_ptr};
+// expected-error@+1 {{no viable conversion from 'multi_ptr_t<}}
+legacy_ptr leg_ptr2 = nonleg_ptr;
+
+const_void_ptr const_void{nullptr};
+
+// expected-error@+1 {{no matching constructor for initialization of 'void_ptr'}}
+void_ptr void_ptr1{const_void};
+// expected-error@+1 {{no viable conversion from 'multi_ptr_t<}}
+void_ptr void_ptr2 = const_void;
+
+global_ptr global{nullptr};
+
+private_ptr private_ptr_instance{nullptr};
+
+// expected-error@+1 {{no matching constructor for initialization of 'local_ptr'}}
+local_ptr local{global};
+
+// expected-error@+1 {{no viable conversion from 'multi_ptr_t<}}
+local_ptr local1 = global;
+
+// expected-error@+1 {{no matching constructor for initialization of 'global_ptr'}}
+global_ptr global1{local};
+
+// expected-error@+1 {{no viable conversion from 'multi_ptr_t<}}
+global_ptr global2 = local;
+
+// expected-error@+1 {{no matching constructor for initialization of 'local_ptr'}}
+local_ptr local_from_private{private_ptr_instance};
+
+// expected-error@+1 {{no viable conversion from 'multi_ptr_t<}}
+local_ptr local_from_private_implicit = private_ptr_instance;
+
+// expected-error@+1 {{no matching constructor for initialization of 'global_ptr'}}
+global_ptr global_from_private{private_ptr_instance};
+
+// expected-error@+1 {{no matching constructor for initialization of 'private_ptr'}}
+private_ptr private_from_local{local};
+
+// expected-error@+1 {{no matching constructor for initialization of 'private_ptr'}}
+private_ptr private_from_global{global};
+
+// expected-warning@+1 2 {{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+bool private_equals_local = private_ptr_instance == local;
+
+// expected-warning@+1 2 {{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+bool private_less_than_local = private_ptr_instance < local;
diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt
index 9a322ae24fbaa..75c1b86aa078a 100644
--- a/sycl/unittests/CMakeLists.txt
+++ b/sycl/unittests/CMakeLists.txt
@@ -51,6 +51,7 @@ add_subdirectory(accessor)
 add_subdirectory(handler)
 add_subdirectory(builtins)
 add_subdirectory(buffer/l0_specific)
+add_subdirectory(multi_ptr)
 
 # Enable compression unit-tests only if zstd is present.
 if (LLVM_ENABLE_ZSTD)
diff --git a/sycl/unittests/multi_ptr/Accessors.cpp b/sycl/unittests/multi_ptr/Accessors.cpp
new file mode 100644
index 0000000000000..38b38681c58ae
--- /dev/null
+++ b/sycl/unittests/multi_ptr/Accessors.cpp
@@ -0,0 +1,127 @@
+#include <sycl/accessor.hpp>
+#include <sycl/multi_ptr.hpp>
+
+#include <gtest/gtest.h>
+
+#include <type_traits>
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+using multi_ptr_t = sycl::multi_ptr<ElementType, Space, IsDecorated>;
+
+TEST(MultiPtrAccessors, AccessorDevice) {
+  using rw_acc = sycl::accessor<int, 1, sycl::access_mode::read_write,
+                                sycl::target::device>;
+  using atomic_acc =
+      sycl::accessor<int, 1, sycl::access_mode::atomic, sycl::target::device>;
+  using ro_acc =
+      sycl::accessor<int, 1, sycl::access_mode::read, sycl::target::device>;
+  using global_multi_ptr =
+      sycl::multi_ptr<int, sycl::access::address_space::global_space>;
+  using generic_multi_ptr =
+      sycl::multi_ptr<int, sycl::access::address_space::generic_space>;
+  using const_multi_ptr =
+      sycl::multi_ptr<const int, sycl::access::address_space::generic_space>;
+
+  using float_global_ptr =
+      sycl::multi_ptr<float, sycl::access::address_space::global_space>;
+
+  EXPECT_TRUE((std::is_constructible_v<global_multi_ptr, rw_acc>));
+  EXPECT_TRUE((std::is_constructible_v<generic_multi_ptr, rw_acc>));
+  EXPECT_TRUE((std::is_constructible_v<global_multi_ptr, atomic_acc>));
+  EXPECT_TRUE((std::is_constructible_v<global_multi_ptr, ro_acc>));
+  EXPECT_TRUE((std::is_constructible_v<generic_multi_ptr, ro_acc>));
+  EXPECT_FALSE((std::is_constructible_v<float_global_ptr, rw_acc>));
+  EXPECT_FALSE((std::is_constructible_v<float_global_ptr, ro_acc>));
+  EXPECT_TRUE((std::is_constructible_v<const_multi_ptr, rw_acc>));
+  EXPECT_TRUE((std::is_constructible_v<const_multi_ptr, ro_acc>));
+}
+
+TEST(MultiPtrAccessors, LocalAccessor) {
+  using local_acc = sycl::local_accessor<int, 1>;
+  using local_multi_ptr =
+      sycl::multi_ptr<int, sycl::access::address_space::local_space>;
+  using generic_multi_ptr =
+      sycl::multi_ptr<int, sycl::access::address_space::generic_space>;
+  using const_multi_ptr =
+      sycl::multi_ptr<const int, sycl::access::address_space::generic_space>;
+  using float_local_ptr =
+      sycl::multi_ptr<float, sycl::access::address_space::local_space>;
+
+  EXPECT_TRUE((std::is_constructible_v<local_multi_ptr, local_acc>));
+  EXPECT_TRUE((std::is_constructible_v<generic_multi_ptr, local_acc>));
+  EXPECT_FALSE((std::is_constructible_v<
+                sycl::multi_ptr<int, sycl::access::address_space::global_space>,
+                local_acc>));
+  EXPECT_FALSE((std::is_constructible_v<float_local_ptr, local_acc>));
+  EXPECT_TRUE((std::is_constructible_v<const_multi_ptr, local_acc>));
+}
+
+TEST(MultiPtrAccessors, AddrSpaceCast) {
+  int Value = 42;
+  const int ConstValue = 17;
+
+  auto PrivateUndecorated =
+      sycl::address_space_cast<address_space::private_space, decorated::no>(
+          &Value);
+  auto GlobalDecorated =
+      sycl::address_space_cast<address_space::global_space, decorated::yes>(
+          &Value);
+  auto LocalUndecorated =
+      sycl::address_space_cast<address_space::local_space, decorated::no>(
+          &Value);
+  auto GenericDecorated =
+      sycl::address_space_cast<address_space::generic_space, decorated::yes>(
+          &Value);
+  auto ConstPrivateDecorated =
+      sycl::address_space_cast<address_space::private_space, decorated::yes>(
+          &ConstValue);
+
+  EXPECT_TRUE((std::is_same_v<
+               decltype(PrivateUndecorated),
+               multi_ptr_t<int, address_space::private_space, decorated::no>>));
+  EXPECT_TRUE((std::is_same_v<
+               decltype(GlobalDecorated),
+               multi_ptr_t<int, address_space::global_space, decorated::yes>>));
+  EXPECT_TRUE((std::is_same_v<
+               decltype(LocalUndecorated),
+               multi_ptr_t<int, address_space::local_space, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_same_v<
+          decltype(GenericDecorated),
+          multi_ptr_t<int, address_space::generic_space, decorated::yes>>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(ConstPrivateDecorated),
+                      multi_ptr_t<const int, address_space::private_space,
+                                  decorated::yes>>));
+
+  EXPECT_EQ(PrivateUndecorated.get_raw(), &Value);
+  EXPECT_EQ(GlobalDecorated.get_raw(), &Value);
+  EXPECT_EQ(LocalUndecorated.get_raw(), &Value);
+  EXPECT_EQ(GenericDecorated.get_raw(), &Value);
+  EXPECT_EQ(ConstPrivateDecorated.get_raw(), &ConstValue);
+
+  EXPECT_FALSE((std::is_same_v<
+                decltype(PrivateUndecorated),
+                multi_ptr_t<int, address_space::global_space, decorated::no>>));
+  EXPECT_FALSE(
+      (std::is_same_v<
+          decltype(PrivateUndecorated),
+          multi_ptr_t<int, address_space::private_space, decorated::yes>>));
+  EXPECT_FALSE(
+      (std::is_same_v<
+          decltype(ConstPrivateDecorated),
+          multi_ptr_t<int, address_space::private_space, decorated::yes>>));
+
+  auto NullPrivate =
+      sycl::address_space_cast<address_space::private_space, decorated::no>(
+          static_cast<int *>(nullptr));
+  auto NullGlobalConst =
+      sycl::address_space_cast<address_space::global_space, decorated::yes>(
+          static_cast<const int *>(nullptr));
+
+  EXPECT_EQ(NullPrivate, nullptr);
+  EXPECT_EQ(NullGlobalConst, nullptr);
+}
diff --git a/sycl/unittests/multi_ptr/CMakeLists.txt b/sycl/unittests/multi_ptr/CMakeLists.txt
new file mode 100644
index 0000000000000..c2f2c88c500d9
--- /dev/null
+++ b/sycl/unittests/multi_ptr/CMakeLists.txt
@@ -0,0 +1,9 @@
+add_sycl_unittest(MultiPtrTests OBJECT
+  Constructors.cpp
+  Operators.cpp
+  Conversion.cpp
+  Accessors.cpp
+  NegativeConversions.cpp
+  VoidSpecialization.cpp
+)
+
diff --git a/sycl/unittests/multi_ptr/Constructors.cpp b/sycl/unittests/multi_ptr/Constructors.cpp
new file mode 100644
index 0000000000000..f36a3064c207b
--- /dev/null
+++ b/sycl/unittests/multi_ptr/Constructors.cpp
@@ -0,0 +1,194 @@
+//===---------------------------- Constructors.cpp ------------===//
+//
+// 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 <gtest/gtest.h>
+
+#include <type_traits>
+#include <utility>
+
+namespace {
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+template <typename ElementType, decorated IsDecorated>
+using private_multi_ptr =
+    sycl::multi_ptr<ElementType, address_space::private_space, IsDecorated>;
+
+template <typename ElementType, decorated IsDecorated>
+using local_multi_ptr =
+    sycl::multi_ptr<ElementType, address_space::local_space, IsDecorated>;
+
+template <typename ElementType, decorated IsDecorated>
+private_multi_ptr<ElementType, IsDecorated>
+makePrivatePtr(std::add_pointer_t<ElementType> Ptr) {
+  return sycl::address_space_cast<address_space::private_space, IsDecorated>(
+      Ptr);
+}
+
+template <decorated IsDecorated> void checkCopyConstruction() {
+  int Value = 17;
+  auto Original = makePrivatePtr<int, IsDecorated>(&Value);
+  private_multi_ptr<int, IsDecorated> Copy{Original};
+
+  EXPECT_EQ(Copy, Original);
+  EXPECT_NE(Copy, nullptr);
+}
+
+template <decorated IsDecorated> void checkMoveConstruction() {
+  int Value = 23;
+  auto Source = makePrivatePtr<int, IsDecorated>(&Value);
+  auto Expected = makePrivatePtr<int, IsDecorated>(&Value);
+  private_multi_ptr<int, IsDecorated> Moved{std::move(Source)};
+
+  EXPECT_EQ(Moved, Expected);
+  EXPECT_NE(Moved, nullptr);
+}
+
+TEST(MultiPtrConstructors, ValidConstructorTraitsAreSatisfied) {
+  EXPECT_TRUE(
+      (std::is_default_constructible_v<private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE((
+      std::is_default_constructible_v<private_multi_ptr<int, decorated::yes>>));
+  EXPECT_TRUE((std::is_constructible_v<private_multi_ptr<int, decorated::no>,
+                                       std::nullptr_t>));
+  EXPECT_TRUE((std::is_constructible_v<private_multi_ptr<int, decorated::yes>,
+                                       std::nullptr_t>));
+  EXPECT_TRUE(
+      (std::is_copy_constructible_v<private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_copy_constructible_v<private_multi_ptr<int, decorated::yes>>));
+  EXPECT_TRUE(
+      (std::is_move_constructible_v<private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_move_constructible_v<private_multi_ptr<int, decorated::yes>>));
+  EXPECT_TRUE(
+      (std::is_constructible_v<private_multi_ptr<const int, decorated::no>,
+                               private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_constructible_v<private_multi_ptr<const int, decorated::yes>,
+                               private_multi_ptr<int, decorated::yes>>));
+  EXPECT_TRUE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::no>,
+                               private_multi_ptr<int, decorated::yes>>));
+  EXPECT_TRUE((std::is_constructible_v<private_multi_ptr<int, decorated::yes>,
+                                       private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_constructible_v<private_multi_ptr<const int, decorated::yes>,
+                               private_multi_ptr<int, decorated::no>>));
+  EXPECT_TRUE(
+      (std::is_constructible_v<private_multi_ptr<const int, decorated::no>,
+                               private_multi_ptr<int, decorated::yes>>));
+}
+
+TEST(MultiPtrConstructors, InvalidConstructorTraitsAreRejected) {
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::no>,
+                               private_multi_ptr<const int, decorated::no>>));
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::yes>,
+                               private_multi_ptr<const int, decorated::yes>>));
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::no>,
+                               private_multi_ptr<float, decorated::no>>));
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::yes>,
+                               private_multi_ptr<float, decorated::yes>>));
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::no>,
+                               local_multi_ptr<float, decorated::no>>));
+  EXPECT_FALSE(
+      (std::is_constructible_v<private_multi_ptr<int, decorated::yes>,
+                               local_multi_ptr<float, decorated::yes>>));
+}
+
+TEST(MultiPtrConstructors, DefaultConstructedPointersAreNull) {
+  private_multi_ptr<int, decorated::no> Undecorated;
+  private_multi_ptr<int, decorated::yes> Decorated;
+
+  EXPECT_EQ(Undecorated, nullptr);
+  EXPECT_EQ(Decorated, nullptr);
+}
+
+TEST(MultiPtrConstructors, NullptrConstructedPointersAreNull) {
+  private_multi_ptr<int, decorated::no> Undecorated{nullptr};
+  private_multi_ptr<int, decorated::yes> Decorated{nullptr};
+
+  EXPECT_EQ(Undecorated, nullptr);
+  EXPECT_EQ(Decorated, nullptr);
+}
+
+TEST(MultiPtrConstructors, CopyConstructionPreservesPointerValue) {
+  checkCopyConstruction<decorated::no>();
+  checkCopyConstruction<decorated::yes>();
+}
+
+TEST(MultiPtrConstructors, MoveConstructionPreservesPointerValue) {
+  checkMoveConstruction<decorated::no>();
+  checkMoveConstruction<decorated::yes>();
+}
+
+TEST(MultiPtrConstructors, ConvertingConstructionAddsConstQualification) {
+  int Value = 31;
+
+  private_multi_ptr<int, decorated::no> MutableUndecorated =
+      makePrivatePtr<int, decorated::no>(&Value);
+  private_multi_ptr<int, decorated::yes> MutableDecorated =
+      makePrivatePtr<int, decorated::yes>(&Value);
+
+  private_multi_ptr<const int, decorated::no> ConstUndecorated{
+      MutableUndecorated};
+  private_multi_ptr<const int, decorated::yes> ConstDecorated{MutableDecorated};
+  auto ExpectedUndecorated = makePrivatePtr<const int, decorated::no>(&Value);
+  auto ExpectedDecorated = makePrivatePtr<const int, decorated::yes>(&Value);
+
+  EXPECT_EQ(ConstUndecorated, ExpectedUndecorated);
+  EXPECT_EQ(ConstDecorated, ExpectedDecorated);
+}
+
+TEST(MultiPtrConstructors,
+     ConvertingConstructionBetweenDecorationModesPreservesPointerValue) {
+  int Value = 47;
+
+  private_multi_ptr<int, decorated::yes> DecoratedPtr =
+      makePrivatePtr<int, decorated::yes>(&Value);
+  private_multi_ptr<int, decorated::no> UndecoratedPtr =
+      makePrivatePtr<int, decorated::no>(&Value);
+
+  private_multi_ptr<int, decorated::no> FromDecorated{DecoratedPtr};
+  private_multi_ptr<int, decorated::yes> FromUndecorated{UndecoratedPtr};
+  auto ExpectedUndecorated = makePrivatePtr<int, decorated::no>(&Value);
+  auto ExpectedDecorated = makePrivatePtr<int, decorated::yes>(&Value);
+
+  EXPECT_EQ(FromDecorated, ExpectedUndecorated);
+  EXPECT_EQ(FromUndecorated, ExpectedDecorated);
+}
+
+TEST(MultiPtrConstructors,
+     ConvertingConstructionCanAddConstAndChangeDecorationTogether) {
+  int Value = 59;
+
+  private_multi_ptr<int, decorated::no> MutableUndecorated =
+      makePrivatePtr<int, decorated::no>(&Value);
+  private_multi_ptr<int, decorated::yes> MutableDecorated =
+      makePrivatePtr<int, decorated::yes>(&Value);
+
+  private_multi_ptr<const int, decorated::yes> ConstDecorated{
+      MutableUndecorated};
+  private_multi_ptr<const int, decorated::no> ConstUndecorated{
+      MutableDecorated};
+  auto ExpectedDecorated = makePrivatePtr<const int, decorated::yes>(&Value);
+  auto ExpectedUndecorated = makePrivatePtr<const int, decorated::no>(&Value);
+
+  EXPECT_EQ(ConstDecorated, ExpectedDecorated);
+  EXPECT_EQ(ConstUndecorated, ExpectedUndecorated);
+}
+
+} // namespace
diff --git a/sycl/unittests/multi_ptr/Conversion.cpp b/sycl/unittests/multi_ptr/Conversion.cpp
new file mode 100644
index 0000000000000..a59082ac5801a
--- /dev/null
+++ b/sycl/unittests/multi_ptr/Conversion.cpp
@@ -0,0 +1,434 @@
+#include <sycl/multi_ptr.hpp>
+
+#include <gtest/gtest.h>
+
+#include <memory>
+#include <type_traits>
+#include <utility>
+
+namespace {
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+struct Record {
+  int Value;
+  int Tag;
+};
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+using multi_ptr_t = sycl::multi_ptr<ElementType, Space, IsDecorated>;
+
+template <address_space Space>
+using access_element_t =
+    std::conditional_t<Space == address_space::constant_space, const Record,
+                       Record>;
+
+template <address_space Space>
+using access_storage_t = std::remove_const_t<access_element_t<Space>>;
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+multi_ptr_t<ElementType, Space, IsDecorated>
+makePtr(std::add_pointer_t<ElementType> Ptr) {
+  if constexpr (IsDecorated == decorated::legacy) {
+    return multi_ptr_t<ElementType, Space, IsDecorated>{Ptr};
+  } else {
+    return sycl::address_space_cast<Space, IsDecorated>(Ptr);
+  }
+}
+
+template <decorated IsDecorated> constexpr decorated toggledDecoration() {
+  static_assert(IsDecorated != decorated::legacy);
+  return IsDecorated == decorated::no ? decorated::yes : decorated::no;
+}
+
+template <address_space Space, decorated IsDecorated>
+void checkPointerAccessAndUnderlyingPointers() {
+  using element_t = access_element_t<Space>;
+  using ptr_t = multi_ptr_t<element_t, Space, IsDecorated>;
+
+  access_storage_t<Space> Values[] = {{11, 111}, {22, 222}};
+  auto *RawPtr = static_cast<std::add_pointer_t<element_t>>(&Values[0]);
+  ptr_t Ptr = makePtr<element_t, Space, IsDecorated>(RawPtr);
+
+  EXPECT_EQ(Ptr->Value, 11);
+  EXPECT_EQ(Ptr->Tag, 111);
+  EXPECT_EQ(Ptr[1].Value, 22);
+  EXPECT_EQ(Ptr[1].Tag, 222);
+
+  decltype(Ptr.get()) PointerValue = Ptr;
+  EXPECT_EQ(PointerValue, Ptr.get());
+  EXPECT_EQ(Ptr.get_raw(), RawPtr);
+
+  if constexpr (IsDecorated == decorated::legacy) {
+    EXPECT_EQ(Ptr.get_decorated(), Ptr.get());
+  } else {
+    auto DecoratedPtr = makePtr<element_t, Space, decorated::yes>(RawPtr);
+    EXPECT_EQ(Ptr.get_decorated(), DecoratedPtr.get_decorated());
+  }
+}
+
+template <decorated SourceDecoration, decorated TargetDecoration>
+void checkGenericExplicitMutableConversions() {
+  using source_t =
+      multi_ptr_t<Record, address_space::generic_space, SourceDecoration>;
+  using private_t =
+      multi_ptr_t<Record, address_space::private_space, TargetDecoration>;
+  using global_t =
+      multi_ptr_t<Record, address_space::global_space, TargetDecoration>;
+  using local_t =
+      multi_ptr_t<Record, address_space::local_space, TargetDecoration>;
+  using const_private_t =
+      multi_ptr_t<const Record, address_space::private_space, TargetDecoration>;
+  using const_global_t =
+      multi_ptr_t<const Record, address_space::global_space, TargetDecoration>;
+  using const_local_t =
+      multi_ptr_t<const Record, address_space::local_space, TargetDecoration>;
+
+  Record Value{31, 131};
+  source_t Source =
+      makePtr<Record, address_space::generic_space, SourceDecoration>(&Value);
+
+  EXPECT_TRUE((std::is_constructible_v<private_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<global_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<local_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<const_private_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<const_global_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<const_local_t, source_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, private_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, global_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, local_t>));
+
+  private_t Private{Source};
+  global_t Global{Source};
+  local_t Local{Source};
+  const_private_t ConstPrivate{Source};
+  const_global_t ConstGlobal{Source};
+  const_local_t ConstLocal{Source};
+
+  EXPECT_EQ(Private.get_raw(), &Value);
+  EXPECT_EQ(Global.get_raw(), &Value);
+  EXPECT_EQ(Local.get_raw(), &Value);
+  EXPECT_EQ(ConstPrivate.get_raw(), &Value);
+  EXPECT_EQ(ConstGlobal.get_raw(), &Value);
+  EXPECT_EQ(ConstLocal.get_raw(), &Value);
+}
+
+template <decorated SourceDecoration, decorated TargetDecoration>
+void checkGenericExplicitConstConversions() {
+  using source_t =
+      multi_ptr_t<const Record, address_space::generic_space, SourceDecoration>;
+  using private_t =
+      multi_ptr_t<const Record, address_space::private_space, TargetDecoration>;
+  using global_t =
+      multi_ptr_t<const Record, address_space::global_space, TargetDecoration>;
+  using local_t =
+      multi_ptr_t<const Record, address_space::local_space, TargetDecoration>;
+
+  const Record Value{32, 132};
+  source_t Source =
+      makePtr<const Record, address_space::generic_space, SourceDecoration>(
+          &Value);
+
+  EXPECT_TRUE((std::is_constructible_v<private_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<global_t, source_t>));
+  EXPECT_TRUE((std::is_constructible_v<local_t, source_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, private_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, global_t>));
+  EXPECT_FALSE((std::is_convertible_v<source_t, local_t>));
+
+  private_t Private{Source};
+  global_t Global{Source};
+  local_t Local{Source};
+
+  EXPECT_EQ(Private.get_raw(), &Value);
+  EXPECT_EQ(Global.get_raw(), &Value);
+  EXPECT_EQ(Local.get_raw(), &Value);
+}
+
+template <address_space Space, decorated SourceDecoration>
+void checkNonLegacyMutableSameSpaceConversions() {
+  static_assert(SourceDecoration != decorated::legacy);
+
+  using source_t = multi_ptr_t<Record, Space, SourceDecoration>;
+  using typed_no_t = multi_ptr_t<Record, Space, decorated::no>;
+  using typed_yes_t = multi_ptr_t<Record, Space, decorated::yes>;
+  using typed_legacy_t = multi_ptr_t<Record, Space, decorated::legacy>;
+  using toggled_t =
+      multi_ptr_t<Record, Space, toggledDecoration<SourceDecoration>()>;
+  using void_same_t = multi_ptr_t<void, Space, SourceDecoration>;
+  using void_toggled_t =
+      multi_ptr_t<void, Space, toggledDecoration<SourceDecoration>()>;
+  using void_no_t = multi_ptr_t<void, Space, decorated::no>;
+  using void_yes_t = multi_ptr_t<void, Space, decorated::yes>;
+  using void_legacy_t = multi_ptr_t<void, Space, decorated::legacy>;
+  using const_no_t = multi_ptr_t<const Record, Space, decorated::no>;
+  using const_yes_t = multi_ptr_t<const Record, Space, decorated::yes>;
+  using const_legacy_t = multi_ptr_t<const Record, Space, decorated::legacy>;
+
+  Record Value{41, 141};
+  source_t Source = makePtr<Record, Space, SourceDecoration>(&Value);
+
+  EXPECT_TRUE((std::is_convertible_v<source_t, toggled_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, void_no_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, void_yes_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, void_legacy_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_no_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_yes_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_legacy_t>));
+  EXPECT_TRUE((std::is_constructible_v<source_t, void_same_t>));
+  EXPECT_FALSE((std::is_convertible_v<void_same_t, source_t>));
+
+  toggled_t Toggled = Source;
+  void_no_t VoidNo = Source;
+  void_yes_t VoidYes = Source;
+  void_legacy_t VoidLegacy = Source;
+  const_no_t ConstNo = Source;
+  const_yes_t ConstYes = Source;
+  const_legacy_t ConstLegacy = Source;
+  void_same_t VoidSame = Source;
+  void_toggled_t VoidToggled = VoidSame;
+  typed_no_t FromVoidNo = static_cast<typed_no_t>(VoidNo);
+  typed_yes_t FromVoidYes = static_cast<typed_yes_t>(VoidYes);
+  typed_legacy_t FromVoidLegacy = static_cast<typed_legacy_t>(VoidLegacy);
+  toggled_t FromVoidToggled = static_cast<toggled_t>(VoidToggled);
+  source_t FromVoid = static_cast<source_t>(VoidSame);
+
+  EXPECT_EQ(Toggled.get_raw(), &Value);
+  EXPECT_EQ(ConstNo.get_raw(), &Value);
+  EXPECT_EQ(ConstYes.get_raw(), &Value);
+  EXPECT_EQ(ConstLegacy.get_raw(), &Value);
+  EXPECT_EQ(FromVoidNo.get_raw(), &Value);
+  EXPECT_EQ(FromVoidYes.get_raw(), &Value);
+  EXPECT_EQ(FromVoidLegacy.get_raw(), &Value);
+  EXPECT_EQ(FromVoidToggled.get_raw(), &Value);
+  EXPECT_EQ(FromVoid.get_raw(), &Value);
+}
+
+template <address_space Space, decorated SourceDecoration>
+void checkNonLegacyConstSameSpaceConversions() {
+  static_assert(SourceDecoration != decorated::legacy);
+
+  using source_t = multi_ptr_t<const Record, Space, SourceDecoration>;
+  using typed_no_t = multi_ptr_t<const Record, Space, decorated::no>;
+  using typed_yes_t = multi_ptr_t<const Record, Space, decorated::yes>;
+  using typed_legacy_t = multi_ptr_t<const Record, Space, decorated::legacy>;
+  using toggled_t =
+      multi_ptr_t<const Record, Space, toggledDecoration<SourceDecoration>()>;
+  using const_void_same_t = multi_ptr_t<const void, Space, SourceDecoration>;
+  using const_void_toggled_t =
+      multi_ptr_t<const void, Space, toggledDecoration<SourceDecoration>()>;
+  using const_void_no_t = multi_ptr_t<const void, Space, decorated::no>;
+  using const_void_yes_t = multi_ptr_t<const void, Space, decorated::yes>;
+  using const_void_legacy_t = multi_ptr_t<const void, Space, decorated::legacy>;
+
+  const Record Value{42, 142};
+  source_t Source = makePtr<const Record, Space, SourceDecoration>(&Value);
+
+  EXPECT_TRUE((std::is_convertible_v<source_t, toggled_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_void_no_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_void_yes_t>));
+  EXPECT_TRUE((std::is_convertible_v<source_t, const_void_legacy_t>));
+  EXPECT_TRUE((std::is_constructible_v<source_t, const_void_same_t>));
+  EXPECT_FALSE((std::is_convertible_v<const_void_same_t, source_t>));
+
+  toggled_t Toggled = Source;
+  const_void_no_t ConstVoidNo = Source;
+  const_void_yes_t ConstVoidYes = Source;
+  const_void_legacy_t ConstVoidLegacy = Source;
+  const_void_same_t ConstVoidSame = Source;
+  const_void_toggled_t ConstVoidToggled = ConstVoidSame;
+  typed_no_t FromConstVoidNo = static_cast<typed_no_t>(ConstVoidNo);
+  typed_yes_t FromConstVoidYes = static_cast<typed_yes_t>(ConstVoidYes);
+  typed_legacy_t FromConstVoidLegacy =
+      static_cast<typed_legacy_t>(ConstVoidLegacy);
+  toggled_t FromConstVoidToggled = static_cast<toggled_t>(ConstVoidToggled);
+  source_t FromConstVoid = static_cast<source_t>(ConstVoidSame);
+
+  EXPECT_EQ(Toggled.get_raw(), &Value);
+  EXPECT_EQ(FromConstVoidNo.get_raw(), &Value);
+  EXPECT_EQ(FromConstVoidYes.get_raw(), &Value);
+  EXPECT_EQ(FromConstVoidLegacy.get_raw(), &Value);
+  EXPECT_EQ(FromConstVoidToggled.get_raw(), &Value);
+  EXPECT_EQ(FromConstVoid.get_raw(), &Value);
+}
+
+template <address_space Space> void checkLegacyConversions() {
+  using mutable_t = multi_ptr_t<Record, Space, decorated::legacy>;
+  using const_t = multi_ptr_t<const Record, Space, decorated::legacy>;
+  using void_t = multi_ptr_t<void, Space, decorated::legacy>;
+  using const_void_t = multi_ptr_t<const void, Space, decorated::legacy>;
+
+  Record MutableValue{51, 151};
+  const Record ConstValue{52, 152};
+
+  mutable_t Mutable = makePtr<Record, Space, decorated::legacy>(&MutableValue);
+  const_t Const = makePtr<const Record, Space, decorated::legacy>(&ConstValue);
+
+  EXPECT_TRUE((std::is_convertible_v<mutable_t, void_t>));
+  EXPECT_TRUE((std::is_convertible_v<mutable_t, const_t>));
+  EXPECT_TRUE((std::is_convertible_v<const_t, const_void_t>));
+  EXPECT_TRUE((std::is_constructible_v<mutable_t, void_t>));
+  EXPECT_TRUE((std::is_constructible_v<const_t, const_void_t>));
+
+  void_t VoidPtr = Mutable;
+  const_t ConstFromMutable = Mutable;
+  const_void_t ConstVoidPtr = Const;
+  const_void_t ConstVoidFromVoid = VoidPtr;
+  mutable_t MutableRoundTrip = static_cast<mutable_t>(VoidPtr);
+  const_t ConstRoundTrip = static_cast<const_t>(ConstVoidPtr);
+
+  EXPECT_EQ(VoidPtr.get_raw(), static_cast<void *>(&MutableValue));
+  EXPECT_EQ(ConstFromMutable.get_raw(), &MutableValue);
+  EXPECT_EQ(ConstVoidPtr.get_raw(), static_cast<const void *>(&ConstValue));
+  EXPECT_EQ(ConstVoidFromVoid.get_raw(),
+            static_cast<const void *>(&MutableValue));
+  EXPECT_EQ(MutableRoundTrip.get_raw(), &MutableValue);
+  EXPECT_EQ(ConstRoundTrip.get_raw(), &ConstValue);
+}
+
+template <decorated IsDecorated> void checkFactoryFunctions() {
+  using ptr_t = multi_ptr_t<int, address_space::private_space, IsDecorated>;
+
+  int Value = 77;
+  auto CastResult =
+      sycl::address_space_cast<address_space::private_space, IsDecorated>(
+          &Value);
+  using cast_result_t = decltype(CastResult);
+  EXPECT_TRUE((std::is_same_v<cast_result_t, ptr_t>));
+  EXPECT_EQ(CastResult.get_raw(), &Value);
+
+  auto NullCast =
+      sycl::address_space_cast<address_space::private_space, IsDecorated>(
+          static_cast<int *>(nullptr));
+  EXPECT_EQ(NullCast, nullptr);
+
+  auto MakeResult =
+      sycl::make_ptr<int, address_space::private_space, IsDecorated>(
+          CastResult.get());
+  using make_result_t = decltype(MakeResult);
+  EXPECT_TRUE((std::is_same_v<make_result_t, ptr_t>));
+  EXPECT_EQ(MakeResult.get_raw(), &Value);
+
+  using underlying_pointer_t = decltype(CastResult.get());
+  underlying_pointer_t NullPointer = nullptr;
+  auto NullMake =
+      sycl::make_ptr<int, address_space::private_space, IsDecorated>(
+          NullPointer);
+  EXPECT_EQ(NullMake, nullptr);
+}
+
+void checkDefaultLegacyMakePtr() {
+  int Value = 88;
+  auto LegacyPtr = sycl::make_ptr<int, address_space::private_space>(&Value);
+  EXPECT_TRUE(
+      (std::is_same_v<
+          decltype(LegacyPtr),
+          multi_ptr_t<int, address_space::private_space, decorated::legacy>>));
+  EXPECT_EQ(LegacyPtr.get_raw(), &Value);
+}
+
+TEST(MultiPtrConversion,
+     PointerAccessPointerConversionAndGetDecoratedCoverAllEnumValues) {
+  checkPointerAccessAndUnderlyingPointers<address_space::global_space,
+                                          decorated::no>();
+  checkPointerAccessAndUnderlyingPointers<address_space::global_space,
+                                          decorated::yes>();
+  checkPointerAccessAndUnderlyingPointers<address_space::global_space,
+                                          decorated::legacy>();
+  checkPointerAccessAndUnderlyingPointers<address_space::local_space,
+                                          decorated::no>();
+  checkPointerAccessAndUnderlyingPointers<address_space::local_space,
+                                          decorated::yes>();
+  checkPointerAccessAndUnderlyingPointers<address_space::local_space,
+                                          decorated::legacy>();
+  checkPointerAccessAndUnderlyingPointers<address_space::constant_space,
+                                          decorated::no>();
+  checkPointerAccessAndUnderlyingPointers<address_space::constant_space,
+                                          decorated::yes>();
+  checkPointerAccessAndUnderlyingPointers<address_space::constant_space,
+                                          decorated::legacy>();
+  checkPointerAccessAndUnderlyingPointers<address_space::private_space,
+                                          decorated::no>();
+  checkPointerAccessAndUnderlyingPointers<address_space::private_space,
+                                          decorated::yes>();
+  checkPointerAccessAndUnderlyingPointers<address_space::private_space,
+                                          decorated::legacy>();
+  checkPointerAccessAndUnderlyingPointers<address_space::generic_space,
+                                          decorated::no>();
+  checkPointerAccessAndUnderlyingPointers<address_space::generic_space,
+                                          decorated::yes>();
+  checkPointerAccessAndUnderlyingPointers<address_space::generic_space,
+                                          decorated::legacy>();
+}
+
+TEST(MultiPtrConversion,
+     GenericSpaceExplicitConversionsSupportAllNonLegacyDecorationPairs) {
+  checkGenericExplicitMutableConversions<decorated::no, decorated::no>();
+  checkGenericExplicitMutableConversions<decorated::no, decorated::yes>();
+  checkGenericExplicitMutableConversions<decorated::yes, decorated::no>();
+  checkGenericExplicitMutableConversions<decorated::yes, decorated::yes>();
+
+  checkGenericExplicitConstConversions<decorated::no, decorated::no>();
+  checkGenericExplicitConstConversions<decorated::no, decorated::yes>();
+  checkGenericExplicitConstConversions<decorated::yes, decorated::no>();
+  checkGenericExplicitConstConversions<decorated::yes, decorated::yes>();
+}
+
+TEST(MultiPtrConversion,
+     NonLegacySameSpaceConversionsCoverVoidConstAndDecorationChanges) {
+  checkNonLegacyMutableSameSpaceConversions<address_space::global_space,
+                                            decorated::no>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::global_space,
+                                            decorated::yes>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::local_space,
+                                            decorated::no>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::local_space,
+                                            decorated::yes>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::private_space,
+                                            decorated::no>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::private_space,
+                                            decorated::yes>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::generic_space,
+                                            decorated::no>();
+  checkNonLegacyMutableSameSpaceConversions<address_space::generic_space,
+                                            decorated::yes>();
+
+  checkNonLegacyConstSameSpaceConversions<address_space::global_space,
+                                          decorated::no>();
+  checkNonLegacyConstSameSpaceConversions<address_space::global_space,
+                                          decorated::yes>();
+  checkNonLegacyConstSameSpaceConversions<address_space::local_space,
+                                          decorated::no>();
+  checkNonLegacyConstSameSpaceConversions<address_space::local_space,
+                                          decorated::yes>();
+  checkNonLegacyConstSameSpaceConversions<address_space::constant_space,
+                                          decorated::no>();
+  checkNonLegacyConstSameSpaceConversions<address_space::constant_space,
+                                          decorated::yes>();
+  checkNonLegacyConstSameSpaceConversions<address_space::private_space,
+                                          decorated::no>();
+  checkNonLegacyConstSameSpaceConversions<address_space::private_space,
+                                          decorated::yes>();
+  checkNonLegacyConstSameSpaceConversions<address_space::generic_space,
+                                          decorated::no>();
+  checkNonLegacyConstSameSpaceConversions<address_space::generic_space,
+                                          decorated::yes>();
+}
+
+TEST(MultiPtrConversion, LegacyConversionsCoverAllAddressSpaces) {
+  checkLegacyConversions<address_space::global_space>();
+  checkLegacyConversions<address_space::local_space>();
+  checkLegacyConversions<address_space::constant_space>();
+  checkLegacyConversions<address_space::private_space>();
+  checkLegacyConversions<address_space::generic_space>();
+}
+
+TEST(MultiPtrConversion, FactoryFunctionsPreservePrivatePointersAndNullptr) {
+  checkFactoryFunctions<decorated::no>();
+  checkFactoryFunctions<decorated::yes>();
+  checkFactoryFunctions<decorated::legacy>();
+  checkDefaultLegacyMakePtr();
+}
+
+} // namespace
diff --git a/sycl/unittests/multi_ptr/NegativeConversions.cpp b/sycl/unittests/multi_ptr/NegativeConversions.cpp
new file mode 100644
index 0000000000000..0d88942517b5a
--- /dev/null
+++ b/sycl/unittests/multi_ptr/NegativeConversions.cpp
@@ -0,0 +1,126 @@
+//===------------------ NegativeConversions.cpp ---------------===//
+//
+// 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 <functional>
+#include <sycl/multi_ptr.hpp>
+
+#include <gtest/gtest.h>
+
+#include <type_traits>
+
+namespace {
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+using multi_ptr_t = sycl::multi_ptr<ElementType, Space, IsDecorated>;
+
+TEST(MultiPtrNegativeConversions, CannotRemoveConstQualification) {
+  using const_ptr =
+      multi_ptr_t<const int, address_space::private_space, decorated::no>;
+  using mutable_ptr =
+      multi_ptr_t<int, address_space::private_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_constructible_v<mutable_ptr, const_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<const_ptr, mutable_ptr>));
+  EXPECT_FALSE((std::is_assignable_v<mutable_ptr, const_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotConvertBetweenIncompatibleTypes) {
+  using int_ptr = multi_ptr_t<int, address_space::private_space, decorated::no>;
+  using float_ptr =
+      multi_ptr_t<float, address_space::private_space, decorated::no>;
+  using double_ptr =
+      multi_ptr_t<double, address_space::private_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_constructible_v<int_ptr, float_ptr>));
+  EXPECT_FALSE((std::is_constructible_v<float_ptr, int_ptr>));
+  EXPECT_FALSE((std::is_constructible_v<int_ptr, double_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<int_ptr, float_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<float_ptr, double_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotConvertToGenericFromNonGeneric) {
+  using private_ptr =
+      multi_ptr_t<int, address_space::private_space, decorated::no>;
+  using local_ptr = multi_ptr_t<int, address_space::local_space, decorated::no>;
+  using global_ptr =
+      multi_ptr_t<int, address_space::global_space, decorated::no>;
+  using generic_ptr =
+      multi_ptr_t<int, address_space::generic_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_convertible_v<private_ptr, generic_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<local_ptr, generic_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<global_ptr, generic_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions,
+     CannotImplicitlyConvertFromGenericToSpecificSpace) {
+  using generic_ptr =
+      multi_ptr_t<int, address_space::generic_space, decorated::no>;
+  using private_ptr =
+      multi_ptr_t<int, address_space::private_space, decorated::no>;
+  using local_ptr = multi_ptr_t<int, address_space::local_space, decorated::no>;
+  using global_ptr =
+      multi_ptr_t<int, address_space::global_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_convertible_v<generic_ptr, private_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<generic_ptr, local_ptr>));
+  EXPECT_FALSE((std::is_convertible_v<generic_ptr, global_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotConvertVoidToTypedWithoutExplicitCast) {
+  using void_ptr =
+      multi_ptr_t<void, address_space::private_space, decorated::no>;
+  using int_ptr = multi_ptr_t<int, address_space::private_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_convertible_v<void_ptr, int_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotConstructMutableFromConstTypedVoid) {
+  using const_void_ptr =
+      multi_ptr_t<const void, address_space::private_space, decorated::no>;
+  using int_ptr = multi_ptr_t<int, address_space::private_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_constructible_v<int_ptr, const_void_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotAssignBetweenIncompatibleSpaces) {
+  using private_ptr =
+      multi_ptr_t<int, address_space::private_space, decorated::no>;
+  using local_ptr = multi_ptr_t<int, address_space::local_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_assignable_v<private_ptr, local_ptr>));
+  EXPECT_FALSE((std::is_assignable_v<local_ptr, private_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, GenericCannotAssignFromConstantSpace) {
+  using generic_ptr =
+      multi_ptr_t<const int, address_space::generic_space, decorated::no>;
+  using constant_ptr =
+      multi_ptr_t<const int, address_space::constant_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_assignable_v<generic_ptr, constant_ptr>));
+}
+
+TEST(MultiPtrNegativeConversions, ConstantSpaceVoidInNonLegacy) {
+  using const_void_constant_legacy =
+      multi_ptr_t<const void, address_space::constant_space, decorated::legacy>;
+  EXPECT_TRUE((std::is_default_constructible_v<const_void_constant_legacy>));
+}
+
+TEST(MultiPtrNegativeConversions, CannotComparePointersOfDifferentTypes) {
+  using int_ptr = multi_ptr_t<int, address_space::private_space, decorated::no>;
+  using float_ptr =
+      multi_ptr_t<float, address_space::private_space, decorated::no>;
+
+  EXPECT_FALSE((std::is_invocable_v<std::equal_to<>, int_ptr, float_ptr>));
+}
+
+} // namespace
diff --git a/sycl/unittests/multi_ptr/Operators.cpp b/sycl/unittests/multi_ptr/Operators.cpp
new file mode 100644
index 0000000000000..e35dc066f2499
--- /dev/null
+++ b/sycl/unittests/multi_ptr/Operators.cpp
@@ -0,0 +1,256 @@
+//===---------------------------- Operators.cpp ------------===//
+//
+// 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 <gtest/gtest.h>
+
+#include <memory>
+#include <type_traits>
+#include <utility>
+
+namespace {
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+template <typename ElementType, decorated IsDecorated>
+using private_multi_ptr =
+    sycl::multi_ptr<ElementType, address_space::private_space, IsDecorated>;
+
+template <typename ElementType, decorated IsDecorated>
+private_multi_ptr<ElementType, IsDecorated>
+makePrivatePtr(std::add_pointer_t<ElementType> Ptr) {
+  return sycl::address_space_cast<address_space::private_space, IsDecorated>(
+      Ptr);
+}
+
+struct Record {
+  int Value;
+  int Tag;
+};
+
+template <decorated IsDecorated> void checkOperatorTraits() {
+  using ptr_t = private_multi_ptr<Record, IsDecorated>;
+  using diff_t = typename ptr_t::difference_type;
+
+  EXPECT_TRUE((std::is_same_v<decltype(*std::declval<const ptr_t &>()),
+                              typename ptr_t::reference>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>()[diff_t{0}]),
+                      typename ptr_t::reference>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>().operator->()),
+                      typename ptr_t::pointer>));
+
+  EXPECT_TRUE((std::is_same_v<decltype(++std::declval<ptr_t &>()), ptr_t &>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<ptr_t &>()++), ptr_t>));
+  EXPECT_TRUE((std::is_same_v<decltype(--std::declval<ptr_t &>()), ptr_t &>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<ptr_t &>()--), ptr_t>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(std::declval<ptr_t &>() += diff_t{0}), ptr_t &>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(std::declval<ptr_t &>() -= diff_t{0}), ptr_t &>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() + diff_t{0}),
+                      ptr_t>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() - diff_t{0}),
+                      ptr_t>));
+
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() ==
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() !=
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() <
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() >
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() <=
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+  EXPECT_TRUE((std::is_same_v<decltype(std::declval<const ptr_t &>() >=
+                                       std::declval<const ptr_t &>()),
+                              bool>));
+
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() == nullptr),
+                      bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() != nullptr),
+                      bool>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(std::declval<const ptr_t &>() < nullptr), bool>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(std::declval<const ptr_t &>() > nullptr), bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() <= nullptr),
+                      bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(std::declval<const ptr_t &>() >= nullptr),
+                      bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(nullptr == std::declval<const ptr_t &>()),
+                      bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(nullptr != std::declval<const ptr_t &>()),
+                      bool>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(nullptr < std::declval<const ptr_t &>()), bool>));
+  EXPECT_TRUE((
+      std::is_same_v<decltype(nullptr > std::declval<const ptr_t &>()), bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(nullptr <= std::declval<const ptr_t &>()),
+                      bool>));
+  EXPECT_TRUE(
+      (std::is_same_v<decltype(nullptr >= std::declval<const ptr_t &>()),
+                      bool>));
+}
+
+template <decorated IsDecorated> void checkAccessOperators() {
+  Record Values[] = {{10, 100}, {20, 200}, {30, 300}};
+  auto Ptr = makePrivatePtr<Record, IsDecorated>(&Values[0]);
+
+  EXPECT_EQ((*Ptr).Value, 10);
+  EXPECT_EQ((*Ptr).Tag, 100);
+  EXPECT_EQ(Ptr->Value, 10);
+  EXPECT_EQ(Ptr->Tag, 100);
+  EXPECT_EQ(Ptr[1].Value, 20);
+  EXPECT_EQ(Ptr[1].Tag, 200);
+  EXPECT_EQ(Ptr[2].Value, 30);
+  EXPECT_EQ(Ptr[2].Tag, 300);
+}
+
+template <decorated IsDecorated> void checkArithmeticOperators() {
+  Record Values[] = {{1, 10}, {2, 20}, {3, 30}, {4, 40}};
+  using ptr_t = private_multi_ptr<Record, IsDecorated>;
+  using diff_t = typename ptr_t::difference_type;
+
+  auto Ptr = makePrivatePtr<Record, IsDecorated>(&Values[0]);
+  auto Expected0 = makePrivatePtr<Record, IsDecorated>(&Values[0]);
+  auto Expected1 = makePrivatePtr<Record, IsDecorated>(&Values[1]);
+  auto Expected2 = makePrivatePtr<Record, IsDecorated>(&Values[2]);
+  auto Expected3 = makePrivatePtr<Record, IsDecorated>(&Values[3]);
+
+  ptr_t &PrefixIncrement = ++Ptr;
+  EXPECT_EQ(std::addressof(PrefixIncrement), std::addressof(Ptr));
+  EXPECT_EQ(Ptr, Expected1);
+
+  ptr_t PostfixIncrement = Ptr++;
+  EXPECT_EQ(PostfixIncrement, Expected1);
+  EXPECT_EQ(Ptr, Expected2);
+
+  ptr_t &PrefixDecrement = --Ptr;
+  EXPECT_EQ(std::addressof(PrefixDecrement), std::addressof(Ptr));
+  EXPECT_EQ(Ptr, Expected1);
+
+  ptr_t PostfixDecrement = Ptr--;
+  EXPECT_EQ(PostfixDecrement, Expected1);
+  EXPECT_EQ(Ptr, Expected0);
+
+  ptr_t &PlusAssign = (Ptr += diff_t{3});
+  EXPECT_EQ(std::addressof(PlusAssign), std::addressof(Ptr));
+  EXPECT_EQ(Ptr, Expected3);
+
+  ptr_t &MinusAssign = (Ptr -= diff_t{2});
+  EXPECT_EQ(std::addressof(MinusAssign), std::addressof(Ptr));
+  EXPECT_EQ(Ptr, Expected1);
+
+  ptr_t PlusResult = Ptr + diff_t{2};
+  EXPECT_EQ(PlusResult, Expected3);
+  EXPECT_EQ(Ptr, Expected1);
+
+  ptr_t MinusResult = Ptr - diff_t{1};
+  EXPECT_EQ(MinusResult, Expected0);
+  EXPECT_EQ(Ptr, Expected1);
+}
+
+template <decorated IsDecorated> void checkRelationalOperators() {
+  int Values[] = {5, 6, 7};
+
+  auto First = makePrivatePtr<int, IsDecorated>(&Values[0]);
+  auto SameAsFirst = makePrivatePtr<int, IsDecorated>(&Values[0]);
+  auto Second = makePrivatePtr<int, IsDecorated>(&Values[1]);
+
+  EXPECT_TRUE(First == SameAsFirst);
+  EXPECT_FALSE(First != SameAsFirst);
+  EXPECT_FALSE(First < SameAsFirst);
+  EXPECT_FALSE(First > SameAsFirst);
+  EXPECT_TRUE(First <= SameAsFirst);
+  EXPECT_TRUE(First >= SameAsFirst);
+
+  EXPECT_FALSE(First == Second);
+  EXPECT_TRUE(First != Second);
+  EXPECT_TRUE(First < Second);
+  EXPECT_FALSE(First > Second);
+  EXPECT_TRUE(First <= Second);
+  EXPECT_FALSE(First >= Second);
+
+  EXPECT_FALSE(Second == First);
+  EXPECT_TRUE(Second != First);
+  EXPECT_FALSE(Second < First);
+  EXPECT_TRUE(Second > First);
+  EXPECT_FALSE(Second <= First);
+  EXPECT_TRUE(Second >= First);
+}
+
+template <decorated IsDecorated> void checkNullptrComparisonOperators() {
+  int Value = 9;
+  auto NullPtr = private_multi_ptr<int, IsDecorated>{nullptr};
+  auto NonNullPtr = makePrivatePtr<int, IsDecorated>(&Value);
+
+  EXPECT_TRUE(NullPtr == nullptr);
+  EXPECT_FALSE(NullPtr != nullptr);
+  EXPECT_FALSE(NullPtr < nullptr);
+  EXPECT_FALSE(NullPtr > nullptr);
+  EXPECT_TRUE(NullPtr <= nullptr);
+  EXPECT_TRUE(NullPtr >= nullptr);
+
+  EXPECT_TRUE(nullptr == NullPtr);
+  EXPECT_FALSE(nullptr != NullPtr);
+  EXPECT_FALSE(nullptr < NullPtr);
+  EXPECT_FALSE(nullptr > NullPtr);
+  EXPECT_TRUE(nullptr <= NullPtr);
+  EXPECT_TRUE(nullptr >= NullPtr);
+
+  EXPECT_FALSE(NonNullPtr == nullptr);
+  EXPECT_TRUE(NonNullPtr != nullptr);
+  EXPECT_FALSE(nullptr == NonNullPtr);
+  EXPECT_TRUE(nullptr != NonNullPtr);
+}
+
+TEST(MultiPtrOperators, OperatorTraitsAreSatisfied) {
+  checkOperatorTraits<decorated::no>();
+  checkOperatorTraits<decorated::yes>();
+}
+
+TEST(MultiPtrOperators, AccessOperatorsProvidePointerLikeSemantics) {
+  checkAccessOperators<decorated::no>();
+  checkAccessOperators<decorated::yes>();
+}
+
+TEST(MultiPtrOperators, ArithmeticOperatorsProvideRandomAccessTraversal) {
+  checkArithmeticOperators<decorated::no>();
+  checkArithmeticOperators<decorated::yes>();
+}
+
+TEST(MultiPtrOperators, RelationalOperatorsCompareUnderlyingLocations) {
+  checkRelationalOperators<decorated::no>();
+  checkRelationalOperators<decorated::yes>();
+}
+
+TEST(MultiPtrOperators, NullptrComparisonOperatorsHandleNullValuesBothWays) {
+  checkNullptrComparisonOperators<decorated::no>();
+  checkNullptrComparisonOperators<decorated::yes>();
+}
+
+} // namespace
diff --git a/sycl/unittests/multi_ptr/VoidSpecialization.cpp b/sycl/unittests/multi_ptr/VoidSpecialization.cpp
new file mode 100644
index 0000000000000..254b20de98c82
--- /dev/null
+++ b/sycl/unittests/multi_ptr/VoidSpecialization.cpp
@@ -0,0 +1,354 @@
+//===------------------- VoidSpecialization.cpp ----------------===//
+//
+// 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 <utility>
+
+#include <gtest/gtest.h>
+
+#include <type_traits>
+
+namespace {
+
+using sycl::access::address_space;
+using sycl::access::decorated;
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+using multi_ptr_t = sycl::multi_ptr<ElementType, Space, IsDecorated>;
+
+template <typename ElementType, address_space Space, decorated IsDecorated>
+multi_ptr_t<ElementType, Space, IsDecorated>
+makePtr(std::add_pointer_t<ElementType> Ptr) {
+  if constexpr (IsDecorated == decorated::legacy) {
+    return multi_ptr_t<ElementType, Space, IsDecorated>{Ptr};
+  } else {
+    return sycl::address_space_cast<Space, IsDecorated>(Ptr);
+  }
+}
+
+struct TestData {
+  int Value;
+  float Score;
+};
+
+template <decorated IsDecorated> void checkVoidPointerConstruction() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+  using const_void_ptr =
+      multi_ptr_t<const void, address_space::private_space, IsDecorated>;
+
+  void_ptr DefaultVoid;
+  const_void_ptr DefaultConstVoid;
+  EXPECT_EQ(DefaultVoid, nullptr);
+  EXPECT_EQ(DefaultConstVoid, nullptr);
+
+  void_ptr NullVoid{nullptr};
+  const_void_ptr NullConstVoid{nullptr};
+  EXPECT_EQ(NullVoid, nullptr);
+  EXPECT_EQ(NullConstVoid, nullptr);
+
+  TestData Data{42, 3.14f};
+  auto TypedPtr =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data);
+
+  void_ptr VoidFromTyped = TypedPtr;
+  EXPECT_NE(VoidFromTyped, nullptr);
+  // Convert back to typed to verify it points to the right data
+  auto TypedCheck = static_cast<
+      multi_ptr_t<TestData, address_space::private_space, IsDecorated>>(
+      VoidFromTyped);
+  EXPECT_EQ(TypedCheck->Value, 42);
+
+  // For const void, need to go through const typed pointer first
+  auto ConstTypedPtr = static_cast<
+      multi_ptr_t<const TestData, address_space::private_space, IsDecorated>>(
+      TypedPtr);
+  const_void_ptr ConstVoidFromTyped = ConstTypedPtr;
+  EXPECT_NE(ConstVoidFromTyped, nullptr);
+  // Convert back to const typed to verify
+  auto ConstTypedCheck = static_cast<
+      multi_ptr_t<const TestData, address_space::private_space, IsDecorated>>(
+      ConstVoidFromTyped);
+  EXPECT_EQ(ConstTypedCheck->Value, 42);
+}
+
+template <decorated IsDecorated> void checkVoidPointerCopyAndMove() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+
+  TestData Data{17, 2.71f};
+  auto TypedPtr =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data);
+  void_ptr VoidPtr = TypedPtr;
+
+  void_ptr CopiedVoid{VoidPtr};
+  EXPECT_EQ(CopiedVoid, VoidPtr);
+  // Verify by converting back to typed
+  auto TypedCheck = static_cast<
+      multi_ptr_t<TestData, address_space::private_space, IsDecorated>>(
+      CopiedVoid);
+  EXPECT_EQ(TypedCheck->Value, 17);
+
+  void_ptr MovedVoid{std::move(VoidPtr)};
+  auto TypedCheck2 = static_cast<
+      multi_ptr_t<TestData, address_space::private_space, IsDecorated>>(
+      MovedVoid);
+  EXPECT_EQ(TypedCheck2->Value, 17);
+}
+
+template <decorated IsDecorated> void checkVoidPointerAssignment() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+
+  TestData Data1{10, 1.0f};
+  TestData Data2{20, 2.0f};
+
+  auto TypedPtr1 =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data1);
+  auto TypedPtr2 =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data2);
+
+  void_ptr VoidPtr1 = TypedPtr1;
+  void_ptr VoidPtr2 = TypedPtr2;
+
+  EXPECT_NE(VoidPtr1, VoidPtr2);
+
+  VoidPtr1 = VoidPtr2;
+  EXPECT_EQ(VoidPtr1, VoidPtr2);
+  // Verify by converting back
+  auto TypedCheck = static_cast<
+      multi_ptr_t<TestData, address_space::private_space, IsDecorated>>(
+      VoidPtr1);
+  EXPECT_EQ(TypedCheck->Value, 20);
+
+  VoidPtr1 = nullptr;
+  EXPECT_EQ(VoidPtr1, nullptr);
+  EXPECT_NE(VoidPtr2, nullptr);
+}
+
+template <decorated IsDecorated>
+void checkVoidPointerExplicitConversionToTyped() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+  using int_ptr = multi_ptr_t<int, address_space::private_space, IsDecorated>;
+  using data_ptr =
+      multi_ptr_t<TestData, address_space::private_space, IsDecorated>;
+
+  TestData Data{99, 9.9f};
+  auto TypedPtr =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data);
+  void_ptr VoidPtr = TypedPtr;
+
+  data_ptr RoundTrip = static_cast<data_ptr>(VoidPtr);
+  EXPECT_EQ(RoundTrip->Value, 99);
+  EXPECT_EQ(RoundTrip->Score, 9.9f);
+
+  int Value = 123;
+  auto IntTypedPtr =
+      makePtr<int, address_space::private_space, IsDecorated>(&Value);
+  void_ptr IntVoidPtr = IntTypedPtr;
+  int_ptr IntRoundTrip = static_cast<int_ptr>(IntVoidPtr);
+  EXPECT_EQ(*IntRoundTrip, 123);
+}
+
+template <decorated IsDecorated> void checkConstVoidPointerConstConversions() {
+  using const_void_ptr =
+      multi_ptr_t<const void, address_space::private_space, IsDecorated>;
+  using const_int_ptr =
+      multi_ptr_t<const int, address_space::private_space, IsDecorated>;
+
+  int Value = 456;
+  auto MutablePtr =
+      makePtr<int, address_space::private_space, IsDecorated>(&Value);
+
+  // Convert mutable to const typed first, then to const void
+  const_int_ptr ConstIntPtr1 = MutablePtr;
+  const_void_ptr ConstVoidPtr = ConstIntPtr1;
+  EXPECT_NE(ConstVoidPtr.get(), nullptr);
+
+  const_int_ptr ConstIntPtr = static_cast<const_int_ptr>(ConstVoidPtr);
+  EXPECT_EQ(*ConstIntPtr, 456);
+
+  const int ConstValue = 789;
+  auto ConstTypedPtr =
+      makePtr<const int, address_space::private_space, IsDecorated>(
+          &ConstValue);
+  const_void_ptr ConstVoidPtr2 = ConstTypedPtr;
+  const_int_ptr ConstIntPtr2 = static_cast<const_int_ptr>(ConstVoidPtr2);
+  EXPECT_EQ(*ConstIntPtr2, 789);
+}
+
+template <decorated IsDecorated> void checkVoidPointerComparisons() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+
+  TestData Data1{1, 1.0f};
+  TestData Data2{2, 2.0f};
+
+  auto TypedPtr1 =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data1);
+  auto TypedPtr2 =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data2);
+
+  void_ptr VoidPtr1 = TypedPtr1;
+  void_ptr VoidPtr2 = TypedPtr2;
+  void_ptr VoidPtr1Copy = TypedPtr1;
+
+  EXPECT_TRUE(VoidPtr1 == VoidPtr1Copy);
+  EXPECT_FALSE(VoidPtr1 != VoidPtr1Copy);
+  EXPECT_TRUE(VoidPtr1 != VoidPtr2);
+  EXPECT_FALSE(VoidPtr1 == VoidPtr2);
+
+  void_ptr NullPtr = nullptr;
+  EXPECT_TRUE(NullPtr == nullptr);
+  EXPECT_FALSE(NullPtr != nullptr);
+  EXPECT_TRUE(nullptr == NullPtr);
+  EXPECT_FALSE(nullptr != NullPtr);
+
+  EXPECT_FALSE(VoidPtr1 == nullptr);
+  EXPECT_TRUE(VoidPtr1 != nullptr);
+}
+
+template <decorated IsDecorated> void checkVoidPointerGetMethods() {
+  using void_ptr = multi_ptr_t<void, address_space::private_space, IsDecorated>;
+
+  TestData Data{50, 5.0f};
+  auto TypedPtr =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data);
+  void_ptr VoidPtr = TypedPtr;
+
+  EXPECT_NE(VoidPtr.get(), nullptr);
+
+  if constexpr (IsDecorated == decorated::legacy) {
+    EXPECT_EQ(VoidPtr.get_raw(), static_cast<void *>(&Data));
+    EXPECT_EQ(VoidPtr.get(), VoidPtr.get_decorated());
+  } else
+    EXPECT_NE(VoidPtr.get(), nullptr);
+}
+
+template <decorated IsDecorated> void checkVoidPointerDecorationConversion() {
+  using void_no =
+      multi_ptr_t<void, address_space::private_space, decorated::no>;
+  using void_yes =
+      multi_ptr_t<void, address_space::private_space, decorated::yes>;
+
+  TestData Data{100, 10.0f};
+  auto TypedPtr =
+      makePtr<TestData, address_space::private_space, IsDecorated>(&Data);
+
+  if constexpr (IsDecorated != decorated::legacy) {
+    auto VoidPtr = static_cast<
+        multi_ptr_t<void, address_space::private_space, IsDecorated>>(TypedPtr);
+
+    if constexpr (IsDecorated == decorated::no) {
+      // Test implicit conversion to decorated::yes
+      void_yes ToDecorated = VoidPtr;
+      EXPECT_NE(ToDecorated.get(), nullptr);
+
+      // Convert back to typed pointer to verify correctness
+      auto TypedCheck = static_cast<
+          multi_ptr_t<TestData, address_space::private_space, decorated::yes>>(
+          ToDecorated);
+      EXPECT_EQ(TypedCheck->Value, 100);
+    } else {
+      // Test implicit conversion to decorated::no
+      void_no ToUndecorated = VoidPtr;
+      EXPECT_NE(ToUndecorated.get(), nullptr);
+
+      // Convert back to typed pointer to verify correctness
+      auto TypedCheck = static_cast<
+          multi_ptr_t<TestData, address_space::private_space, decorated::no>>(
+          ToUndecorated);
+      EXPECT_EQ(TypedCheck->Value, 100);
+    }
+  }
+}
+
+TEST(MultiPtrVoidSpecialization,
+     DefaultAndNullptrConstructionCreatesNullPointers) {
+  checkVoidPointerConstruction<decorated::no>();
+  checkVoidPointerConstruction<decorated::yes>();
+  checkVoidPointerConstruction<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, CopyAndMoveConstructionPreservesPointerValue) {
+  checkVoidPointerCopyAndMove<decorated::no>();
+  checkVoidPointerCopyAndMove<decorated::yes>();
+  checkVoidPointerCopyAndMove<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, AssignmentOperatorsWorkCorrectly) {
+  checkVoidPointerAssignment<decorated::no>();
+  checkVoidPointerAssignment<decorated::yes>();
+  checkVoidPointerAssignment<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, ExplicitConversionToTypedPointerWorks) {
+  checkVoidPointerExplicitConversionToTyped<decorated::no>();
+  checkVoidPointerExplicitConversionToTyped<decorated::yes>();
+  checkVoidPointerExplicitConversionToTyped<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, ConstVoidPointerHandlesConstCorrectly) {
+  checkConstVoidPointerConstConversions<decorated::no>();
+  checkConstVoidPointerConstConversions<decorated::yes>();
+  checkConstVoidPointerConstConversions<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, ComparisonOperatorsWorkCorrectly) {
+  checkVoidPointerComparisons<decorated::no>();
+  checkVoidPointerComparisons<decorated::yes>();
+  checkVoidPointerComparisons<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, GetMethodsReturnCorrectPointers) {
+  checkVoidPointerGetMethods<decorated::no>();
+  checkVoidPointerGetMethods<decorated::yes>();
+  checkVoidPointerGetMethods<decorated::legacy>();
+}
+
+TEST(MultiPtrVoidSpecialization, DecorationConversionsBetweenModesWork) {
+  checkVoidPointerDecorationConversion<decorated::no>();
+  checkVoidPointerDecorationConversion<decorated::yes>();
+}
+
+TEST(MultiPtrVoidSpecialization, VoidPointerWorksWithDifferentAddressSpaces) {
+  using global_void =
+      multi_ptr_t<void, address_space::global_space, decorated::no>;
+  using local_void =
+      multi_ptr_t<void, address_space::local_space, decorated::no>;
+  using generic_void =
+      multi_ptr_t<void, address_space::generic_space, decorated::no>;
+
+  EXPECT_TRUE((std::is_default_constructible_v<global_void>));
+  EXPECT_TRUE((std::is_default_constructible_v<local_void>));
+  EXPECT_TRUE((std::is_default_constructible_v<generic_void>));
+
+  global_void GlobalNull;
+  local_void LocalNull;
+  generic_void GenericNull;
+
+  EXPECT_EQ(GlobalNull, nullptr);
+  EXPECT_EQ(LocalNull, nullptr);
+  EXPECT_EQ(GenericNull, nullptr);
+}
+
+TEST(MultiPtrVoidSpecialization, LegacyVoidPointerConstructsFromTypedPointer) {
+  using void_legacy =
+      multi_ptr_t<void, address_space::private_space, decorated::legacy>;
+  using int_legacy =
+      multi_ptr_t<int, address_space::private_space, decorated::legacy>;
+
+  int Value = 999;
+  int_legacy IntPtr{&Value};
+
+  void_legacy VoidPtr{IntPtr};
+  // Legacy void has get_raw()
+  EXPECT_EQ(VoidPtr.get_raw(), static_cast<void *>(&Value));
+
+  int_legacy RoundTrip = static_cast<int_legacy>(VoidPtr);
+  EXPECT_EQ(RoundTrip.get_raw(), &Value);
+  EXPECT_EQ(*RoundTrip, 999);
+}
+
+} // namespace