diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp
index 7cf893778394c..bd59361ff2eeb 100644
--- a/sycl/include/sycl/detail/generic_type_traits.hpp
+++ b/sycl/include/sycl/detail/generic_type_traits.hpp
@@ -663,6 +663,14 @@ template <> struct ConvertToOpenCLTypeImpl<Boolean<1>> {
   // Or should it be "int"?
   using type = Boolean<1>;
 };
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+// TODO: It seems we only use this to convert a pointer's element type. As such,
+// although it doesn't look very clean, it should be ok having this case handled
+// explicitly until further refactoring of this area.
+template <> struct ConvertToOpenCLTypeImpl<std::byte> {
+  using type = uint8_t;
+};
+#endif
 #endif
 
 template <typename T> struct ConvertToOpenCLTypeImpl<T *> {
@@ -700,8 +708,30 @@ convertDataToType(FROM t) {
 // Now fuse the above into a simpler helper that's easy to use.
 // TODO: That should probably be moved outside of "type_traits".
 template <typename T> auto convertToOpenCLType(T &&x) {
-  using OpenCLType = ConvertToOpenCLType_t<std::remove_reference_t<T>>;
-  return convertDataToType<T, OpenCLType>(std::forward<T>(x));
+  using no_ref = std::remove_reference_t<T>;
+  if constexpr (is_multi_ptr_v<no_ref>) {
+    return convertToOpenCLType(x.get_decorated());
+  } else if constexpr (std::is_pointer_v<no_ref>) {
+    // TODO: Below ignores volatile, but we didn't have a need for it yet.
+    using elem_type = remove_decoration_t<std::remove_pointer_t<no_ref>>;
+    using converted_elem_type_no_cv =
+        ConvertToOpenCLType_t<std::remove_const_t<elem_type>>;
+    using converted_elem_type =
+        std::conditional_t<std::is_const_v<elem_type>,
+                           const converted_elem_type_no_cv,
+                           converted_elem_type_no_cv>;
+#ifdef __SYCL_DEVICE_ONLY__
+    using result_type =
+        typename DecoratedType<converted_elem_type,
+                               deduce_AS<no_ref>::value>::type *;
+#else
+    using result_type = converted_elem_type *;
+#endif
+    return reinterpret_cast<result_type>(x);
+  } else {
+    using OpenCLType = ConvertToOpenCLType_t<no_ref>;
+    return convertDataToType<T, OpenCLType>(std::forward<T>(x));
+  }
 }
 
 template <typename To, typename From> auto convertFromOpenCLTypeFor(From &&x) {
diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp
index ceb0c58dcf99c..faeacd10f1998 100644
--- a/sycl/include/sycl/group.hpp
+++ b/sycl/include/sycl/group.hpp
@@ -315,12 +315,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
                                                        global_ptr<dataT> src,
                                                        size_t numElements,
                                                        size_t srcStride) const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
     __ocl_event_t E = __SYCL_OpGroupAsyncCopyGlobalToLocal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, srcStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
+        detail::convertToOpenCLType(src), numElements, srcStride, 0);
     return device_event(E);
   }
 
@@ -337,12 +334,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
                                                        size_t numElements,
                                                        size_t destStride)
       const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
     __ocl_event_t E = __SYCL_OpGroupAsyncCopyLocalToGlobal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, destStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
+        detail::convertToOpenCLType(src), numElements, destStride, 0);
     return device_event(E);
   }
 
@@ -359,12 +353,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
   async_work_group_copy(decorated_local_ptr<DestDataT> dest,
                         decorated_global_ptr<SrcDataT> src, size_t numElements,
                         size_t srcStride) const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
     __ocl_event_t E = __SYCL_OpGroupAsyncCopyGlobalToLocal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, srcStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
+        detail::convertToOpenCLType(src), numElements, srcStride, 0);
     return device_event(E);
   }
 
