From c29709b3444e644f9cf05b65d461145f5869ce05 Mon Sep 17 00:00:00 2001 From: Corey Lowman Date: Sat, 2 Dec 2023 13:57:02 -0500 Subject: [PATCH] Update dfdx-core/src/tensor_ops/utilities/compatibility.cuh --- dfdx-core/src/tensor_ops/utilities/compatibility.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/dfdx-core/src/tensor_ops/utilities/compatibility.cuh b/dfdx-core/src/tensor_ops/utilities/compatibility.cuh index 01d75427..48387eb9 100644 --- a/dfdx-core/src/tensor_ops/utilities/compatibility.cuh +++ b/dfdx-core/src/tensor_ops/utilities/compatibility.cuh @@ -172,8 +172,8 @@ __device__ __forceinline__ double atomicMinf(double * addr, double value) { #if __CUDA_ARCH__ < 530 //On older GPU there are no hcos/hsin/hlog/hexp. -__device__ __forceinline__ __half hcos(__half a) { return cosf(a); } -__device__ __forceinline__ __half hsin(__half a) { return sinf(a); } -__device__ __forceinline__ __half hlog(__half a) { return logf(a); } -__device__ __forceinline__ __half hexp(__half a) { return expf(a); } +__device__ __forceinline__ __half hcos(__half a) { return __float2half(cosf(__half2float(a))); } +__device__ __forceinline__ __half hsin(__half a) { return __float2half(sinf(__half2float(a))); } +__device__ __forceinline__ __half hlog(__half a) { return __float2half(logf(__half2float(a))); } +__device__ __forceinline__ __half hexp(__half a) { return __float2half(expf(__half2float(a))); } #endif \ No newline at end of file