On older GPU there are no hcos/hsin/hlog/hexp

coreylowman · Nov 8, 2023 · 27c217c · 27c217c
1 parent 4476b5e
commit 27c217c
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 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); }
+#endif