@@ -381,12 +372,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
   async_work_group_copy(decorated_global_ptr<DestDataT> dest,
                         decorated_local_ptr<SrcDataT> src, size_t numElements,
                         size_t destStride) const {
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
     __ocl_event_t E = __SYCL_OpGroupAsyncCopyLocalToGlobal(
-        __spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
-        numElements, destStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
+        detail::convertToOpenCLType(src), numElements, destStride, 0);
     return device_event(E);
   }
 
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
index 8cfd4ea9b555b..908b5d4aebff0 100644
--- a/sycl/include/sycl/sub_group.hpp
+++ b/sycl/include/sycl/sub_group.hpp
@@ -42,6 +42,24 @@ namespace sub_group {
 template <typename T>
 using SelectBlockT = select_cl_scalar_integral_unsigned_t<T>;
 
+template <typename MultiPtrTy> auto convertToBlockPtr(MultiPtrTy MultiPtr) {
+  static_assert(is_multi_ptr_v<MultiPtrTy>);
+  auto DecoratedPtr = convertToOpenCLType(MultiPtr);
+  using DecoratedPtrTy = decltype(DecoratedPtr);
+  using ElemTy = remove_decoration_t<std::remove_pointer_t<DecoratedPtrTy>>;
+
+  using TargetElemTy = SelectBlockT<ElemTy>;
+  // TODO: Handle cv qualifiers.
+#ifdef __SYCL_DEVICE_ONLY__
+  using ResultTy =
+      typename DecoratedType<TargetElemTy,
+                             deduce_AS<DecoratedPtrTy>::value>::type *;
+#else
+  using ResultTy = TargetElemTy *;
+#endif
+  return reinterpret_cast<ResultTy>(DecoratedPtr);
+}
+
 template <typename T, access::address_space Space>
 using AcceptableForGlobalLoadStore =
     std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
@@ -57,11 +75,7 @@ template <typename T, access::address_space Space,
           access::decorated DecorateAddress>
 T load(const multi_ptr<T, Space, DecorateAddress> src) {
   using BlockT = SelectBlockT<T>;
-  using PtrT = sycl::detail::ConvertToOpenCLType_t<
-      const multi_ptr<BlockT, Space, DecorateAddress>>;
-
-  BlockT Ret =
-      __spirv_SubgroupBlockReadINTEL<BlockT>(reinterpret_cast<PtrT>(src.get()));
+  BlockT Ret = __spirv_SubgroupBlockReadINTEL<BlockT>(convertToBlockPtr(src));
 
   return sycl::bit_cast<T>(Ret);
 }
@@ -71,11 +85,7 @@ template <int N, typename T, access::address_space Space,
 vec<T, N> load(const multi_ptr<T, Space, DecorateAddress> src) {
   using BlockT = SelectBlockT<T>;
   using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
-  using PtrT = sycl::detail::ConvertToOpenCLType_t<
-      const multi_ptr<BlockT, Space, DecorateAddress>>;
-
-  VecT Ret =
-      __spirv_SubgroupBlockReadINTEL<VecT>(reinterpret_cast<PtrT>(src.get()));
+  VecT Ret = __spirv_SubgroupBlockReadINTEL<VecT>(convertToBlockPtr(src));
 
   return sycl::bit_cast<typename vec<T, N>::vector_t>(Ret);
 }
@@ -84,10 +94,8 @@ template <typename T, access::address_space Space,
           access::decorated DecorateAddress>
 void store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) {
   using BlockT = SelectBlockT<T>;
-  using PtrT = sycl::detail::ConvertToOpenCLType_t<
-      multi_ptr<BlockT, Space, DecorateAddress>>;
 
-  __spirv_SubgroupBlockWriteINTEL(reinterpret_cast<PtrT>(dst.get()),
+  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
                                   sycl::bit_cast<BlockT>(x));
 }
 
@@ -96,10 +104,8 @@ template <int N, typename T, access::address_space Space,
 void store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) {
   using BlockT = SelectBlockT<T>;
   using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
-  using PtrT = sycl::detail::ConvertToOpenCLType_t<
-      const multi_ptr<BlockT, Space, DecorateAddress>>;
 
-  __spirv_SubgroupBlockWriteINTEL(reinterpret_cast<PtrT>(dst.get()),
+  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
                                   sycl::bit_cast<VecT>(x));
 }
 #endif // __SYCL_DEVICE_ONLY__