Skip to content

[cuda] fix MSVC/Windows CUDA compilation errors#31

Open
Shraman123 wants to merge 1 commit into
BelixRogner:masterfrom
Shraman123:windows-msvc-cuda-build-fixes
Open

[cuda] fix MSVC/Windows CUDA compilation errors#31
Shraman123 wants to merge 1 commit into
BelixRogner:masterfrom
Shraman123:windows-msvc-cuda-build-fixes

Conversation

@Shraman123

Copy link
Copy Markdown

Summary

Three independent issues blocked building the CUDA backend with MSVC + nvcc on Windows (found while getting a single-GPU Windows CUDA build working, alongside #30):

  1. fmt's use_utf8 static_assert (fmt/base.h) detects /utf-8 by checking how the compiler encodes a "§" literal. Passing -Xcompiler=/utf-8 reaches cl.exe for nvcc's host-side compile, but not nvcc's own front-end parse of .cu files, so any .cu file pulling in fmt headers failed to compile. Defines FMT_UNICODE=0 for CUDA translation units to skip the assertion - it only gates an unused Windows-console UTF-8 print codepath, not core formatting.

  2. cuda_single_gpu_tree_learner.cpp called std::max(static_cast<size_t>(...), 1UL). unsigned long is 32-bit on LLP64 platforms (MSVC/Windows) but size_t is 64-bit there, so std::max's matching-type template deduction failed to compile (works on Linux, where unsigned long is 64-bit). Uses static_cast<size_t>(1) instead.

  3. kMinScore/kMaxScore/kEpsilon/kZeroThreshold (meta.h) are used directly inside __global__ CUDA kernels (e.g. cuda_best_split_finder.cu), but were declared as plain const rather than constexpr. Whether a const global with a constant initializer is usable in device code without a device-side symbol is toolchain-dependent, and MSVC's nvcc front-end rejected it ("undefined in device code") where other host compilers accept it. constexpr makes this unambiguous on every toolchain.

Test plan

  • Builds cleanly with MSVC 2022 (Build Tools) + CUDA 12.5
  • Ran real single-GPU CUDA training end-to-end on a GTX 1650 (examples/binary_classification, device_type=cuda) - logloss decreased and AUC improved over 10 iterations as expected

Three independent issues blocked building the CUDA backend with
MSVC + nvcc on Windows (found while getting a single-GPU Windows
CUDA build working):

1. fmt's use_utf8 static_assert (fmt/base.h) detects /utf-8 by
   checking how the compiler encodes a "§" literal. Passing
   -Xcompiler=/utf-8 reaches cl.exe for nvcc's host-side compile,
   but not nvcc's own front-end parse of .cu files, so any .cu file
   pulling in fmt headers failed to compile. Define FMT_UNICODE=0
   for CUDA translation units to skip the assertion; it only gates
   an unused Windows-console UTF-8 print codepath, not core
   formatting.

2. cuda_single_gpu_tree_learner.cpp called
   std::max(static_cast<size_t>(...), 1UL). unsigned long is 32-bit
   on LLP64 platforms (MSVC/Windows) but size_t is 64-bit there, so
   std::max's matching-type template deduction failed to compile
   (works on Linux, where unsigned long is 64-bit). Use
   static_cast<size_t>(1) instead.

3. kMinScore/kMaxScore/kEpsilon/kZeroThreshold (meta.h) are used
   directly inside __global__ CUDA kernels (e.g.
   cuda_best_split_finder.cu), but were declared as plain const
   rather than constexpr. Whether a const global with a constant
   initializer is usable in device code without a device-side
   symbol is toolchain-dependent, and MSVC's nvcc front-end rejected
   it ("undefined in device code") where other host compilers
   accept it. constexpr makes this unambiguous on every toolchain.

Verified: builds cleanly with MSVC 2022 + CUDA 12.5, and single-GPU
CUDA training runs end-to-end on a GTX 1650 (binary classification
example, 10 iterations, logloss/AUC converge as expected).
@BelixRogner

Copy link
Copy Markdown
Owner

Verdict: mergeable as-is pending gates (compile-only change surface; no GPU behavior change).

Reviewed the three fixes against current master:

  • include/LightGBM/meta.h constconstexpr: correct and principled — device-code usability of a namespace-scope const with constant initializer is genuinely toolchain-dependent, and constexpr states the intent. Verified the four constants are still plain const on master (meta.h:51-57), so this applies cleanly.
  • std::max(..., static_cast<size_t>(1)) in AllocateBitset: correct LLP64 fix — 1UL is 32-bit on MSVC while size_t is 64-bit, so template deduction fails there. The 1UL sites are still present on master (cuda_single_gpu_tree_learner.cpp:2733-2734), applies cleanly.
  • -Xcompiler=/utf-8 -DFMT_UNICODE=0 for MSVC+CUDA: the comment explaining why -Xcompiler=/utf-8 alone doesn't satisfy fmt's use_utf8 check (nvcc's own front-end parse vs the host cl.exe it shells out to) is exactly the kind of definition-not-workaround rationale we want in CMake. FMT_UNICODE=0 only disables the Windows-console UTF-8 print path, which we don't use.

Caveats:

  • We have no Windows CUDA CI, so this is review-verified only. The Linux gate run tomorrow will confirm the constexpr/static_cast changes are behavior-neutral there (they should be bit-neutral). If you can attach a successful local Windows build log (nvcc + MSVC version) to the PR, that closes the loop.
  • No CI checks ran on this PR (fork workflow approval); needs a maintainer approval-to-run or a re-push after [cuda] make NCCL optional for single-GPU builds (USE_NCCL) #30 lands so both go through together — the two PRs are natural companions (this one makes MSVC compile, [cuda] make NCCL optional for single-GPU builds (USE_NCCL) #30 removes the NCCL hard dependency Windows can't satisfy).

Gate runs tomorrow: standard Linux CUDA suite + one covtype md5 confirm (expected: unchanged).

🤖 Review drafted with Claude Code

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants