diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__tuple_dir/vector_types.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__tuple_dir/vector_types.h
new file mode 100644
index 00000000000..ceeb789be15
--- /dev/null
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__tuple_dir/vector_types.h
@@ -0,0 +1,318 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___TUPLE_VECTOR_TYPES_H
+#define _LIBCUDACXX___TUPLE_VECTOR_TYPES_H
+
+#ifndef __cuda_std__
+#  include <__config>
+#endif // __cuda_std__
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#ifdef __cuda_std__
+#  if defined(_CCCL_CUDA_COMPILER)
+
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wmismatched-tags")
+
+#    if !defined(_CCCL_COMPILER_NVRTC)
+// Fetch utility to get primary template for ::std::tuple_size necessary for the specialization of
+// ::std::tuple_size<cuda::std::tuple> to enable structured bindings.
+// See https://github.com/NVIDIA/libcudacxx/issues/316
+#      include <utility>
+#    endif
+
+#    include "../__fwd/get.h"
+#    include "../__tuple_dir/structured_bindings.h"
+#    include "../__tuple_dir/tuple_element.h"
+#    include "../__tuple_dir/tuple_size.h"
+#    include "../__type_traits/integral_constant.h"
+#    include "../__type_traits/enable_if.h"
+#    include "../__utility/forward.h"
+#    include "../__utility/move.h"
+
+#    define _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE(__name, __type, __size)                             \
+      template <>                                                                                      \
+      struct tuple_size<__name##__size> : _CUDA_VSTD::integral_constant<size_t, __size>                \
+      {};                                                                                              \
+      template <>                                                                                      \
+      struct tuple_size<const __name##__size> : _CUDA_VSTD::integral_constant<size_t, __size>          \
+      {};                                                                                              \
+      template <>                                                                                      \
+      struct tuple_size<volatile __name##__size> : _CUDA_VSTD::integral_constant<size_t, __size>       \
+      {};                                                                                              \
+      template <>                                                                                      \
+      struct tuple_size<const volatile __name##__size> : _CUDA_VSTD::integral_constant<size_t, __size> \
+      {};                                                                                              \
+                                                                                                       \
+      template <size_t _Ip>                                                                            \
+      struct tuple_element<_Ip, __name##__size>                                                        \
+      {                                                                                                \
+        static_assert(_Ip < __size, "tuple_element index out of range");                               \
+        using type = __type;                                                                           \
+      };                                                                                               \
+      template <size_t _Ip>                                                                            \
+      struct tuple_element<_Ip, const __name##__size>                                                  \
+      {                                                                                                \
+        static_assert(_Ip < __size, "tuple_element index out of range");                               \
+        using type = const __type;                                                                     \
+      };                                                                                               \
+      template <size_t _Ip>                                                                            \
+      struct tuple_element<_Ip, volatile __name##__size>                                               \
+      {                                                                                                \
+        static_assert(_Ip < __size, "tuple_element index out of range");                               \
+        using type = volatile __type;                                                                  \
+      };                                                                                               \
+      template <size_t _Ip>                                                                            \
+      struct tuple_element<_Ip, const volatile __name##__size>                                         \
+      {                                                                                                \
+        static_assert(_Ip < __size, "tuple_element index out of range");                               \
+        using type = const volatile __type;                                                            \
+      };
+
+#    define _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(__name, __type) \
+      _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE(__name, __type, 1)           \
+      _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE(__name, __type, 2)           \
+      _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE(__name, __type, 3)           \
+      _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE(__name, __type, 4)
+
+#    define _LIBCUDACXX_SPECIALIZE_GET(__name, __base_type)                                                           \
+      template <size_t _Ip>                                                                                           \
+      _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __base_type& get(__name& __val) noexcept        \
+      {                                                                                                               \
+        return _CUDA_VSTD::__get_element<_Ip>::template get<__name, __base_type>(__val);                              \
+      }                                                                                                               \
+      template <size_t _Ip>                                                                                           \
+      _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __base_type& get(                         \
+        const __name& __val) noexcept                                                                                 \
+      {                                                                                                               \
+        return _CUDA_VSTD::__get_element<_Ip>::template get<__name, __base_type>(__val);                              \
+      }                                                                                                               \
+      template <size_t _Ip>                                                                                           \
+      _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __base_type&& get(__name&& __val) noexcept      \
+      {                                                                                                               \
+        return _CUDA_VSTD::__get_element<_Ip>::template get<__name, __base_type>(static_cast<__name&&>(__val));       \
+      }                                                                                                               \
+      template <size_t _Ip>                                                                                           \
+      _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __base_type&& get(                        \
+        const __name&& __val) noexcept                                                                                \
+      {                                                                                                               \
+        return _CUDA_VSTD::__get_element<_Ip>::template get<__name, __base_type>(static_cast<const __name&&>(__val)); \
+      }
+
+#    define _LIBCUDACXX_SPECIALIZE_GET_VECTOR(__name, __base_type) \
+      _LIBCUDACXX_SPECIALIZE_GET(__name##1, __base_type)           \
+      _LIBCUDACXX_SPECIALIZE_GET(__name##2, __base_type)           \
+      _LIBCUDACXX_SPECIALIZE_GET(__name##3, __base_type)           \
+      _LIBCUDACXX_SPECIALIZE_GET(__name##4, __base_type)
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(char, signed char)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(uchar, unsigned char)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(short, short)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ushort, unsigned short)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(int, int)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(uint, unsigned int)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(long, long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ulong, unsigned long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(longlong, long long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ulonglong, unsigned long long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(float, float)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(double, double)
+
+template <size_t _Ip>
+struct __get_element;
+
+template <>
+struct __get_element<0>
+{
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType& get(_Vec& __val) noexcept
+  {
+    return __val.x;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType& get(const _Vec& __val) noexcept
+  {
+    return __val.x;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType&& get(_Vec&& __val) noexcept
+  {
+    return static_cast<_BaseType&&>(__val.x);
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType&&
+  get(const _Vec&& __val) noexcept
+  {
+    return static_cast<const _BaseType&&>(__val.x);
+  }
+};
+
+template <>
+struct __get_element<1>
+{
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType& get(_Vec& __val) noexcept
+  {
+    return __val.y;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType& get(const _Vec& __val) noexcept
+  {
+    return __val.y;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType&& get(_Vec&& __val) noexcept
+  {
+    return static_cast<_BaseType&&>(__val.y);
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType&&
+  get(const _Vec&& __val) noexcept
+  {
+    return static_cast<const _BaseType&&>(__val.y);
+  }
+};
+template <>
+struct __get_element<2>
+{
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType& get(_Vec& __val) noexcept
+  {
+    return __val.z;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType& get(const _Vec& __val) noexcept
+  {
+    return __val.z;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType&& get(_Vec&& __val) noexcept
+  {
+    return static_cast<_BaseType&&>(__val.z);
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType&&
+  get(const _Vec&& __val) noexcept
+  {
+    return static_cast<const _BaseType&&>(__val.z);
+  }
+};
+
+template <>
+struct __get_element<3>
+{
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType& get(_Vec& __val) noexcept
+  {
+    return __val.w;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType& get(const _Vec& __val) noexcept
+  {
+    return __val.w;
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _BaseType&& get(_Vec&& __val) noexcept
+  {
+    return static_cast<_BaseType&&>(__val.w);
+  }
+
+  template <class _Vec, class _BaseType>
+  static _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const _BaseType&&
+  get(const _Vec&& __val) noexcept
+  {
+    return static_cast<const _BaseType&&>(__val.w);
+  }
+};
+
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(char, signed char)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(uchar, unsigned char)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(short, short)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ushort, unsigned short)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(int, int)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(uint, unsigned int)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(long, long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ulong, unsigned long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(longlong, long long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ulonglong, unsigned long long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(float, float)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(double, double)
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+// Those need to be defined in the global namespace because we need ADL to find them
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(char, signed char)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(uchar, unsigned char)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(short, short)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ushort, unsigned short)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(int, int)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(uint, unsigned int)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(long, long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ulong, unsigned long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(longlong, long long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(ulonglong, unsigned long long)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(float, float)
+_LIBCUDACXX_SPECIALIZE_GET_VECTOR(double, double)
+
+// This is a workaround for the fact that structured bindings require that the specializations of
+// `tuple_size` and `tuple_element` reside in namespace std (https://eel.is/c++draft/dcl.struct.bind#4).
+// See https://github.com/NVIDIA/libcudacxx/issues/316 for a short discussion
+#    if _CCCL_STD_VER >= 2017
+namespace std
+{
+
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(char, signed char)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(uchar, unsigned char)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(short, short)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ushort, unsigned short)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(int, int)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(uint, unsigned int)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(long, long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ulong, unsigned long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(longlong, long long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(ulonglong, unsigned long long)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(float, float)
+_LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR(double, double)
+
+} // namespace std
+
+#    endif // _CCCL_STD_VER >= 2017
+
+#    undef _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE
+#    undef _LIBCUDACXX_SPECIALIZE_TUPLE_INTERFACE_VECTOR
+#    undef _LIBCUDACXX_SPECIALIZE_GET
+#    undef _LIBCUDACXX_SPECIALIZE_GET_VECTOR
+
+_CCCL_DIAG_POP
+
+#  endif // _CCCL_CUDA_COMPILER
+#endif // __cuda_std__
+
+#endif // _LIBCUDACXX___TUPLE_VECTOR_TYPES_H
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
index 7a5a1b0657e..2b01c6eae8c 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple
@@ -170,6 +170,7 @@ template <class... Types>
 #include "__tuple_dir/tuple_like.h"
 #include "__tuple_dir/tuple_size.h"
 #include "__tuple_dir/tuple_types.h"
+#include "__tuple_dir/vector_types.h"
 #include "__type_traits/maybe_const.h"
 #include "__utility/forward.h"
 #include "__utility/integer_sequence.h"
diff --git a/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_get.pass.cpp b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_get.pass.cpp
new file mode 100644
index 00000000000..a38455b9b3b
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_get.pass.cpp
@@ -0,0 +1,155 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the libcu++ 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
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/cassert>
+#include <cuda/std/tuple>
+
+#include "test_macros.h"
+
+template <class VType, class BaseType, size_t VSize>
+struct get_val;
+
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 1> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 2> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 3> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337),
+                 static_cast<BaseType>(-1)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 4> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337),
+                 static_cast<BaseType>(-1), static_cast<BaseType>(0)};
+  }
+};
+
+template <class BaseType, size_t Size>
+struct get_expected;
+
+template <class BaseType>
+struct get_expected<BaseType, 0> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(42)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 1> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(1337)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 2> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(-1)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 3> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(0)};
+  }
+};
+
+template <class VType, class BaseType, size_t VSize, size_t Index,
+          cuda::std::__enable_if_t<(Index < VSize), int> = 0>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  { // & overload
+    VType val = get_val<VType, BaseType, VSize>::create();
+    auto&& ret = cuda::std::get<Index>(val);
+    //    static_assert(cuda::std::is_same<decltype(ret), BaseType&>::value, "");
+
+    const BaseType expected = get_expected<BaseType, Index>::create();
+    assert(ret == expected);
+  }
+
+  { // const& overload
+    const VType val = get_val<const VType, BaseType, VSize>::create();
+    auto&& ret = cuda::std::get<Index>(val);
+    //static_assert(cuda::std::is_same<decltype(ret), const BaseType&>::value,"");
+
+    const BaseType expected = get_expected<BaseType, Index>::create();
+    assert(ret == expected);
+  }
+
+  { // && overload
+    VType val = get_val<VType, BaseType, VSize>::create();
+    auto&& ret = cuda::std::get<Index>(cuda::std::move(val));
+    //    static_assert(cuda::std::is_same<decltype(ret), BaseType&&>::value, "");
+
+    const BaseType expected = get_expected<BaseType, Index>::create();
+    assert(ret == expected);
+  }
+
+  { // const && overload
+    const VType val = get_val<const VType, BaseType, VSize>::create();
+    auto&& ret = cuda::std::get<Index>(cuda::std::move(val));
+    //    static_assert(cuda::std::is_same<decltype(ret), const BaseType&&>::value,"");
+
+    const BaseType expected = get_expected<const BaseType, Index>::create();
+    assert(ret == expected);
+  }
+}
+
+template <class VType, class BaseType, size_t VSize, size_t Index,
+          cuda::std::__enable_if_t<(Index >= VSize), int> = 0>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {}
+
+template <class VType, class BaseType, size_t VSize>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  test<VType, BaseType, VSize, 0>();
+  test<VType, BaseType, VSize, 1>();
+  test<VType, BaseType, VSize, 2>();
+  test<VType, BaseType, VSize, 3>();
+}
+
+#define EXPAND_VECTOR_TYPE(Type, BaseType)                                               \
+  test<Type##1, BaseType, 1>();                                                    \
+  test<Type##2, BaseType, 2>();                                                    \
+  test<Type##3, BaseType, 3>();                                                    \
+  test<Type##4, BaseType, 4>();
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test() {
+  EXPAND_VECTOR_TYPE(char, signed char);
+  EXPAND_VECTOR_TYPE(uchar, unsigned char);
+  EXPAND_VECTOR_TYPE(short, short);
+  EXPAND_VECTOR_TYPE(ushort, unsigned short);
+  EXPAND_VECTOR_TYPE(int, int);
+  EXPAND_VECTOR_TYPE(uint, unsigned int);
+  EXPAND_VECTOR_TYPE(long, long);
+  EXPAND_VECTOR_TYPE(ulong, unsigned long);
+  EXPAND_VECTOR_TYPE(longlong, long long);
+  EXPAND_VECTOR_TYPE(ulonglong, unsigned long long);
+  EXPAND_VECTOR_TYPE(float, float);
+  EXPAND_VECTOR_TYPE(double, double);
+
+  return true;
+}
+
+int main(int arg, char** argv) {
+  test();
+#if TEST_STD_VER >= 2014
+  static_assert(test(), "");
+#endif // TEST_STD_VER >= 2014
+
+  return 0;
+}
\ No newline at end of file
diff --git a/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_structured_bindings.pass.cpp b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_structured_bindings.pass.cpp
new file mode 100644
index 00000000000..998265548c9
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_structured_bindings.pass.cpp
@@ -0,0 +1,290 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the libcu++ 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
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++11, c++14
+
+#include <cuda/std/cassert>
+#include <cuda/std/tuple>
+
+#include "test_macros.h"
+
+template <class VType, class BaseType, size_t VSize>
+struct get_val;
+
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 1> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 2> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 3> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337),
+                 static_cast<BaseType>(-1)};
+  }
+};
+template <class VType, class BaseType>
+struct get_val<VType, BaseType, 4> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 VType create() {
+    return VType{static_cast<BaseType>(42), static_cast<BaseType>(1337),
+                 static_cast<BaseType>(-1), static_cast<BaseType>(0)};
+  }
+};
+
+template <class BaseType, size_t Size>
+struct get_expected;
+
+template <class BaseType>
+struct get_expected<BaseType, 0> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(42)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 1> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(1337)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 2> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(-1)};
+  }
+};
+template <class BaseType>
+struct get_expected<BaseType, 3> {
+  __host__ __device__ static TEST_CONSTEXPR_CXX14 BaseType create() {
+    return BaseType{static_cast<BaseType>(0)};
+  }
+};
+
+template <class BaseType, class VType1, class VType2, class VType3,
+          class VType4>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  {   // & overload
+    { // vec1 structured bindings
+      VType1 val = get_val<VType1, BaseType, 1>::create();
+      auto&& [ret] = val;
+      static_assert(cuda::std::is_same<decltype(ret), BaseType>::value, "");
+
+      assert(ret == (get_expected<BaseType, 0>::create()));
+    }
+
+    { // vec2 structured bindings
+      VType2 val = get_val<VType2, BaseType, 2>::create();
+      auto&& [ret1, ret2] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+    }
+
+    { // vec3 structured bindings
+      VType3 val = get_val<VType3, BaseType, 3>::create();
+      auto&& [ret1, ret2, ret3] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+    }
+
+    { // vec4 structured bindings
+      VType4 val = get_val<VType4, BaseType, 4>::create();
+      auto&& [ret1, ret2, ret3, ret4] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret4), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+      assert(ret4 == (get_expected<BaseType, 3>::create()));
+    }
+  }
+
+  {   // const & overload
+    { // vec1 structured bindings
+      const VType1 val = get_val<VType1, BaseType, 1>::create();
+      auto&& [ret] = val;
+      static_assert(cuda::std::is_same<decltype(ret), const BaseType>::value,
+                    "");
+
+      assert(ret == (get_expected<BaseType, 0>::create()));
+    }
+
+    { // vec2 structured bindings
+      const VType2 val = get_val<VType2, BaseType, 2>::create();
+      auto&& [ret1, ret2] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value,
+                    "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+    }
+
+    { // vec3 structured bindings
+      const VType3 val = get_val<VType3, BaseType, 3>::create();
+      auto&& [ret1, ret2, ret3] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret3), const BaseType>::value,
+                    "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+    }
+
+    { // vec4 structured bindings
+      const VType4 val = get_val<VType4, BaseType, 4>::create();
+      auto&& [ret1, ret2, ret3, ret4] = val;
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret3), const BaseType>::value,
+                    "");
+      static_assert(cuda::std::is_same<decltype(ret4), const BaseType>::value,
+                    "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+      assert(ret4 == (get_expected<BaseType, 3>::create()));
+    }
+  }
+
+  {   // && overload
+    { // vec1 structured bindings
+      auto&& [ret] = get_val<VType1, BaseType, 1>::create();
+      static_assert(cuda::std::is_same<decltype(ret), BaseType>::value, "");
+
+      assert(ret == (get_expected<BaseType, 0>::create()));
+    }
+
+    { // vec2 structured bindings
+      auto&& [ret1, ret2] = get_val<VType2, BaseType, 2>::create();
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+    }
+
+    { // vec3 structured bindings
+      auto&& [ret1, ret2, ret3] = get_val<VType3, BaseType, 3>::create();
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+    }
+
+    { // vec4 structured bindings
+      auto&& [ret1, ret2, ret3, ret4] = get_val<VType4, BaseType, 4>::create();
+      static_assert(cuda::std::is_same<decltype(ret1), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret4), BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+      assert(ret4 == (get_expected<BaseType, 3>::create()));
+    }
+  }
+
+  {   // const&& overload
+    { // vec1 structured bindings
+      auto&& [ret] = const_cast<const VType1&&>(get_val<VType1, BaseType, 1>::create());
+      static_assert(cuda::std::is_same<decltype(ret), const BaseType>::value, "");
+
+      assert(ret == (get_expected<BaseType, 0>::create()));
+    }
+
+    { // vec2 structured bindings
+      auto&& [ret1, ret2] = const_cast<const VType2&&>(get_val<VType2, BaseType, 2>::create());
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+    }
+
+    { // vec3 structured bindings
+      auto&& [ret1, ret2, ret3] = const_cast<const VType3&&>(get_val<VType3, BaseType, 3>::create());
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), const BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+    }
+
+    { // vec4 structured bindings
+      auto&& [ret1, ret2, ret3, ret4] = const_cast<const VType4&&>(get_val<VType4, BaseType, 4>::create());
+      static_assert(cuda::std::is_same<decltype(ret1), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret2), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret3), const BaseType>::value, "");
+      static_assert(cuda::std::is_same<decltype(ret4), const BaseType>::value, "");
+
+      assert(ret1 == (get_expected<BaseType, 0>::create()));
+      assert(ret2 == (get_expected<BaseType, 1>::create()));
+      assert(ret3 == (get_expected<BaseType, 2>::create()));
+      assert(ret4 == (get_expected<BaseType, 3>::create()));
+    }
+  }
+}
+
+#define EXPAND_VECTOR_TYPE(Type, BaseType)                                     \
+  test<BaseType, Type##1, Type##2, Type##3, Type##4>();
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test() {
+  EXPAND_VECTOR_TYPE(char, signed char);
+  EXPAND_VECTOR_TYPE(uchar, unsigned char);
+  EXPAND_VECTOR_TYPE(short, short);
+  EXPAND_VECTOR_TYPE(ushort, unsigned short);
+  EXPAND_VECTOR_TYPE(int, int);
+  EXPAND_VECTOR_TYPE(uint, unsigned int);
+  EXPAND_VECTOR_TYPE(long, long);
+  EXPAND_VECTOR_TYPE(ulong, unsigned long);
+  EXPAND_VECTOR_TYPE(longlong, long long);
+  EXPAND_VECTOR_TYPE(ulonglong, unsigned long long);
+  EXPAND_VECTOR_TYPE(float, float);
+  EXPAND_VECTOR_TYPE(double, double);
+
+  return true;
+}
+
+int main(int arg, char** argv) {
+  test();
+  //static_assert(test(), "");
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_element.pass.cpp b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_element.pass.cpp
new file mode 100644
index 00000000000..c6a4eda4116
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_element.pass.cpp
@@ -0,0 +1,69 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the libcu++ 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
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/cassert>
+#include <cuda/std/tuple>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class VType, class BaseType, size_t Index>
+using expected_type = cuda::std::is_same<typename cuda::std::tuple_element<Index, VType>::type, BaseType>;
+
+template <class VType, class BaseType, size_t VSize, size_t Index, cuda::std::__enable_if_t<(Index < VSize), int> = 0>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  assert((expected_type<VType, BaseType, Index>::value));
+  assert((expected_type<const VType, const BaseType, Index>::value));
+  assert((expected_type<volatile VType, volatile BaseType, Index>::value));
+  assert((expected_type<const volatile VType, const volatile BaseType, Index>::value));
+}
+
+template <class VType, class BaseType, size_t VSize, size_t Index, cuda::std::__enable_if_t<(Index >= VSize), int> = 0>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {}
+
+template <class VType, class BaseType, size_t VSize>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  test<VType, BaseType, VSize, 0>();
+  test<VType, BaseType, VSize, 1>();
+  test<VType, BaseType, VSize, 2>();
+  test<VType, BaseType, VSize, 3>();
+}
+
+#define EXPAND_VECTOR_TYPE(Type, BaseType) \
+  test<Type##1, BaseType, 1>();      \
+  test<Type##2, BaseType, 2>();      \
+  test<Type##3, BaseType, 3>();      \
+  test<Type##4, BaseType, 4>();      \
+
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test() {
+  EXPAND_VECTOR_TYPE(char, signed char);
+  EXPAND_VECTOR_TYPE(uchar, unsigned char);
+  EXPAND_VECTOR_TYPE(short, short);
+  EXPAND_VECTOR_TYPE(ushort, unsigned short);
+  EXPAND_VECTOR_TYPE(int, int);
+  EXPAND_VECTOR_TYPE(uint, unsigned int);
+  EXPAND_VECTOR_TYPE(long, long);
+  EXPAND_VECTOR_TYPE(ulong, unsigned long);
+  EXPAND_VECTOR_TYPE(longlong, long long);
+  EXPAND_VECTOR_TYPE(ulonglong, unsigned long long);
+  EXPAND_VECTOR_TYPE(float, float);
+  EXPAND_VECTOR_TYPE(double, double);
+
+  return true;
+}
+
+int main(int arg, char** argv) {
+    test();
+#if TEST_STD_VER >= 2014
+    static_assert(test(), "");
+#endif // TEST_STD_VER >= 2014
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_size.pass.cpp b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_size.pass.cpp
new file mode 100644
index 00000000000..87a692ee771
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/tuple/vector_types_tuple_size.pass.cpp
@@ -0,0 +1,54 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the libcu++ 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
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/cassert>
+#include <cuda/std/tuple>
+
+#include "test_macros.h"
+
+template <class VType, size_t Size>
+__host__ __device__ TEST_CONSTEXPR_CXX14 void test() {
+  assert(cuda::std::tuple_size<VType>::value == Size);
+  assert(cuda::std::tuple_size<const VType>::value == Size);
+  assert(cuda::std::tuple_size<volatile VType>::value == Size);
+  assert(cuda::std::tuple_size<const volatile VType>::value == Size);
+}
+
+#define EXPAND_VECTOR_TYPE(Type) \
+  test<Type##1, 1>();            \
+  test<Type##2, 2>();            \
+  test<Type##3, 3>();            \
+  test<Type##4, 4>();
+
+
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool test() {
+  EXPAND_VECTOR_TYPE(char);
+  EXPAND_VECTOR_TYPE(uchar);
+  EXPAND_VECTOR_TYPE(short);
+  EXPAND_VECTOR_TYPE(ushort);
+  EXPAND_VECTOR_TYPE(int);
+  EXPAND_VECTOR_TYPE(uint);
+  EXPAND_VECTOR_TYPE(long);
+  EXPAND_VECTOR_TYPE(ulong);
+  EXPAND_VECTOR_TYPE(longlong);
+  EXPAND_VECTOR_TYPE(ulonglong);
+  EXPAND_VECTOR_TYPE(float);
+  EXPAND_VECTOR_TYPE(double);
+
+  return true;
+}
+
+int main(int arg, char** argv) {
+    test();
+#if TEST_STD_VER >= 2014
+    static_assert(test(), "");
+#endif // TEST_STD_VER >= 2014
+
+  return 0;
+}