Add compatibility for hcos/hsin/hlog/hexp on older GPUs (#887)

* On older GPU there are no hcos/hsin/hlog/hexp * Update dfdx-core/src/tensor_ops/utilities/compatibility.cuh --------- Co-authored-by: Corey Lowman <[email protected]>
coreylowman · Dec 2, 2023 · 2ff7b60 · 2ff7b60
1 parent a522c7a
commit 2ff7b60
Showing 1 changed file with 9 additions and 1 deletion.
diff --git a/dfdx-core/src/tensor_ops/utilities/compatibility.cuh b/dfdx-core/src/tensor_ops/utilities/compatibility.cuh
@@ -168,4 +168,12 @@ __device__ __forceinline__ double atomicMinf(double * addr, double value) {
     } else {
         return __longlong_as_double(atomicMin((long long int *)addr, __double_as_longlong(value)));
     }
-}
+}
+
+#if __CUDA_ARCH__ < 530
+//On older GPU there are no hcos/hsin/hlog/hexp.
+__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