CUDA: Restore support for sm_50..sm_53 architectures libeigen/eigen!3445 Co-authored-by: Antonio Sanchez <cantonios@google.com>

diff --git a/Eigen/src/Core/arch/Default/Half.h b/Eigen/src/Core/arch/Default/Half.h
index 65cc133..509c84d 100644
--- a/Eigen/src/Core/arch/Default/Half.h
+++ b/Eigen/src/Core/arch/Default/Half.h

@@ -447,7 +447,7 @@
 
 namespace half_impl {
 
-#if defined(EIGEN_GPU_COMPILE_PHASE)
+#if (defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || defined(EIGEN_HIP_DEVICE_COMPILE)
 // Note: We deliberately do *not* define this to 1 even if we have Arm's native
 // fp16 type since GPU half types are rather different from native CPU half types.
 #define EIGEN_HAS_NATIVE_GPU_FP16
@@ -907,7 +907,7 @@
 
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool(isinf)(const half& a) { return (raw_half_as_uint16(a) & 0x7fff) == 0x7c00; }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool(isnan)(const half& a) {
-#if defined(EIGEN_GPU_COMPILE_PHASE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return __hisnan(a);
 #else
   return (raw_half_as_uint16(a) & 0x7fff) > 0x7c00;
@@ -925,14 +925,14 @@
 #endif
 }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp(const half& a) {
-#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hexp(::__half(a)));
 #else
   return half(numext::exp(float(a)));
 #endif
 }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp2(const half& a) {
-#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hexp2(::__half(a)));
 #else
   return half(numext::exp2(float(a)));
@@ -944,7 +944,7 @@
   return half(numext::ldexp(float(a), exponent));
 }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log(const half& a) {
-#if defined(EIGEN_GPU_COMPILE_PHASE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hlog(::__half(a)));
 #else
   return half(numext::log(float(a)));
@@ -955,7 +955,7 @@
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log2(const half& a) { return half(numext::log2(float(a))); }
 
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sqrt(const half& a) {
-#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hsqrt(::__half(a)));
 #else
   return half(numext::sqrt(float(a)));
@@ -979,14 +979,14 @@
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half atan(const half& a) { return half(numext::atan(float(a))); }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half atanh(const half& a) { return half(numext::atanh(float(a))); }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half floor(const half& a) {
-#if (defined(EIGEN_CUDA_ARCH)) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hfloor(::__half(a)));
 #else
   return half(numext::floor(float(a)));
 #endif
 }
 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half ceil(const half& a) {
-#if (defined(EIGEN_CUDA_ARCH)) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   return half(hceil(::__half(a)));
 #else
   return half(numext::ceil(float(a)));
@@ -1026,7 +1026,7 @@
 #if defined(EIGEN_HAS_NATIVE_GPU_FP16)
   // Rounded once in hardware.
   return __hfma(::__half(a), ::__half(b), ::__half(c));
-#elif defined(__SYCL_DEVICE_ONLY__)
+#elif defined(EIGEN_GPU_COMPILE_PHASE) || defined(__SYCL_DEVICE_ONLY__)
   // Rounds twice; see fma_round_to_odd.
   return half(numext::fma(static_cast<float>(a), static_cast<float>(b), static_cast<float>(c)));
 #elif defined(EIGEN_HAS_ARM64_FP16_SCALAR_ARITHMETIC)

diff --git a/Eigen/src/Core/arch/GPU/PacketMath.h b/Eigen/src/Core/arch/GPU/PacketMath.h
index 12dfd21..058bbbb 100644
--- a/Eigen/src/Core/arch/GPU/PacketMath.h
+++ b/Eigen/src/Core/arch/GPU/PacketMath.h

@@ -647,6 +647,10 @@
 // ---------------------------------------------------------------------------------------------------------------
 // half2, the native two-lane packet.
 
+#if (defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530) || defined(EIGEN_HIP_DEVICE_COMPILE)
+#define EIGEN_GPU_HAS_FP16_ARITHMETIC 1
+#endif
+
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) {
   return __half2half2(from);
@@ -754,7 +758,11 @@
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plset<half2>(const Eigen::half& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __halves2half2(a, __hadd(a, __float2half(1.0f)));
+#else
+  return __halves2half2(a, __float2half(__half2float(a) + 1.0f));
+#endif
 }
 
 template <>
@@ -802,32 +810,57 @@
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2& a, const half2& b) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hadd2(a, b);
+#else
+  return __floats2half2_rn(__low2float(a) + __low2float(b), __high2float(a) + __high2float(b));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub<half2>(const half2& a, const half2& b) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hsub2(a, b);
+#else
+  return __floats2half2_rn(__low2float(a) - __low2float(b), __high2float(a) - __high2float(b));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hneg2(a);
+#else
+  return __floats2half2_rn(-__low2float(a), -__high2float(a));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2& a, const half2& b) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hmul2(a, b);
+#else
+  return __floats2half2_rn(__low2float(a) * __low2float(b), __high2float(a) * __high2float(b));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd<half2>(const half2& a, const half2& b, const half2& c) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hfma2(a, b, c);
+#else
+  return __floats2half2_rn(fmaf(__low2float(a), __low2float(b), __low2float(c)),
+                           fmaf(__high2float(a), __high2float(b), __high2float(c)));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2& a, const half2& b) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __h2div(a, b);
+#else
+  return __floats2half2_rn(__low2float(a) / __low2float(b), __high2float(a) / __high2float(b));
+#endif
 }
 
 // Compared through float, and every comparison against a NaN is false, so the result is b. This differs from
