From 0771130fd67f552b799e2c6d7f232b39496a67a9 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 15 Oct 2024 09:26:53 -0700
Subject: [PATCH 1/2] [SYCL] Remove `generic_type_lists.hpp`

We don't need to instantiate all the `vec`/`marray` types to implement
`generic_type_traits.hpp`.
---
 .../sycl/detail/generic_type_lists.hpp        | 477 ------------------
 .../sycl/detail/generic_type_traits.hpp       |  60 ++-
 sycl/include/sycl/detail/type_traits.hpp      |   1 -
 .../detail/type_traits/vec_marray_traits.hpp  |   4 +
 .../sycl/ext/oneapi/bf16_storage_builtins.hpp |   1 -
 .../ext/oneapi/experimental/cuda/builtins.hpp |  36 +-
 sycl/include/sycl/types.hpp                   |   1 -
 sycl/include/sycl/vector.hpp                  |   3 +-
 .../DotProduct/dot_product_int_test.cpp       |   1 -
 .../DotProduct/dot_product_vec_test.cpp       |   1 -
 sycl/test/include_deps/sycl_accessor.hpp.cpp  |   1 -
 .../include_deps/sycl_detail_core.hpp.cpp     |   1 -
 12 files changed, 49 insertions(+), 538 deletions(-)
 delete mode 100644 sycl/include/sycl/detail/generic_type_lists.hpp

