#832 - Unable to build on Windows with VS and CUDA 12.6

[misureaudio]

Minimum reproducible example

cargo build --release --features cuda

Error

Build error candle 0.7.2

error trace:

error: failed to run custom build command for candle-kernels v0.7.2 (https://github.com/EricLBuehler/candle.git?rev=60eb251#60eb251f)
note: To improve backtraces for build dependencies, set the CARGO_PROFILE_RELEASE_BUILD_OVERRIDE_DEBUG=true environment variable to enable debug information generation.

Caused by:
process didn't exit successfully: C:\Users\MATTIA\Desktop\rustsrc\mistral.rs.0.3.1.0832\target\release\build\candle-kernels-89221a871de49b16\build-script-build (exit code: 101)
--- stdout
cargo:rerun-if-changed=build.rs
cargo:rerun-if-changed=src/compatibility.cuh
cargo:rerun-if-changed=src/cuda_utils.cuh
cargo:rerun-if-changed=src/binary_op_macros.cuh
cargo:info=["/usr", "/usr/local/cuda", "/opt/cuda", "/usr/lib/cuda", "C:/Program Files/NVIDIA GPU Computing Toolkit", "C:/CUDA"]
cargo:rerun-if-env-changed=CUDA_COMPUTE_CAP
cargo:rustc-env=CUDA_COMPUTE_CAP=89
cargo:info=Builder { cuda_root: Some("C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.6"), kernel_paths: ["src\affine.cu", "src\binary.cu", "src\cast.cu", "src\conv.cu", "src\fill.cu", "src\fused_rms_norm.cu", "src\fused_rope.cu", "src\indexing.cu", "src\kvconcat.cu", "src\quantized.cu", "src\reduce.cu", "src\sort.cu", "src\ternary.cu", "src\unary.cu"], watch: [], include_paths: ["src\binary_op_macros.cuh", "src\compatibility.cuh", "src\cuda_utils.cuh"], compute_cap: Some(89), out_dir: "C:\Users\MATTIA\Desktop\rustsrc\mistral.rs.0.3.1.0832\target\release\build\candle-kernels-1e018c63da86e142\out", extra_args: [] }
cargo:rustc-env=CUDA_INCLUDE_DIR=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.6\include
cargo:rerun-if-changed=src\binary_op_macros.cuh
cargo:rerun-if-changed=src\compatibility.cuh
cargo:rerun-if-changed=src\cuda_utils.cuh
cargo:rerun-if-env-changed=NVCC_CCBIN
cargo:rerun-if-changed=src\affine.cu
cargo:rerun-if-changed=src\binary.cu
cargo:rerun-if-changed=src\indexing.cu
cargo:rerun-if-changed=src\kvconcat.cu
cargo:rerun-if-changed=src\conv.cu
cargo:rerun-if-changed=src\reduce.cu
cargo:rerun-if-changed=src\fused_rms_norm.cu
cargo:rerun-if-changed=src\quantized.cu
cargo:rerun-if-changed=src\cast.cu
cargo:rerun-if-changed=src\fused_rope.cu
cargo:rerun-if-changed=src\sort.cu
cargo:rerun-if-changed=src\fill.cu
cargo:rerun-if-changed=src\unary.cu
cargo:rerun-if-changed=src\ternary.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(371): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_argmin(size_t, size_t, size_t, const size_t *, const T *, uint32_t *) [with T=uint32_t]" at line 616

Remark: The warnings can be suppressed with "-diag-suppress "

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(331): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_min(size_t, size_t, size_t, const size_t *, const T *, T *) [with T=uint32_t]" at line 616

1 error detected in the compilation of "src/indexing.cu".
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(371): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_argmin(size_t, size_t, size_t, const size_t *, const T *, uint32_t *) [with T=int16_t]" at line 617

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(331): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_min(size_t, size_t, size_t, const size_t *, const T *, T *) [with T=int16_t]" at line 617

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(371): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_argmin(size_t, size_t, size_t, const size_t *, const T *, uint32_t *) [with T=int32_t]" at line 618

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(331): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_min(size_t, size_t, size_t, const size_t *, const T *, T *) [with T=int32_t]" at line 618

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(371): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_argmin(size_t, size_t, size_t, const size_t *, const T *, uint32_t *) [with T=int64_t]" at line 619

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(331): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_min(size_t, size_t, size_t, const size_t *, const T *, T *) [with T=int64_t]" at line 619

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(371): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_argmin(size_t, size_t, size_t, const size_t *, const T *, uint32_t *) [with T=uint8_t]" at line 620

C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\reduce.cu(331): warning #173-D: floating-point value does not fit in required integral type
shr[tid] = ((float)(1e+300 * 1e+300));
^
detected during instantiation of "void fast_min(size_t, size_t, size_t, const size_t *, const T *, T *) [with T=uint8_t]" at line 620

1 error detected in the compilation of "src/reduce.cu".
indexing.cu
reduce.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/cast.cu".
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

cast.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/unary.cu".
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/binary.cu".
1 error detected in the compilation of "src/conv.cu".
unary.cu
binary.cu
conv.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/kvconcat.cu".
kvconcat.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/sort.cu".
sort.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/ternary.cu".
ternary.cu
C:\Users\MATTIA.cargo\git\checkouts\candle-c6a149c3b35a488f\60eb251\candle-kernels\src\cuda_utils.cuh(238): error: identifier "isnanf" is undefined
__declspec(device) __forceinline bool isnang(__nv_fp8_e4m3 a) { return isnanf(__half2float(__nv_cvt_fp8_to_halfraw(a.__x, __NV_E4M3))); }
^

1 error detected in the compilation of "src/affine.cu".
affine.cu

--- stderr
thread 'main' panicked at C:\Users\MATTIA.cargo\registry\src\index.crates.io-6f17d22bba15001f\bindgen_cuda-0.1.5\src\lib.rs:391:13:
nvcc error while compiling "src\affine.cu":

CLI "nvcc" "--gpu-architecture=sm_89" "--ptx" "--default-stream" "per-thread" "--output-directory" "C:\Users\MATTIA\Desktop\rustsrc\mistral.rs.0.3.1.0832\target\release\build\candle-kernels-1e018c63da86e142\out" "-Isrc" "-IC:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.6\include" "src\affine.cu"

stdout

stderr

stack backtrace:
0: std::panicking::begin_panic_handler
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\panicking.rs:665
1: core::panicking::panic_fmt
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\core\src\panicking.rs:74
2: bindgen_cuda::Builder::build_ptx
3: std::rt::lang_start
4: std::rt::lang_start
5: std::rt::lang_start
6: std::rt::lang_start
7: std::rt::lang_start_internal::closure$2
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\rt.rs:141
8: std::panicking::try::do_call
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\panicking.rs:557
9: std::panicking::try
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\panicking.rs:521
10: std::panic::catch_unwind
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\panic.rs:350
11: std::rt::lang_start_internal
at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library\std\src\rt.rs:141
12: std::rt::lang_start
13: main
14: invoke_main
at D:\a_work\1\s\src\vctools\crt\vcstartup\src\startup\exe_common.inl:78
15: __scrt_common_main_seh
at D:\a_work\1\s\src\vctools\crt\vcstartup\src\startup\exe_common.inl:288
16: BaseThreadInitThunk
17: RtlUserThreadStart
note: Some details are omitted, run with RUST_BACKTRACE=full for a verbose backtrace.
warning: build failed, waiting for other jobs to finish...

Other information

Please specify:
Windows 11, VS 2022 17.11.4 - rustc 1.81.0

• GPU or accelerator information
  nvcc: NVIDIA (R) Cuda compiler driver
  Copyright (c) 2005-2024 NVIDIA Corporation
  Built on Wed_Aug_14_10:26:51_Pacific_Daylight_Time_2024
  Cuda compilation tools, release 12.6, V12.6.68
  Build cuda_12.6.r12.6/compiler.34714021_0

Latest commit or version

0.3.1.0832

no problem with preceding versions

[gfxenjoyer]

Checking Cuda Math API and cmath, there is no isnanf only isnan. It compiles if you remove the f.

In your %HOME%\.cargo\git\checkouts\candle-<some-hash>\<git-hash>\candle-kernels\src\cuda-utils.cuh remove the f in line 238 like so:

__device__ __forceinline__ bool isnang(__nv_fp8_e4m3 a) { return isnan(F8E4M3_TO_FLOAT(a)); }

I was able to build and run after that.
@EricLBuehler can you remove that f and push a hotfix?

[misureaudio]

Checking Cuda Math API and cmath, there is no isnanf only isnan. It compiles if you remove the f.

In your %HOME%\.cargo\git\checkouts\candle-<some-hash>\<git-hash>\candle-kernels\src\cuda-utils.cuh remove the f in line 238 like so:

__device__ __forceinline__ bool isnang(__nv_fp8_e4m3 a) { return isnan(F8E4M3_TO_FLOAT(a)); }

I was able to build and run after that. @EricLBuehler can you remove that f and push a hotfix?

It works. Confirmed. Many thanks!

[EricLBuehler]

@misureaudio @gfxenjoyer I pushed a fix in #859, can you please confirm and I can merge?

[misureaudio]

I managed to compile with the suggested mod. Confirmed. Thank you!

[EricLBuehler]

@misureaudio I merged #859!