summaryrefslogtreecommitdiff
path: root/candle-kernels/src
diff options
context:
space:
mode:
authorNicolas Patry <patry.nicolas@protonmail.com>2023-08-10 17:46:47 +0200
committerNicolas Patry <patry.nicolas@protonmail.com>2023-08-10 17:46:47 +0200
commit4a95d34c83453085c0d9dcecb19df9566024025e (patch)
tree52be70f6b66d2836b92557302f00c62fee9b8ec4 /candle-kernels/src
parentc8039579a5886f1df55a961b98fef3185a560b65 (diff)
downloadcandle-4a95d34c83453085c0d9dcecb19df9566024025e.tar.gz
candle-4a95d34c83453085c0d9dcecb19df9566024025e.tar.bz2
candle-4a95d34c83453085c0d9dcecb19df9566024025e.zip
Compat windows.
Diffstat (limited to 'candle-kernels/src')
-rw-r--r--candle-kernels/src/compatibility.cuh9
1 files changed, 9 insertions, 0 deletions
diff --git a/candle-kernels/src/compatibility.cuh b/candle-kernels/src/compatibility.cuh
index 5a22f4bc..d0791749 100644
--- a/candle-kernels/src/compatibility.cuh
+++ b/candle-kernels/src/compatibility.cuh
@@ -6,6 +6,15 @@
// FIXME: the minimum compute capabilities are just guesses since the table is not specific enough
+#if (__CUDACC_VER_MAJOR__ < 12 || __CUDACC_VER_MINOR__ < 2) && __CUDA_ARCH__ < 800
+__device__ __forceinline__ __half __hmax_nan(__half a, __half b) {
+ return __hisnan(a) ? a : (__hisnan(b) ? b : __hmax(a, b));
+}
+__device__ __forceinline__ __half __hmin_nan(__half a, __half b) {
+ return __hisnan(a) ? a : (__hisnan(b) ? b : __hmin(a, b));
+}
+#endif
+
#if __CUDA_ARCH__ < 600
// Copied from https://docs.nvidia.com/cuda/cuda-c-programming-guide/#atomic-functions
__device__ double atomicAdd(double* address, double val) {