@@ -849,7 +882,11 @@
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hadd(__low2half(a), __high2half(a));
+#else
+  return __float2half(__low2float(a) + __high2float(a));
+#endif
 }
 
 template <>
@@ -859,41 +896,69 @@
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   const __half first = __low2half(a);
   const __half second = __high2half(a);
   return __hgt(first, second) ? first : second;
+#else
+  return __low2float(a) > __high2float(a) ? __low2half(a) : __high2half(a);
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   const __half first = __low2half(a);
   const __half second = __high2half(a);
   return __hlt(first, second) ? first : second;
+#else
+  return __low2float(a) < __high2float(a) ? __low2half(a) : __high2half(a);
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return __hmul(__low2half(a), __high2half(a));
+#else
+  return __float2half(__low2float(a) * __high2float(a));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return h2log(a);
+#else
+  return __floats2half2_rn(logf(__low2float(a)), logf(__high2float(a)));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return h2exp(a);
+#else
+  return __floats2half2_rn(expf(__low2float(a)), expf(__high2float(a)));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return h2sqrt(a);
+#else
+  return __floats2half2_rn(sqrtf(__low2float(a)), sqrtf(__high2float(a)));
+#endif
 }
 
 template <>
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt<half2>(const half2& a) {
+#if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
   return h2rsqrt(a);
+#else
+  return __floats2half2_rn(rsqrtf(__low2float(a)), rsqrtf(__high2float(a)));
+#endif
 }
 
 // No native h2log1p/h2expm1; both go through float.
@@ -1050,8 +1115,9 @@
 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 plset<Packet4h2>(const Eigen::half& a) {
   // Add each offset once: (a + 2*k) + 1 can round twice. Preserve a's sign in lane zero.
   const half2 base = pset1<half2>(a);
-  return make_packet4h2(plset<half2>(a), __hadd2(base, __floats2half2_rn(2.0f, 3.0f)),
-                        __hadd2(base, __floats2half2_rn(4.0f, 5.0f)), __hadd2(base, __floats2half2_rn(6.0f, 7.0f)));
+  return make_packet4h2(plset<half2>(a), padd<half2>(base, __floats2half2_rn(2.0f, 3.0f)),
+                        padd<half2>(base, __floats2half2_rn(4.0f, 5.0f)),
+                        padd<half2>(base, __floats2half2_rn(6.0f, 7.0f)));
 }
 
 template <>
@@ -1131,6 +1197,8 @@
   return predux_mul<half2>(product);
 }
 
+#undef EIGEN_GPU_HAS_FP16_ARITHMETIC
+
 #endif  // defined(EIGEN_GPU_COMPILE_PHASE)
 
 }  // end namespace internal

diff --git a/Eigen/src/Core/util/Macros.h b/Eigen/src/Core/util/Macros.h
index d71162e..ec30a27 100644
--- a/Eigen/src/Core/util/Macros.h
+++ b/Eigen/src/Core/util/Macros.h

