Repository: coreylowman/dfdx (language: rust)

Issue:
Wrong CUDA_ARCH check in compatibility.cuh
I have a GPU with CUDA_ARCH=750, and I was getting `__hmax_nan undefined` errors. If I change this:
```
#if __CUDA_ARCH__ < 700
__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
```
to this:
```
#if __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
```
it works. I guess __hmax_nan doesn't exist at 750.


Discussion / hints:
The boundaries for when each function is defined are not well documented, so when adding these guards I just guessed. I think your proposed change is completely reasonable.
Maybe you can make a pull request?
Sure will do tomorrow
Yeah thanks for doing this, wish nvidia was more clear 😓 

Fix patch:
diff --git a/src/tensor_ops/utilities/compatibility.cuh b/src/tensor_ops/utilities/compatibility.cuh
index 93195a7f4..de11cae13 100644
--- a/src/tensor_ops/utilities/compatibility.cuh
+++ b/src/tensor_ops/utilities/compatibility.cuh
@@ -14,7 +14,7 @@ __device__ __forceinline__ __half __hmin(__half a, __half b) {
 }
 #endif
 
-#if __CUDA_ARCH__ < 700
+#if __CUDA_ARCH__ < 800
 __device__ __forceinline__ __half __hmax_nan(__half a, __half b) {
     return __hisnan(a) ? a : (__hisnan(b) ? b : __hmax(a, b));
 }


Test patch: