Use the float flavors of the cmath functions in the extended floating point fallbacks (#2106)

miscco · web-flow · commit bddcd2024022 · 2024-07-31T09:58:09.000+02:00
Fixes #2078
diff --git a/libcudacxx/include/cuda/std/__cuda/cmath_nvbf16.h b/libcudacxx/include/cuda/std/__cuda/cmath_nvbf16.h
@@ -37,47 +37,47 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 // trigonometric functions
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sin(__nv_bfloat16 __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __float2bfloat16(::sin(__bfloat162float(__v)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __float2bfloat16(::sinf(__bfloat162float(__v)));))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sinh(__nv_bfloat16 __v)
 {
-  return __float2bfloat16(::sinh(__bfloat162float(__v)));
+  return __float2bfloat16(::sinhf(__bfloat162float(__v)));
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cos(__nv_bfloat16 __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __float2bfloat16(::cos(__bfloat162float(__v)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __float2bfloat16(::cosf(__bfloat162float(__v)));))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cosh(__nv_bfloat16 __v)
 {
-  return __float2bfloat16(::cosh(__bfloat162float(__v)));
+  return __float2bfloat16(::coshf(__bfloat162float(__v)));
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 exp(__nv_bfloat16 __v)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __float2bfloat16(::exp(__bfloat162float(__v)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __float2bfloat16(::expf(__bfloat162float(__v)));))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 hypot(__nv_bfloat16 __x, __nv_bfloat16 __y)
 {
-  return __float2bfloat16(::hypot(__bfloat162float(__x), __bfloat162float(__y)));
+  return __float2bfloat16(::hypotf(__bfloat162float(__x), __bfloat162float(__y)));
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 atan2(__nv_bfloat16 __x, __nv_bfloat16 __y)
 {
-  return __float2bfloat16(::atan2(__bfloat162float(__x), __bfloat162float(__y)));
+  return __float2bfloat16(::atan2f(__bfloat162float(__x), __bfloat162float(__y)));
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 log(__nv_bfloat16 __x)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::log(__bfloat162float(__x)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::logf(__bfloat162float(__x)));))
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sqrt(__nv_bfloat16 __x)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrt(__bfloat162float(__x)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrtf(__bfloat162float(__x)));))
 }
 
 // floating point helper
diff --git a/libcudacxx/include/cuda/std/__cuda/cmath_nvfp16.h b/libcudacxx/include/cuda/std/__cuda/cmath_nvfp16.h
@@ -36,7 +36,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)
 {
   NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (return ::hsin(__v);), ({
                       float __vf            = __half2float(__v);
-                      __vf                  = ::sin(__vf);
+                      __vf                  = ::sinf(__vf);
                       __half_raw __ret_repr = ::__float2half_rn(__vf);
 
                       uint16_t __repr = __half_raw(__v).x;
@@ -61,7 +61,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half sinh(__half __v)
 {
-  return __float2half(::sinh(__half2float(__v)));
+  return __float2half(::sinhf(__half2float(__v)));
 }
 
 // clang-format off
@@ -72,7 +72,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)
   ), (
     {
       float __vf            = __half2float(__v);
-      __vf                  = ::cos(__vf);
+      __vf                  = ::cosf(__vf);
       __half_raw __ret_repr = ::__float2half_rn(__vf);
 
       uint16_t __repr = __half_raw(__v).x;
@@ -94,7 +94,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half cosh(__half __v)
 {
-  return __float2half(::cosh(__half2float(__v)));
+  return __float2half(::coshf(__half2float(__v)));
 }
 
 // clang-format off
@@ -105,7 +105,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)
   ), (
     {
       float __vf            = __half2float(__v);
-      __vf                  = ::exp(__vf);
+      __vf                  = ::expf(__vf);
       __half_raw __ret_repr = ::__float2half_rn(__vf);
 
       uint16_t __repr = __half_raw(__v).x;
@@ -127,12 +127,12 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half hypot(__half __x, __half __y)
 {
-  return __float2half(::hypot(__half2float(__x), __half2float(__y)));
+  return __float2half(::hypotf(__half2float(__x), __half2float(__y)));
 }
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half atan2(__half __x, __half __y)
 {
-  return __float2half(::atan2(__half2float(__x), __half2float(__y)));
+  return __float2half(::atan2f(__half2float(__x), __half2float(__y)));
 }
 
 // clang-format off
@@ -143,7 +143,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)
   ), (
     {
       float __vf            = __half2float(__x);
-      __vf                  = ::log(__vf);
+      __vf                  = ::logf(__vf);
       __half_raw __ret_repr = ::__float2half_rn(__vf);
 
       uint16_t __repr = __half_raw(__x).x;
@@ -164,7 +164,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)
 
 inline _LIBCUDACXX_INLINE_VISIBILITY __half sqrt(__half __x)
 {
-  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrt(__half2float(__x)));))
+  NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrtf(__half2float(__x)));))
 }
 
 // floating point helper

