From 27c217cb377cd15d1a3aedfffbd514c7002db0d5 Mon Sep 17 00:00:00 2001 From: optman Date: Wed, 8 Nov 2023 15:37:19 +0800 Subject: [PATCH] On older GPU there are no hcos/hsin/hlog/hexp --- dfdx-core/src/tensor_ops/utilities/compatibility.cuh | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/dfdx-core/src/tensor_ops/utilities/compatibility.cuh b/dfdx-core/src/tensor_ops/utilities/compatibility.cuh index 117c1c49..01d75427 100644 --- 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))); } -} \ No newline at end of file +} + +#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 \ No newline at end of file