danielamadori opened a new issue, #822:
URL: https://github.com/apache/tvm-ffi/issues/822
### Summary
On Windows, `_generate_ninja_build` emits a CUDA build line that `nvcc`
refuses, and a link line that cannot resolve the CUDA runtime. Four separate
problems in the same `if IS_WINDOWS:` branch of
`python/tvm_ffi/cpp/extension.py`; each one is only reachable after the
previous is fixed, which is why they are reported together.
The Linux branch of the same function is correct in every one of these
places.
Measured with **apache-tvm-ffi 0.1.13.post3**, CUDA 12.6, MSVC, Windows 11.
**0.1.14.post1 carries the identical code** in all four places, so upgrading
does not help.
### 1. `-Xcompiler` takes a single argument, so `/O2` becomes an input file
```python
default_cuda_cflags = ["-Xcompiler", "/std:c++17", "/O2"]
```
`-Xcompiler` forwards **one** argument to the host compiler. `nvcc` hands
`/std:c++17` to `cl.exe` and is left holding `/O2`, which it reads as a second
input file:
```
nvcc fatal : A single input file is required for a non-link phase when an
outputfile is specified
```
Minimal reproduction, with a `.cu` containing only `__global__ void k(){}`:
```
> nvcc -Xcompiler /std:c++17 /O2 -c k.cu -o k.obj
nvcc fatal : A single input file is required for a non-link phase when an
outputfile is specified
> nvcc -Xcompiler /std:c++17,/O2 -c k.cu -o k.obj
(passes argument analysis)
```
Suggested fix: `["-Xcompiler", "/std:c++17,/O2"]`, or two separate
`-Xcompiler` groups.
### 2. that `/std:c++17` then overrides a caller asking for C++20
Not a defect on its own, but worth knowing because it bites immediately
after fixing (1). `nvcc` places its own `-std=` on the host line **first** and
every `-Xcompiler` group **after**, so tvm-ffi's `/std:c++17` ends up the last
`/std:` on `cl`'s command line and wins over `extra_cuda_cflags` passed by the
caller. Headers that need C++20 (`<concepts>`, `<source_location>`) then fail
with `C2039`.
A caller can work around it by appending `-Xcompiler
/std:c++20,/Zc:__cplusplus` (the `/Zc:__cplusplus` is needed because MSVC
otherwise reports `__cplusplus` as `199711L` even when compiling C++20, and
`<source_location>` is guarded on that), but a default that silently downgrades
the caller's language standard is surprising.
### 3. the CUDA runtime library is never linked on Windows
`cudart` appears exactly once in `extension.py`, inside the non-Windows
branch:
```python
if with_cuda:
default_ldflags += [
"-L{}".format(str(Path(_find_cuda_home()) / "lib64")),
"-lcudart",
]
```
The Windows branch sets `default_ldflags` and never adds a CUDA library path
or the runtime, so the link fails on the symbols that live in it:
```
cuda_0.o : error LNK2019: unresolved external symbol __cudaRegisterFatBinary
cuda_0.o : error LNK2019: unresolved external symbol
__cudaRegisterFatBinaryEnd
cuda_0.o : error LNK2019: unresolved external symbol
__cudaUnregisterFatBinary
cuda_0.o : error LNK2019: unresolved external symbol __cudaRegisterFunction
<module>.dll : fatal error LNK1120: 8 unresolved externals
```
Suggested fix, in the Windows branch:
```python
if with_cuda:
default_ldflags += [
'/LIBPATH:"{}"'.format(Path(_find_cuda_home()) / "lib" / "x64"),
"cudart.lib",
]
```
Note `lib/x64`, not `lib64`: on a Windows CUDA install
`v12.6/lib/x64/cudart.lib` exists and `v12.6/lib64` does not.
### 4. `/LIBPATH:` is not quoted, and the default CUDA path contains a space
With (3) applied but unquoted, the linker splits the standard CUDA location
`C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.6` on its space:
```
LINK : fatal error LNK1181: cannot open input file 'Files\NVIDIA.obj'
```
The same applies to the existing `f"/LIBPATH:{tvm_ffi_lib_path}"` one line
above. That one has never been observed failing, because the path lives inside
the virtualenv — but it would fail the same way on a machine whose user name
contains a space. Quoting a path that has no spaces is harmless, so both are
worth quoting.
### What works after all four
With the four fixed, a CUDA kernel built through `load_inline` compiles,
links and loads on Windows, and the engine using it serves generations normally.
### Environment
```
apache-tvm-ffi 0.1.13.post3 (0.1.14.post1 identical in these four places)
CUDA 12.6
OS Windows 11
compiler MSVC (cl.exe), nvcc 12.6
GPU NVIDIA RTX 3050 Laptop (sm_86)
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
