Fix FP16 traits

NVIDIA · Jan 27, 2025 · 23c7c18 · 23c7c18
1 parent c7e64a7
commit 23c7c18
Show file tree

Hide file tree

Showing 2 changed files with 14 additions and 4 deletions.
diff --git a/c2h/include/c2h/bfloat16.cuh b/c2h/include/c2h/bfloat16.cuh
@@ -251,12 +251,17 @@ class __numeric_limits_impl<bfloat16_t, __numeric_limits_type::__floating_point>
 public:
   static __host__ __device__ __forceinline__ bfloat16_t max()
   {
-    return bfloat16_t::max();
+    return bfloat16_t(numeric_limits<__nv_bfloat16>::max());
+  }
+
+  static __host__ __device__ __forceinline__ bfloat16_t min()
+  {
+    return bfloat16_t(numeric_limits<__nv_bfloat16>::min());
   }
 
   static __host__ __device__ __forceinline__ bfloat16_t lowest()
   {
-    return bfloat16_t::lowest();
+    return bfloat16_t(numeric_limits<__nv_bfloat16>::lowest());
   }
 };
 _LIBCUDACXX_END_NAMESPACE_STD

diff --git a/c2h/include/c2h/half.cuh b/c2h/include/c2h/half.cuh
@@ -347,12 +347,17 @@ class __numeric_limits_impl<half_t, __numeric_limits_type::__floating_point>
 public:
   static __host__ __device__ __forceinline__ half_t max()
   {
-    return (half_t::max)();
+    return half_t(numeric_limits<__half>::max());
+  }
+
+  static __host__ __device__ __forceinline__ half_t min()
+  {
+    return half_t(numeric_limits<__half>::min());
   }
 
   static __host__ __device__ __forceinline__ half_t lowest()
   {
-    return half_t::lowest();
+    return half_t(numeric_limits<__half>::lowest());
   }
 };
 _LIBCUDACXX_END_NAMESPACE_STD