Original file line number	Diff line number	Diff line change
`@@ -37,47 +37,47 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD`
`37`	`37`	`// trigonometric functions`
`38`	`38`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sin(__nv_bfloat16 __v)`
`39`	`39`	`{`
`40`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __float2bfloat16(::sin(__bfloat162float(__v)));))`
	`40`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsin(__v);), (return __float2bfloat16(::sinf(__bfloat162float(__v)));))`
`41`	`41`	`}`
`42`	`42`
`43`	`43`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sinh(__nv_bfloat16 __v)`
`44`	`44`	`{`
`45`		`- return __float2bfloat16(::sinh(__bfloat162float(__v)));`
	`45`	`+ return __float2bfloat16(::sinhf(__bfloat162float(__v)));`
`46`	`46`	`}`
`47`	`47`
`48`	`48`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cos(__nv_bfloat16 __v)`
`49`	`49`	`{`
`50`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __float2bfloat16(::cos(__bfloat162float(__v)));))`
	`50`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hcos(__v);), (return __float2bfloat16(::cosf(__bfloat162float(__v)));))`
`51`	`51`	`}`
`52`	`52`
`53`	`53`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 cosh(__nv_bfloat16 __v)`
`54`	`54`	`{`
`55`		`- return __float2bfloat16(::cosh(__bfloat162float(__v)));`
	`55`	`+ return __float2bfloat16(::coshf(__bfloat162float(__v)));`
`56`	`56`	`}`
`57`	`57`
`58`	`58`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 exp(__nv_bfloat16 __v)`
`59`	`59`	`{`
`60`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __float2bfloat16(::exp(__bfloat162float(__v)));))`
	`60`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hexp(__v);), (return __float2bfloat16(::expf(__bfloat162float(__v)));))`
`61`	`61`	`}`
`62`	`62`
`63`	`63`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 hypot(__nv_bfloat16 __x, __nv_bfloat16 __y)`
`64`	`64`	`{`
`65`		`- return __float2bfloat16(::hypot(__bfloat162float(__x), __bfloat162float(__y)));`
	`65`	`+ return __float2bfloat16(::hypotf(__bfloat162float(__x), __bfloat162float(__y)));`
`66`	`66`	`}`
`67`	`67`
`68`	`68`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 atan2(__nv_bfloat16 __x, __nv_bfloat16 __y)`
`69`	`69`	`{`
`70`		`- return __float2bfloat16(::atan2(__bfloat162float(__x), __bfloat162float(__y)));`
	`70`	`+ return __float2bfloat16(::atan2f(__bfloat162float(__x), __bfloat162float(__y)));`
`71`	`71`	`}`
`72`	`72`
`73`	`73`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 log(__nv_bfloat16 __x)`
`74`	`74`	`{`
`75`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::log(__bfloat162float(__x)));))`
	`75`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::logf(__bfloat162float(__x)));))`
`76`	`76`	`}`
`77`	`77`
`78`	`78`	`inline _LIBCUDACXX_INLINE_VISIBILITY __nv_bfloat16 sqrt(__nv_bfloat16 __x)`
`79`	`79`	`{`
`80`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrt(__bfloat162float(__x)));))`
	`80`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrtf(__bfloat162float(__x)));))`
