CUDA kernels missing __hmin and __hmax

[opfromthestart]

This is weird since the exact opposite thing happened in #762, but it appears that now __hmin and __hmax are not being found. I tested driver versions 525 and 535 and neither worked to fix it.

error: failed to run custom build command for `dfdx-core v0.13.0 (https://github.com/coreylowman/dfdx.git#f4f7165d)`

Caused by:
  process didn't exit successfully: `/home/opfromthestart/rust/ai/stock-rnn/target/release/build/dfdx-core-0e67dc48f7c4f1de/build-script-build` (exit status: 101)
  --- stdout
  cargo:rerun-if-changed=build.rs
  cargo:rustc-cfg=feature="nightly"
  cargo:rustc-env=CUDA_INCLUDE_DIR=/usr/include
  cargo:rerun-if-changed=src/tensor_ops/utilities/binary_op_macros.cuh
  cargo:rerun-if-changed=src/tensor_ops/utilities/compatibility.cuh
  cargo:rerun-if-changed=src/tensor_ops/utilities/cuda_utils.cuh
  cargo:rerun-if-changed=src/tensor_ops/utilities/unary_op_macros.cuh
  cargo:rerun-if-env-changed=CUDA_COMPUTE_CAP
  cargo:rustc-env=CUDA_COMPUTE_CAP=sm_75
  cargo:rerun-if-changed=src/tensor_ops/abs/abs.cu
  cargo:rerun-if-changed=src/tensor_ops/accurate_gelu/accurate_gelu.cu
  cargo:rerun-if-changed=src/tensor_ops/adam/adam.cu
  cargo:rerun-if-changed=src/tensor_ops/add/binary_add.cu
  cargo:rerun-if-changed=src/tensor_ops/add/scalar_add.cu
  cargo:rerun-if-changed=src/tensor_ops/attention_reshape/attention_reshape.cu
  cargo:rerun-if-changed=src/tensor_ops/axpy/axpy.cu
  cargo:rerun-if-changed=src/tensor_ops/bce/bce.cu
  cargo:rerun-if-changed=src/tensor_ops/boolean/boolean.cu
  cargo:rerun-if-changed=src/tensor_ops/choose/choose.cu
  cargo:rerun-if-changed=src/tensor_ops/clamp/clamp.cu
  cargo:rerun-if-changed=src/tensor_ops/cmp/cmp.cu
  cargo:rerun-if-changed=src/tensor_ops/conv1d/conv1d.cu
  cargo:rerun-if-changed=src/tensor_ops/conv2d/conv2d.cu
  cargo:rerun-if-changed=src/tensor_ops/convtrans2d/convtrans2d.cu
  cargo:rerun-if-changed=src/tensor_ops/cos/cos.cu
  cargo:rerun-if-changed=src/tensor_ops/div/binary_div.cu
  cargo:rerun-if-changed=src/tensor_ops/div/scalar_div.cu
  cargo:rerun-if-changed=src/tensor_ops/dropout/dropout.cu
  cargo:rerun-if-changed=src/tensor_ops/exp/exp.cu
  cargo:rerun-if-changed=src/tensor_ops/fast_gelu/fast_gelu.cu
  cargo:rerun-if-changed=src/tensor_ops/huber_error/huber_error.cu
  cargo:rerun-if-changed=src/tensor_ops/ln/ln.cu
  cargo:rerun-if-changed=src/tensor_ops/max_to/max_to.cu
  cargo:rerun-if-changed=src/tensor_ops/maximum/maximum.cu
  cargo:rerun-if-changed=src/tensor_ops/min_to/min_to.cu
  cargo:rerun-if-changed=src/tensor_ops/minimum/minimum.cu
  cargo:rerun-if-changed=src/tensor_ops/mul/binary_mul.cu
  cargo:rerun-if-changed=src/tensor_ops/mul/scalar_mul.cu
  cargo:rerun-if-changed=src/tensor_ops/nans_to/nans_to.cu
  cargo:rerun-if-changed=src/tensor_ops/negate/negate.cu
  cargo:rerun-if-changed=src/tensor_ops/pool2d/pool2d.cu
  cargo:rerun-if-changed=src/tensor_ops/pow/pow.cu
  cargo:rerun-if-changed=src/tensor_ops/recip/recip.cu
  cargo:rerun-if-changed=src/tensor_ops/relu/relu.cu
  cargo:rerun-if-changed=src/tensor_ops/rmsprop/rmsprop.cu
  cargo:rerun-if-changed=src/tensor_ops/roll/roll.cu
  cargo:rerun-if-changed=src/tensor_ops/select_and_gather/gather.cu
  cargo:rerun-if-changed=src/tensor_ops/select_and_gather/select.cu
  cargo:rerun-if-changed=src/tensor_ops/sgd/sgd.cu
  cargo:rerun-if-changed=src/tensor_ops/sigmoid/sigmoid.cu
  cargo:rerun-if-changed=src/tensor_ops/sin/sin.cu
  cargo:rerun-if-changed=src/tensor_ops/slice/slice.cu
  cargo:rerun-if-changed=src/tensor_ops/sqrt/sqrt.cu
  cargo:rerun-if-changed=src/tensor_ops/square/square.cu
  cargo:rerun-if-changed=src/tensor_ops/sub/binary_sub.cu
  cargo:rerun-if-changed=src/tensor_ops/sub/scalar_sub.cu
  cargo:rerun-if-changed=src/tensor_ops/sum_to/sum_to.cu
  cargo:rerun-if-changed=src/tensor_ops/tanh/tanh.cu
  cargo:rerun-if-changed=src/tensor_ops/upscale2d/upscale2d.cu

  --- stderr
  thread 'main' panicked at /home/opfromthestart/.cargo/git/checkouts/dfdx-318e6e5ad83eea79/f4f7165/dfdx-core/build.rs:200:17:
  nvcc error while compiling "src/tensor_ops/abs/abs.cu":

  # stdout


  # stderr
  src/tensor_ops/utilities/compatibility.cuh(19): error: identifier "__hmax" is undefined

  src/tensor_ops/utilities/compatibility.cuh(19): error: ambiguous "?" operation: second operand of type "__half" can be converted to third operand type "<error-type>", and vice versa

  src/tensor_ops/utilities/compatibility.cuh(19): error: ambiguous "?" operation: second operand of type "__half" can be converted to third operand type "<error-type>", and vice versa

  src/tensor_ops/utilities/compatibility.cuh(22): error: identifier "__hmin" is undefined

  src/tensor_ops/utilities/compatibility.cuh(22): error: ambiguous "?" operation: second operand of type "__half" can be converted to third operand type "<error-type>", and vice versa

  src/tensor_ops/utilities/compatibility.cuh(22): error: ambiguous "?" operation: second operand of type "__half" can be converted to third operand type "<error-type>", and vice versa

  src/tensor_ops/abs/abs.cu(5): warning #177-D: variable "y" was declared but never referenced

  src/tensor_ops/abs/abs.cu(9): warning #177-D: variable "y" was declared but never referenced

  src/tensor_ops/abs/abs.cu(13): warning #177-D: variable "y" was declared but never referenced

  6 errors detected in the compilation of "src/tensor_ops/abs/abs.cu".

[coreylowman]

Yeah there are so many weird cases with these two I just can't figure it out lol