diff --git a/include/camp/defines.hpp b/include/camp/defines.hpp
index c37e3bc..dea5d69 100644
--- a/include/camp/defines.hpp
+++ b/include/camp/defines.hpp
@@ -98,10 +98,39 @@ namespace camp
 #define CAMP_HAVE_OMP_OFFLOAD 1
 #endif
 
-#if defined(__has_builtin)
-#if __has_builtin(__make_integer_seq)
+// This works for clang, nvcc 10 and higher using clang as a host compiler
+#define CAMP_USE_MAKE_INTEGER_SEQ 0
+#define CAMP_USE_TYPE_PACK_ELEMENT 0
+#if defined(_MSC_FULL_VER) && _MSC_FULL_VER >= 191125507
+  // __has_builtin exists but does not always expose this
+#undef CAMP_USE_MAKE_INTEGER_SEQ
 #define CAMP_USE_MAKE_INTEGER_SEQ 1
+#undef CAMP_USE_TYPE_PACK_ELEMENT
+#define CAMP_USE_TYPE_PACK_ELEMENT 1
+#elif defined(__has_builtin)
+#if __has_builtin(__make_integer_seq) && ((!defined(__NVCC__) || __CUDACC_VER_MAJOR >= 10))
+#undef CAMP_USE_MAKE_INTEGER_SEQ
+#define CAMP_USE_MAKE_INTEGER_SEQ 1
+#undef CAMP_USE_TYPE_PACK_ELEMENT
+#define CAMP_USE_TYPE_PACK_ELEMENT 1
+#endif
 #endif
+
+// This works for:
+//   GCC >= 8
+//   intel 19+ in GCC 8 or higher mode
+//   nvcc 10+ in GCC 8 or higher mode
+//   PGI 19+ in GCC 8 or higher mode
+#if __GNUC__ >= 8 && (\
+    /* intel compiler in gcc 8+ mode */ \
+    ((!defined(__INTEL_COMPILER)) || __INTEL_COMPILER >= 1900) \
+    /* nvcc in gcc 8+ mode */ \
+  ||((!defined(__NVCC__)) || __CUDACC_VER_MAJOR >= 10) \
+  ||((!defined(__PGIC__)) || __PGIC__ >= 19) \
+    )
+#define CAMP_USE_INTEGER_PACK 1
+#else
+#define CAMP_USE_INTEGER_PACK 0
 #endif
 
 // libstdc++ from GCC below version 5 lacks the type trait
diff --git a/include/camp/list/at.hpp b/include/camp/list/at.hpp
index 83aaf34..8374f3d 100644
--- a/include/camp/list/at.hpp
+++ b/include/camp/list/at.hpp
@@ -22,6 +22,15 @@ namespace camp
 
 namespace detail
 {
+  template <typename T, idx_t Idx>
+  struct _at;
+
+  #if CAMP_USE_TYPE_PACK_ELEMENT
+  template <idx_t Idx, template <class...> class T, typename... Pack>
+  struct _at<T<Pack...>, Idx> {
+    using type = __type_pack_element<Idx, Pack...>;
+  };
+  #else
   // Lookup from metal::at machinery
   template <idx_t, typename>
   struct entry {
@@ -45,8 +54,6 @@ namespace detail
       : decltype(_lookup_impl<Idx>(declptr<entries<indices, vals>>())) {
   };
 
-  template <typename T, idx_t Idx>
-  struct _at;
   template <template <class...> class T, typename X, typename... Rest>
   struct _at<T<X, Rest...>, 0> {
     using type = X;
@@ -65,6 +72,7 @@ namespace detail
                                   make_idx_seq_t<sizeof...(Rest)>,
                                   Idx>::type;
   };
+  #endif
 }  // namespace detail
 
 // TODO: document
diff --git a/include/camp/number.hpp b/include/camp/number.hpp
index dbda392..0ec25d8 100644
--- a/include/camp/number.hpp
+++ b/include/camp/number.hpp
@@ -36,11 +36,16 @@ namespace detail
 {
   template <typename T, typename N>
   struct gen_seq;
-#if defined(CAMP_USE_MAKE_INTEGER_SEQ) && !__NVCC__
+#if CAMP_USE_MAKE_INTEGER_SEQ
   template <typename T, T N>
   struct gen_seq<T, integral_constant<T, N>> {
     using type = __make_integer_seq<int_seq, T, N>;
   };
+#elif CAMP_USE_INTEGER_PACK
+  template <typename T, T N>
+  struct gen_seq<T, integral_constant<T, N>> {
+    using type = int_seq<T, __integer_pack(N)...>;
+  };
 #else
   template <typename T, typename S1, typename S2>
   struct concat;