`81`	`81`	`}`
`82`	`82`
`83`	`83`	`// floating point helper`
Original file line number	Diff line number	Diff line change
`@@ -36,7 +36,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)`
`36`	`36`	`{`
`37`	`37`	`NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (return ::hsin(__v);), ({`
`38`	`38`	`float __vf = __half2float(__v);`
`39`		`- __vf = ::sin(__vf);`
	`39`	`+ __vf = ::sinf(__vf);`
`40`	`40`	`__half_raw __ret_repr = ::__float2half_rn(__vf);`
`41`	`41`
`42`	`42`	`uint16_t __repr = __half_raw(__v).x;`
`@@ -61,7 +61,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half sin(__half __v)`
`61`	`61`
`62`	`62`	`inline _LIBCUDACXX_INLINE_VISIBILITY __half sinh(__half __v)`
`63`	`63`	`{`
`64`		`- return __float2half(::sinh(__half2float(__v)));`
	`64`	`+ return __float2half(::sinhf(__half2float(__v)));`
`65`	`65`	`}`
`66`	`66`
`67`	`67`	`// clang-format off`
`@@ -72,7 +72,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)`
`72`	`72`	`), (`
`73`	`73`	`{`
`74`	`74`	`float __vf = __half2float(__v);`
`75`		`- __vf = ::cos(__vf);`
	`75`	`+ __vf = ::cosf(__vf);`
`76`	`76`	`__half_raw __ret_repr = ::__float2half_rn(__vf);`
`77`	`77`
`78`	`78`	`uint16_t __repr = __half_raw(__v).x;`
`@@ -94,7 +94,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half cos(__half __v)`
`94`	`94`
`95`	`95`	`inline _LIBCUDACXX_INLINE_VISIBILITY __half cosh(__half __v)`
`96`	`96`	`{`
`97`		`- return __float2half(::cosh(__half2float(__v)));`
	`97`	`+ return __float2half(::coshf(__half2float(__v)));`
`98`	`98`	`}`
`99`	`99`
`100`	`100`	`// clang-format off`
`@@ -105,7 +105,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)`
`105`	`105`	`), (`
`106`	`106`	`{`
`107`	`107`	`float __vf = __half2float(__v);`
`108`		`- __vf = ::exp(__vf);`
	`108`	`+ __vf = ::expf(__vf);`
`109`	`109`	`__half_raw __ret_repr = ::__float2half_rn(__vf);`
`110`	`110`
`111`	`111`	`uint16_t __repr = __half_raw(__v).x;`
`@@ -127,12 +127,12 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half exp(__half __v)`
`127`	`127`
`128`	`128`	`inline _LIBCUDACXX_INLINE_VISIBILITY __half hypot(__half __x, __half __y)`
`129`	`129`	`{`
`130`		`- return __float2half(::hypot(__half2float(__x), __half2float(__y)));`
	`130`	`+ return __float2half(::hypotf(__half2float(__x), __half2float(__y)));`
`131`	`131`	`}`
`132`	`132`
`133`	`133`	`inline _LIBCUDACXX_INLINE_VISIBILITY __half atan2(__half __x, __half __y)`
`134`	`134`	`{`
`135`		`- return __float2half(::atan2(__half2float(__x), __half2float(__y)));`
	`135`	`+ return __float2half(::atan2f(__half2float(__x), __half2float(__y)));`
`136`	`136`	`}`
`137`	`137`
`138`	`138`	`// clang-format off`
`@@ -143,7 +143,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)`
`143`	`143`	`), (`
`144`	`144`	`{`
`145`	`145`	`float __vf = __half2float(__x);`
`146`		`- __vf = ::log(__vf);`
	`146`	`+ __vf = ::logf(__vf);`
`147`	`147`	`__half_raw __ret_repr = ::__float2half_rn(__vf);`
`148`	`148`
`149`	`149`	`uint16_t __repr = __half_raw(__x).x;`
`@@ -164,7 +164,7 @@ inline _LIBCUDACXX_INLINE_VISIBILITY __half log(__half __x)`
`164`	`164`
`165`	`165`	`inline _LIBCUDACXX_INLINE_VISIBILITY __half sqrt(__half __x)`
`166`	`166`	`{`
`167`		`- NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrt(__half2float(__x)));))`
	`167`	`+ NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrtf(__half2float(__x)));))`
`168`	`168`	`}`
`169`	`169`
`170`	`170`	`// floating point helper`