@@ -585,12 +585,10 @@
 #error "Eigen requires CUDA 11.8 or later."
 #endif
 
-// Native FP16 packet math intrinsics (e.g. __hfma2, h2exp, h2log) are only
-// declared by the CUDA headers when __CUDA_ARCH__ >= 530. Guard the device
-// pass with a clear error rather than surfacing as "identifier `__hfma2` is
-// undefined" deep inside PacketMath.h.
-#if defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH < 600
-#error "Eigen requires CUDA compute capability >= 6.0 (sm_60). Compile with -arch=sm_60 or higher."
+// Guard the CUDA device pass with a clear error when targeting architectures
+// older than Maxwell (sm_50).
+#if defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH < 500
+#error "Eigen requires CUDA compute capability >= 5.0 (sm_50). Compile with -arch=sm_50 or higher."
 #endif
 
 #if defined(__HIPCC__) && !defined(EIGEN_NO_HIP) && !defined(__SYCL_DEVICE_ONLY__)

diff --git a/ci/build.linux.gitlab-ci.yml b/ci/build.linux.gitlab-ci.yml
index dcff393..c5f5cb6 100644
--- a/ci/build.linux.gitlab-ci.yml
+++ b/ci/build.linux.gitlab-ci.yml

@@ -409,9 +409,9 @@
     EIGEN_CI_CUDA_COMPUTE_ARCH: "75"
     EIGEN_CI_CUDA_CXX_FLAGS: "-Wno-unknown-cuda-version"
 
-# ---- Latest CUDA + newest compilers, sm_89;sm_90 ---------------------------
+# ---- CUDA 12.6 + newest compilers, sm_50;sm_53;sm_89;sm_90 -----------------
 # These jobs run natively on the L4 at sm_89 and retain compile-only coverage
-# for sm_90.
+# for sm_50, sm_53, and sm_90 (CUDA 13 dropped sm_50..sm_70).
 
 # GCC-13, CUDA-12.6 (nvcc)
 # CUDA 12.6 nvcc only supports gcc up to 13.
@@ -421,7 +421,7 @@
   variables:
     EIGEN_CI_C_COMPILER: gcc-13
     EIGEN_CI_CXX_COMPILER: g++-13
-    EIGEN_CI_CUDA_COMPUTE_ARCH: "89;90"
+    EIGEN_CI_CUDA_COMPUTE_ARCH: "50;53;89;90"
 
 # Clang-19, CUDA-12.6 (clang as CUDA compiler)
 # Clang CUDA support lags behind nvcc; clang-19 supports up to CUDA 12.x.
@@ -433,7 +433,7 @@
     EIGEN_CI_C_COMPILER: clang-19
     EIGEN_CI_CXX_COMPILER: clang++-19
     EIGEN_CI_TEST_CUDA_CLANG: "on"
-    EIGEN_CI_CUDA_COMPUTE_ARCH: "89;90"
+    EIGEN_CI_CUDA_COMPUTE_ARCH: "50;53;89;90"
 
 # ---- CUDA 13 ----------------------------------------------------------------
 # CUDA 13 dropped offline compilation for sm_70, relocated the CCCL headers and

diff --git a/contrib/Eigen/src/Tensor/TensorReductionGpu.h b/contrib/Eigen/src/Tensor/TensorReductionGpu.h
index cc4bb38..1dc0bf9 100644
--- a/contrib/Eigen/src/Tensor/TensorReductionGpu.h
+++ b/contrib/Eigen/src/Tensor/TensorReductionGpu.h

@@ -156,10 +156,12 @@
   atomicAdd(output, accum);
 }
 
+#if !defined(EIGEN_CUDA_ARCH) || EIGEN_CUDA_ARCH >= 600
 template <>
 __device__ inline void atomicReduce(double* output, double accum, SumReducer<double>&) {
   atomicAdd(output, accum);
 }
+#endif
 
 template <typename CoeffType, typename Index>
 __global__ EIGEN_HIP_LAUNCH_BOUNDS_1024 void ReductionInitKernel(const CoeffType val, Index num_preserved_coeffs,