Fix __half for older architectures (#1543)

Older architectures do not support all mathematical functions for `__half` For that we provide backfills through a conversion to `float` Fixes nvbug4563212
NVIDIA · Mar 18, 2024 · 4ba5eb4 · 4ba5eb4
1 parent 6af3b1e
commit 4ba5eb4
Show file tree

Hide file tree

Showing 2 changed files with 30 additions and 29 deletions.
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/cmath_nvbf16.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/cmath_nvbf16.h
@@ -50,7 +50,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sinh(__nv_bfloat16 __v)
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cos(__nv_bfloat16 __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return hcos(__v);), (return __nv_bfloat16(::cos(float(__v)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __nv_bfloat16(::cos(float(__v)));))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cosh(__nv_bfloat16 __v)

diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/cmath_nvfp16.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/cmath_nvfp16.h
@@ -36,29 +36,33 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 // trigonometric functions
 inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return hsin(__v);), ({
-                      float __vf            = __v;
-                      __vf                  = ::sin(__vf);
-                      __half_raw __ret_repr = ::__float2half_rn(__vf);
+  NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (
+    return ::hsin(__v);
+  ), (
+    {
+      float __vf            = __v;
+      __vf                  = ::sin(__vf);
+      __half_raw __ret_repr = ::__float2half_rn(__vf);
 
-                      uint16_t __repr = __half_raw(__v).x;
-                      switch (__repr)
-                      {
-                        case 12979:
-                        case 45747:
-                          __ret_repr.x += 1;
-                          break;
+      uint16_t __repr = __half_raw(__v).x;
+      switch (__repr)
+      {
+        case 12979:
+        case 45747:
+          __ret_repr.x += 1;
+          break;
 
-                        case 23728:
-                        case 56496:
-                          __ret_repr.x -= 1;
-                          break;
+        case 23728:
+        case 56496:
+          __ret_repr.x -= 1;
+          break;
 
-                        default:;
-                      }
+        default:;
+      }
 
-                      return __ret_repr;
-                    }))
+      return __ret_repr;
+  }
+))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half sinh(__half __v)
@@ -69,10 +73,9 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sinh(__half __v)
 // clang-format off
 inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE,
-  (
+  NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (
     return ::hcos(__v);
-  ),(
+  ), (
     {
       float __vf            = __v;
       __vf                  = ::cos(__vf);
@@ -103,10 +106,9 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cosh(__half __v)
 // clang-format off
 inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE,
-  (
+  NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (
     return ::hexp(__v);
-  ),(
+  ), (
     {
       float __vf            = __v;
       __vf                  = ::exp(__vf);
@@ -142,10 +144,9 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half atan2(__half __x, __half __y)
 // clang-format off
 inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE,
-  (
+  NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (
     return ::hlog(__x);
-  ),(
+  ), (
     {
       float __vf            = __x;
       __vf                  = ::log(__vf);