diff --git a/sycl/include/sycl/detail/generic_type_lists.hpp b/sycl/include/sycl/detail/generic_type_lists.hpp
deleted file mode 100644
index f212fd1df1c4f..0000000000000
--- a/sycl/include/sycl/detail/generic_type_lists.hpp
+++ /dev/null
@@ -1,477 +0,0 @@
-//==-------- generic_type_lists.hpp - SYCL Generic type lists --------------==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
-#pragma once
-
-#include <sycl/access/access.hpp>    // for address_space
-#include <sycl/detail/type_list.hpp> // for type_list, address_space_list
-
-#include <cstddef>     // for byte, size_t
-#include <type_traits> // for conditional_t, is_signed_v, is_...
-
-// Generic type name description, which serves as a description for all valid
-// types of parameters to kernel functions
-
-// Forward declarations
-namespace sycl {
-inline namespace _V1 {
-template <typename T, int N> class __SYCL_EBO vec;
-template <typename Type, std::size_t NumElements> class marray;
-
-namespace detail {
-namespace half_impl {
-class half;
-}
-} // namespace detail
-using half = detail::half_impl::half;
-
-namespace ext::oneapi {
-class bfloat16;
-}
-namespace detail {
-namespace gtl {
-// floating point types
-using scalar_half_list = type_list<half>;
-
-using vector_half_list = type_list<vec<half, 1>, vec<half, 2>, vec<half, 3>,
-                                   vec<half, 4>, vec<half, 8>, vec<half, 16>>;
-
-using marray_half_list =
-    type_list<marray<half, 1>, marray<half, 2>, marray<half, 3>,
-              marray<half, 4>, marray<half, 8>, marray<half, 16>>;
-
-using scalar_vector_half_list = tl_append<scalar_half_list, vector_half_list>;
-
-using half_list =
-    tl_append<scalar_half_list, vector_half_list, marray_half_list>;
-
-using scalar_bfloat16_list = type_list<sycl::ext::oneapi::bfloat16>;
-
-using vector_bfloat16_list = type_list<
-    vec<sycl::ext::oneapi::bfloat16, 1>, vec<sycl::ext::oneapi::bfloat16, 2>,
-    vec<sycl::ext::oneapi::bfloat16, 3>, vec<sycl::ext::oneapi::bfloat16, 4>,
-    vec<sycl::ext::oneapi::bfloat16, 8>, vec<sycl::ext::oneapi::bfloat16, 16>>;
-
-using marray_bfloat16_list = type_list<marray<sycl::ext::oneapi::bfloat16, 1>,
-                                       marray<sycl::ext::oneapi::bfloat16, 2>,
-                                       marray<sycl::ext::oneapi::bfloat16, 3>,
-                                       marray<sycl::ext::oneapi::bfloat16, 4>,
-                                       marray<sycl::ext::oneapi::bfloat16, 8>,
-                                       marray<sycl::ext::oneapi::bfloat16, 16>>;
-
-using scalar_vector_bfloat16_list =
-    tl_append<scalar_bfloat16_list, vector_bfloat16_list>;
-
-using bfloat16_list =
-    tl_append<scalar_bfloat16_list, vector_bfloat16_list, marray_bfloat16_list>;
-
-using scalar_float_list = type_list<float>;
-
-using vector_float_list =
-    type_list<vec<float, 1>, vec<float, 2>, vec<float, 3>, vec<float, 4>,
-              vec<float, 8>, vec<float, 16>>;
-
-using marray_float_list =
-    type_list<marray<float, 1>, marray<float, 2>, marray<float, 3>,
-              marray<float, 4>, marray<float, 8>, marray<float, 16>>;
-
-using scalar_vector_float_list =
-    tl_append<scalar_float_list, vector_float_list>;
-
-using float_list =
-    tl_append<scalar_float_list, vector_float_list, marray_float_list>;
-
-using scalar_double_list = type_list<double>;
-
-using vector_double_list =
-    type_list<vec<double, 1>, vec<double, 2>, vec<double, 3>, vec<double, 4>,
-              vec<double, 8>, vec<double, 16>>;
-
-using marray_double_list =
-    type_list<marray<double, 1>, marray<double, 2>, marray<double, 3>,
-              marray<double, 4>, marray<double, 8>, marray<double, 16>>;
-
-using scalar_vector_double_list =
-    tl_append<scalar_double_list, vector_double_list>;
-
-using double_list =
-    tl_append<scalar_double_list, vector_double_list, marray_double_list>;
-
-using scalar_floating_list = tl_append<scalar_float_list, scalar_double_list,
-                                       scalar_half_list, scalar_bfloat16_list>;
-
-using vector_floating_list = tl_append<vector_float_list, vector_double_list,
-                                       vector_half_list, vector_bfloat16_list>;
-
-using marray_floating_list = tl_append<marray_float_list, marray_double_list,
-                                       marray_half_list, marray_bfloat16_list>;
-
-using scalar_vector_floating_list =
-    tl_append<scalar_floating_list, vector_floating_list>;
-
-using floating_list =
-    tl_append<scalar_floating_list, vector_floating_list, marray_floating_list>;
-
-using scalar_default_char_list = type_list<char>;
-
-using vector_default_char_list =
-    type_list<vec<char, 1>, vec<char, 2>, vec<char, 3>, vec<char, 4>,
-              vec<char, 8>, vec<char, 16>>;
-
-using marray_default_char_list =
-    type_list<marray<char, 1>, marray<char, 2>, marray<char, 3>,
-              marray<char, 4>, marray<char, 8>, marray<char, 16>>;
-
-using default_char_list =
-    tl_append<scalar_default_char_list, vector_default_char_list,
-              marray_default_char_list>;
-
-using scalar_signed_char_list = type_list<signed char>;
-
-using vector_signed_char_list =
-    type_list<vec<signed char, 1>, vec<signed char, 2>, vec<signed char, 3>,
-              vec<signed char, 4>, vec<signed char, 8>, vec<signed char, 16>>;
-
-using marray_signed_char_list =
-    type_list<marray<signed char, 1>, marray<signed char, 2>,
-              marray<signed char, 3>, marray<signed char, 4>,
-              marray<signed char, 8>, marray<signed char, 16>>;
-
-using scalar_unsigned_char_list = type_list<unsigned char>;
-
-using vector_unsigned_char_list =
-    type_list<vec<unsigned char, 1>, vec<unsigned char, 2>,
-              vec<unsigned char, 3>, vec<unsigned char, 4>,
-              vec<unsigned char, 8>, vec<unsigned char, 16>>;
-
-using marray_unsigned_char_list =
-    type_list<marray<unsigned char, 1>, marray<unsigned char, 2>,
-              marray<unsigned char, 3>, marray<unsigned char, 4>,
-              marray<unsigned char, 8>, marray<unsigned char, 16>>;
-
-// short int types
-using scalar_signed_short_list = type_list<signed short>;
-
-using vector_signed_short_list =
-    type_list<vec<signed short, 1>, vec<signed short, 2>, vec<signed short, 3>,
-              vec<signed short, 4>, vec<signed short, 8>,
-              vec<signed short, 16>>;
-
-using marray_signed_short_list =
-    type_list<marray<signed short, 1>, marray<signed short, 2>,
-              marray<signed short, 3>, marray<signed short, 4>,
-              marray<signed short, 8>, marray<signed short, 16>>;
-
-using scalar_unsigned_short_list = type_list<unsigned short>;
-
-using vector_unsigned_short_list =
-    type_list<vec<unsigned short, 1>, vec<unsigned short, 2>,
-              vec<unsigned short, 3>, vec<unsigned short, 4>,
-              vec<unsigned short, 8>, vec<unsigned short, 16>>;
-
-using marray_unsigned_short_list =
-    type_list<marray<unsigned short, 1>, marray<unsigned short, 2>,
-              marray<unsigned short, 3>, marray<unsigned short, 4>,
-              marray<unsigned short, 8>, marray<unsigned short, 16>>;
-
-using unsigned_short_list =
-    tl_append<scalar_unsigned_short_list, vector_unsigned_short_list,
-              marray_unsigned_short_list>;
-
-using scalar_short_list =
-    tl_append<scalar_signed_short_list, scalar_unsigned_short_list>;
-
-using vector_short_list =
-    tl_append<vector_signed_short_list, vector_unsigned_short_list,
-              marray_unsigned_short_list>;
-
-using short_list = tl_append<scalar_short_list, vector_short_list>;
-
-// int types
-using scalar_signed_int_list = type_list<signed int>;
-
-using vector_signed_int_list =
-    type_list<vec<signed int, 1>, vec<signed int, 2>, vec<signed int, 3>,
-              vec<signed int, 4>, vec<signed int, 8>, vec<signed int, 16>>;
-
-using marray_signed_int_list =
-    type_list<marray<signed int, 1>, marray<signed int, 2>,
-              marray<signed int, 3>, marray<signed int, 4>,
-              marray<signed int, 8>, marray<signed int, 16>>;
-
-using signed_int_list =
-    tl_append<scalar_signed_int_list, vector_signed_int_list,
-              marray_signed_int_list>;
-
-using scalar_unsigned_int_list = type_list<unsigned int>;
-
-using vector_unsigned_int_list =
-    type_list<vec<unsigned int, 1>, vec<unsigned int, 2>, vec<unsigned int, 3>,
-              vec<unsigned int, 4>, vec<unsigned int, 8>,
-              vec<unsigned int, 16>>;
-
-using marray_unsigned_int_list =
-    type_list<marray<unsigned int, 1>, marray<unsigned int, 2>,
-              marray<unsigned int, 3>, marray<unsigned int, 4>,
-              marray<unsigned int, 8>, marray<unsigned int, 16>>;
-
-using unsigned_int_list =
-    tl_append<scalar_unsigned_int_list, vector_unsigned_int_list,
-              marray_unsigned_int_list>;
-
-using scalar_int_list =
-    tl_append<scalar_signed_int_list, scalar_unsigned_int_list>;
-
-using vector_int_list =
-    tl_append<vector_signed_int_list, vector_unsigned_int_list>;
-
-using marray_int_list =
-    tl_append<marray_signed_int_list, marray_unsigned_int_list>;
-
-using int_list = tl_append<scalar_int_list, vector_int_list, marray_int_list>;
-
-// long types
-using scalar_signed_long_list = type_list<signed long>;
-
-using vector_signed_long_list =
-    type_list<vec<signed long, 1>, vec<signed long, 2>, vec<signed long, 3>,
-              vec<signed long, 4>, vec<signed long, 8>, vec<signed long, 16>>;
-
-using marray_signed_long_list =
-    type_list<marray<signed long, 1>, marray<signed long, 2>,
-              marray<signed long, 3>, marray<signed long, 4>,
-              marray<signed long, 8>, marray<signed long, 16>>;
-
-using signed_long_list =
-    tl_append<scalar_signed_long_list, vector_signed_long_list,
-              marray_signed_long_list>;
-
-using scalar_unsigned_long_list = type_list<unsigned long>;
-
-using vector_unsigned_long_list =
-    type_list<vec<unsigned long, 1>, vec<unsigned long, 2>,
-              vec<unsigned long, 3>, vec<unsigned long, 4>,
-              vec<unsigned long, 8>, vec<unsigned long, 16>>;
-
-using marray_unsigned_long_list =
-    type_list<marray<unsigned long, 1>, marray<unsigned long, 2>,
-              marray<unsigned long, 3>, marray<unsigned long, 4>,
-              marray<unsigned long, 8>, marray<unsigned long, 16>>;
-
-using unsigned_long_list =
-    tl_append<scalar_unsigned_long_list, vector_unsigned_long_list,
-              marray_unsigned_long_list>;
-
-using scalar_long_list =
-    tl_append<scalar_signed_long_list, scalar_unsigned_long_list>;
-
-using vector_long_list =
-    tl_append<vector_signed_long_list, vector_unsigned_long_list>;
-
-using marray_long_list =
-    tl_append<marray_signed_long_list, marray_unsigned_long_list>;
-
-using long_list =
-    tl_append<scalar_long_list, vector_long_list, marray_long_list>;
-
-// long long types
-using scalar_signed_longlong_list = type_list<signed long long>;
-
-using vector_signed_longlong_list =
-    type_list<vec<signed long long, 1>, vec<signed long long, 2>,
-              vec<signed long long, 3>, vec<signed long long, 4>,
-              vec<signed long long, 8>, vec<signed long long, 16>>;
-
-using marray_signed_longlong_list =
-    type_list<marray<signed long long, 1>, marray<signed long long, 2>,
-              marray<signed long long, 3>, marray<signed long long, 4>,
-              marray<signed long long, 8>, marray<signed long long, 16>>;
-
-using signed_longlong_list =
-    tl_append<scalar_signed_longlong_list, vector_signed_longlong_list,
-              marray_signed_longlong_list>;
-
-using scalar_unsigned_longlong_list = type_list<unsigned long long>;
-
-using vector_unsigned_longlong_list =
-    type_list<vec<unsigned long long, 1>, vec<unsigned long long, 2>,
-              vec<unsigned long long, 3>, vec<unsigned long long, 4>,
-              vec<unsigned long long, 8>, vec<unsigned long long, 16>>;
-
-using marray_unsigned_longlong_list =
-    type_list<marray<unsigned long long, 1>, marray<unsigned long long, 2>,
-              marray<unsigned long long, 3>, marray<unsigned long long, 4>,
-              marray<unsigned long long, 8>, marray<unsigned long long, 16>>;
-
-using unsigned_longlong_list =
-    tl_append<scalar_unsigned_longlong_list, vector_unsigned_longlong_list,
-              marray_unsigned_longlong_list>;
-
-using scalar_longlong_list =
-    tl_append<scalar_signed_longlong_list, scalar_unsigned_longlong_list>;
-
-using vector_longlong_list =
-    tl_append<vector_signed_longlong_list, vector_unsigned_longlong_list>;
-
-using marray_longlong_list =
-    tl_append<marray_signed_longlong_list, marray_unsigned_longlong_list>;
-
-using longlong_list =
-    tl_append<scalar_longlong_list, vector_longlong_list, marray_longlong_list>;
-
-// long integer types
-using scalar_signed_long_integer_list =
-    tl_append<scalar_signed_long_list, scalar_signed_longlong_list>;
-
-using vector_signed_long_integer_list =
-    tl_append<vector_signed_long_list, vector_signed_longlong_list>;
-
-using marray_signed_long_integer_list =
-    tl_append<marray_signed_long_list, marray_signed_longlong_list>;
-
-using signed_long_integer_list =
-    tl_append<scalar_signed_long_integer_list, vector_signed_long_integer_list,
-              marray_signed_long_integer_list>;
-
-using scalar_unsigned_long_integer_list =
-    tl_append<scalar_unsigned_long_list, scalar_unsigned_longlong_list>;
-
-using vector_unsigned_long_integer_list =
-    tl_append<vector_unsigned_long_list, vector_unsigned_longlong_list>;
-
-using marray_unsigned_long_integer_list =
-    tl_append<marray_unsigned_long_list, marray_unsigned_longlong_list>;
-
-using unsigned_long_integer_list = tl_append<scalar_unsigned_long_integer_list,
-                                             vector_unsigned_long_integer_list,
-                                             marray_unsigned_long_integer_list>;
-
-using scalar_long_integer_list = tl_append<scalar_signed_long_integer_list,
-                                           scalar_unsigned_long_integer_list>;
-
-using vector_long_integer_list = tl_append<vector_signed_long_integer_list,
-                                           vector_unsigned_long_integer_list>;
-
-using marray_long_integer_list = tl_append<marray_signed_long_integer_list,
-                                           marray_unsigned_long_integer_list>;
-
-using long_integer_list =
-    tl_append<scalar_long_integer_list, vector_long_integer_list,
-              marray_long_integer_list>;
-
-#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-// std::byte
-using scalar_byte_list = type_list<std::byte>;
-
-using vector_byte_list =
-    type_list<vec<std::byte, 1>, vec<std::byte, 2>, vec<std::byte, 3>,
-              vec<std::byte, 4>, vec<std::byte, 8>, vec<std::byte, 16>>;
-
-using marray_byte_list = type_list<marray<std::byte, 1>, marray<std::byte, 2>,
-                                   marray<std::byte, 3>, marray<std::byte, 4>,
-                                   marray<std::byte, 8>, marray<std::byte, 16>>;
-#endif
-
-// integer types
-using scalar_signed_integer_list =
-    tl_append<std::conditional_t<
-                  std::is_signed_v<char>,
-                  tl_append<scalar_default_char_list, scalar_signed_char_list>,
-                  scalar_signed_char_list>,
-              scalar_signed_short_list, scalar_signed_int_list,
-              scalar_signed_long_list, scalar_signed_longlong_list>;
-
-using vector_signed_integer_list =
-    tl_append<std::conditional_t<
-                  std::is_signed_v<char>,
-                  tl_append<vector_default_char_list, vector_signed_char_list>,
-                  vector_signed_char_list>,
-              vector_signed_short_list, vector_signed_int_list,
-              vector_signed_long_list, vector_signed_longlong_list>;
-
-using marray_signed_integer_list =
-    tl_append<std::conditional_t<
-                  std::is_signed_v<char>,
-                  tl_append<marray_default_char_list, marray_signed_char_list>,
-                  marray_signed_char_list>,
-              marray_signed_short_list, marray_signed_int_list,
-              marray_signed_long_list, marray_signed_longlong_list>;
-
-using signed_integer_list =
-    tl_append<scalar_signed_integer_list, vector_signed_integer_list,
-              marray_signed_integer_list>;
-
-using scalar_unsigned_integer_list =
-    tl_append<std::conditional_t<std::is_unsigned_v<char>,
-                                 tl_append<scalar_default_char_list,
-                                           scalar_unsigned_char_list>,
-                                 scalar_unsigned_char_list>,
-              scalar_unsigned_short_list, scalar_unsigned_int_list,
-              scalar_unsigned_long_list, scalar_unsigned_longlong_list
-#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-              ,
-              scalar_byte_list
-#endif
-              >;
-
-using vector_unsigned_integer_list =
-    tl_append<std::conditional_t<std::is_unsigned_v<char>,
-                                 tl_append<vector_default_char_list,
-                                           vector_unsigned_char_list>,
-                                 vector_unsigned_char_list>,
-              vector_unsigned_short_list, vector_unsigned_int_list,
-              vector_unsigned_long_list, vector_unsigned_longlong_list
-#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-              ,
-              vector_byte_list
-#endif
-              >;
-
-using marray_unsigned_integer_list =
-    tl_append<std::conditional_t<std::is_unsigned_v<char>,
-                                 tl_append<marray_default_char_list,
-                                           marray_unsigned_char_list>,
-                                 marray_unsigned_char_list>,
-              marray_unsigned_short_list, marray_unsigned_int_list,
-              marray_unsigned_long_list, marray_unsigned_longlong_list
-#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-              ,
-              marray_byte_list
-#endif
-              >;
-
-using unsigned_integer_list =
-    tl_append<scalar_unsigned_integer_list, vector_unsigned_integer_list,
-              marray_unsigned_integer_list>;
-
-using scalar_integer_list =
-    tl_append<scalar_signed_integer_list, scalar_unsigned_integer_list>;
-
-using vector_integer_list =
-    tl_append<vector_signed_integer_list, vector_unsigned_integer_list>;
-
-using marray_integer_list =
-    tl_append<marray_signed_integer_list, marray_unsigned_integer_list>;
-
-using integer_list =
-    tl_append<scalar_integer_list, vector_integer_list, marray_integer_list>;
-
-// bool types
-
-using marray_bool_list =
-    type_list<marray<bool, 1>, marray<bool, 2>, marray<bool, 3>,
-              marray<bool, 4>, marray<bool, 8>, marray<bool, 16>>;
-
-using scalar_bool_list = type_list<bool>;
-
-using bool_list = tl_append<scalar_bool_list, marray_bool_list>;
-
-} // namespace gtl
-} // namespace detail
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp
index 4f05d84bf864e..6d00a39304d0c 100644
--- a/sycl/include/sycl/detail/generic_type_traits.hpp
+++ b/sycl/include/sycl/detail/generic_type_traits.hpp
@@ -10,7 +10,6 @@
 
 #include <sycl/access/access.hpp>             // for decorated, address_space
 #include <sycl/aliases.hpp>                   // for half, cl_char, cl_double
-#include <sycl/detail/generic_type_lists.hpp> // for nonconst_address_space...
 #include <sycl/detail/helpers.hpp>            // for marray
 #include <sycl/detail/type_list.hpp>          // for is_contained, find_sam...
 #include <sycl/detail/type_traits.hpp>        // for is_gen_based_on_type_s...
@@ -28,51 +27,66 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 template <typename T>
-inline constexpr bool is_svgenfloatf_v =
-    is_contained_v<T, gtl::scalar_vector_float_list>;
+using is_byte = typename
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+    std::is_same<T, std::byte>;
+#else
+    std::false_type;
+#endif
+
+template <typename T> inline constexpr bool is_byte_v = is_byte<T>::value;
 
 template <typename T>
-inline constexpr bool is_svgenfloath_v =
-    is_contained_v<T, gtl::scalar_vector_half_list>;
+inline constexpr bool is_svgenfloatf_v =
+    std::is_same_v<T, float> ||
+    (is_vec_v<T> && std::is_same_v<element_type_t<T>, float>);
 
 template <typename T>
-inline constexpr bool is_genfloat_v = is_contained_v<T, gtl::floating_list>;
+inline constexpr bool is_svgenfloath_v =
+    std::is_same_v<T, half> ||
+    (is_vec_v<T> && std::is_same_v<element_type_t<T>, half>);
 
 template <typename T>
 inline constexpr bool is_sgenfloat_v =
-    is_contained_v<T, gtl::scalar_floating_list>;
+    check_type_in_v<T, float, double, half, ext::oneapi::bfloat16>;
 
 template <typename T>
 inline constexpr bool is_vgenfloat_v =
-    is_contained_v<T, gtl::vector_floating_list>;
+    is_vec_v<T> && is_sgenfloat_v<element_type_t<T>>;
 
 template <typename T>
-inline constexpr bool is_geninteger_v = is_contained_v<T, gtl::integer_list>;
-
-template <typename T>
-inline constexpr bool is_sgeninteger_v =
-    is_contained_v<T, gtl::scalar_integer_list>;
+inline constexpr bool is_genfloat_v =
+    is_sgenfloat_v<T> || is_vgenfloat_v<T> ||
+    (is_marray_v<T> && is_sgenfloat_v<element_type_t<T>> &&
+     is_allowed_vec_size_v<num_elements_v<T>>);
 
 template <typename T>
 inline constexpr bool is_sigeninteger_v =
-    is_contained_v<T, gtl::scalar_signed_integer_list>;
+    check_type_in_v<T, signed char, short, int, long, long long> ||
+    (std::is_same_v<T, char> && std::is_signed_v<char>);
 
 template <typename T>
 inline constexpr bool is_sugeninteger_v =
-    is_contained_v<T, gtl::scalar_unsigned_integer_list>;
+    check_type_in_v<T, unsigned char, unsigned short, unsigned int,
+                    unsigned long, unsigned long long> ||
+    (std::is_same_v<T, char> && std::is_unsigned_v<char>) || is_byte_v<T>;
 
 template <typename T>
-inline constexpr bool is_genbool_v = is_contained_v<T, gtl::bool_list>;
+inline constexpr bool is_sgeninteger_v =
+    is_sigeninteger_v<T> || is_sugeninteger_v<T>;
 
 template <typename T>
-using is_byte = typename
-#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-    std::is_same<T, std::byte>;
-#else
-    std::false_type;
-#endif
+inline constexpr bool is_geninteger_v =
+    is_sgeninteger_v<T> ||
+    (is_vec_v<T> && is_sgeninteger_v<element_type_t<T>>) ||
+    (is_marray_v<T> && is_sgeninteger_v<element_type_t<T>> &&
+     is_allowed_vec_size_v<num_elements_v<T>>);
 
-template <typename T> inline constexpr bool is_byte_v = is_byte<T>::value;
+template <typename T>
+inline constexpr bool is_genbool_v =
+    std::is_same_v<T, bool> ||
+    (is_marray_v<T> && std::is_same_v<element_type_t<T>, bool> &&
+     is_allowed_vec_size_v<num_elements_v<T>>);
 
 template <int Size>
 using fixed_width_unsigned = std::conditional_t<
diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp
index 17eb9faba79c4..df5f9ea53a045 100644
--- a/sycl/include/sycl/detail/type_traits.hpp
+++ b/sycl/include/sycl/detail/type_traits.hpp
@@ -11,7 +11,6 @@
 #include <sycl/detail/type_traits/vec_marray_traits.hpp>
 
 #include <sycl/access/access.hpp>             // for decorated, address_space
-#include <sycl/detail/generic_type_lists.hpp> // for vec, marray, integer_list
 #include <sycl/detail/type_list.hpp>          // for is_contained, find_twi...
 
 #include <array>       // for array
diff --git a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
index cac13bc73de33..8097451352e75 100644
--- a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
+++ b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp
@@ -99,6 +99,10 @@ struct element_type<T __attribute__((ext_vector_type(N)))> {
 #endif
 template <typename T> using element_type_t = typename element_type<T>::type;
 
+template <int N>
+inline constexpr bool is_allowed_vec_size_v =
+    N == 1 || N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp b/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
index ee1bea39cae69..4352705693730 100644
--- a/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
+++ b/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
@@ -11,7 +11,6 @@
 #include <sycl/__spirv/spirv_ops.hpp>
 #include <sycl/builtins.hpp>
 #include <sycl/detail/builtins/builtins.hpp>
-#include <sycl/detail/generic_type_lists.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/detail/type_traits.hpp>
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
index 23b0eeeb814e2..c85bbcc84f721 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
@@ -25,37 +25,15 @@ namespace oneapi {
 namespace experimental {
 namespace cuda {
 
-namespace detail {
-using ldg_vector_types = sycl::detail::type_list<
-    sycl::vec<char, 2>, sycl::vec<char, 3>, sycl::vec<char, 4>,
-    sycl::vec<signed char, 2>, sycl::vec<signed char, 3>,
-    sycl::vec<signed char, 4>, sycl::vec<short, 2>, sycl::vec<short, 3>,
-    sycl::vec<short, 4>, sycl::vec<int, 2>, sycl::vec<int, 3>,
-    sycl::vec<int, 4>, sycl::vec<long, 2>, sycl::vec<long, 3>,
-    sycl::vec<long, 4>, sycl::vec<long long, 2>, sycl::vec<long long, 3>,
-    sycl::vec<long long, 4>, sycl::vec<unsigned char, 2>,
-    sycl::vec<unsigned char, 3>, sycl::vec<unsigned char, 4>,
-    sycl::vec<unsigned short, 2>, sycl::vec<unsigned short, 3>,
-    sycl::vec<unsigned short, 4>, sycl::vec<unsigned int, 2>,
-    sycl::vec<unsigned int, 3>, sycl::vec<unsigned int, 4>,
-    sycl::vec<unsigned long, 2>, sycl::vec<unsigned long, 3>,
-    sycl::vec<unsigned long, 4>, sycl::vec<unsigned long long, 2>,
-    sycl::vec<unsigned long long, 3>, sycl::vec<unsigned long long, 4>,
-    sycl::vec<half, 2>, sycl::vec<half, 3>, sycl::vec<half, 4>,
-    sycl::vec<float, 2>, sycl::vec<float, 3>, sycl::vec<float, 4>,
-    sycl::vec<double, 2>, sycl::vec<double, 3>, sycl::vec<double, 4>>;
-
-using ldg_types =
-    sycl::detail::tl_append<ldg_vector_types,
-                            sycl::detail::gtl::scalar_floating_list,
-                            sycl::detail::gtl::scalar_signed_integer_list,
-                            sycl::detail::gtl::scalar_unsigned_integer_list>;
-} // namespace detail
-
 template <typename T>
 inline __SYCL_ALWAYS_INLINE std::enable_if_t<
-    sycl::detail::is_contained<
-        T, sycl::ext::oneapi::experimental::cuda::detail::ldg_types>::value,
+    detail::check_type_in_v<detail::element_type_t<T>, char, signed char, short,
+                            int, long, long long, unsigned char, unsigned short,
+                            unsigned int, unsigned long, unsigned long long,
+                            half, float, double> &&
+        (std::is_same_v<detail::element_type_t<T>, T> ||
+         (detail::is_vec_v<T> && detail::num_elements_v<T> >= 2 &&
+          detail::num_elements_v<T> <= 4)),
     T>
 ldg(const T *ptr) {
 #if defined(__SYCL_DEVICE_ONLY__)
diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp
index 87698c835a3ee..599ba89819241 100644
--- a/sycl/include/sycl/types.hpp
+++ b/sycl/include/sycl/types.hpp
@@ -12,7 +12,6 @@
 #include <sycl/aliases.hpp>                    // for half, cl_char, cl_int
 #include <sycl/detail/common.hpp>              // for ArrayCreator, RepeatV...
 #include <sycl/detail/defines_elementary.hpp>  // for __SYCL2020_DEPRECATED
-#include <sycl/detail/generic_type_lists.hpp>  // for vector_basic_list
 #include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
 #include <sycl/detail/is_device_copyable.hpp>
 #include <sycl/detail/type_list.hpp>           // for is_contained
diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp
index baf8737f6e0d1..01e70f639e7b5 100644
--- a/sycl/include/sycl/vector.hpp
+++ b/sycl/include/sycl/vector.hpp
@@ -137,8 +137,7 @@ class __SYCL_EBO vec
   static_assert(std::is_same_v<DataT, std::remove_cv_t<DataT>>,
                 "DataT must be cv-unqualified");
 
-  static_assert(NumElements == 1 || NumElements == 2 || NumElements == 3 ||
-                    NumElements == 4 || NumElements == 8 || NumElements == 16,
+  static_assert(detail::is_allowed_vec_size_v<NumElements>,
                 "Invalid number of elements for sycl::vec: only 1, 2, 3, 4, 8 "
                 "or 16 are supported");
   static_assert(sizeof(bool) == sizeof(uint8_t), "bool size is not 1 byte");
diff --git a/sycl/test-e2e/DotProduct/dot_product_int_test.cpp b/sycl/test-e2e/DotProduct/dot_product_int_test.cpp
index 41749efd13055..0a0542149d866 100644
--- a/sycl/test-e2e/DotProduct/dot_product_int_test.cpp
+++ b/sycl/test-e2e/DotProduct/dot_product_int_test.cpp
@@ -15,7 +15,6 @@ static int testCount = 4;
 static int passCount;
 
 using namespace sycl;
-using namespace sycl::detail::gtl;
 using namespace sycl::ext::oneapi;
 
 constexpr int RangeLength = 100;
diff --git a/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp b/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp
index a83774d712e40..28ac47cc66993 100644
--- a/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp
+++ b/sycl/test-e2e/DotProduct/dot_product_vec_test.cpp
@@ -15,7 +15,6 @@ static int testCount = 4;
 static int passCount;
 
 using namespace sycl;
-using namespace sycl::detail::gtl;
 using namespace sycl::ext::oneapi;
 
 constexpr int RangeLength = 100;
diff --git a/sycl/test/include_deps/sycl_accessor.hpp.cpp b/sycl/test/include_deps/sycl_accessor.hpp.cpp
index fd5f7b86ad79a..622ea90da7006 100644
--- a/sycl/test/include_deps/sycl_accessor.hpp.cpp
+++ b/sycl/test/include_deps/sycl_accessor.hpp.cpp
@@ -25,7 +25,6 @@
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: detail/type_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: detail/generic_type_lists.hpp
 // CHECK-NEXT: detail/type_list.hpp
 // CHECK-NEXT: detail/boost/mp11/algorithm.hpp
 // CHECK-NEXT: detail/boost/mp11/list.hpp
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index 16d3cd5cfbbca..d7141579c0f48 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -26,7 +26,6 @@
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: detail/type_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
-// CHECK-NEXT: detail/generic_type_lists.hpp
 // CHECK-NEXT: detail/type_list.hpp
 // CHECK-NEXT: detail/boost/mp11/algorithm.hpp
 // CHECK-NEXT: detail/boost/mp11/list.hpp

From 097a0728c625b8c6388bc3bb87c655a59b544c4f Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Wed, 16 Oct 2024 08:45:35 -0700
Subject: [PATCH 2/2] Fix check-sycl

---
 sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
index c85bbcc84f721..067e238c2e36c 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/cuda/builtins.hpp
@@ -24,6 +24,9 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 namespace cuda {
+namespace detail {
+using namespace sycl::detail;
+}
 
 template <typename T>
 inline __SYCL_ALWAYS_INLINE std::enable_if_t<