[cuda] fix MSVC/Windows CUDA compilation errors#31
Open
Shraman123 wants to merge 1 commit into
Open
Conversation
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).
Owner
|
Verdict: mergeable as-is pending gates (compile-only change surface; no GPU behavior change). Reviewed the three fixes against current
Caveats:
Gate runs tomorrow: standard Linux CUDA suite + one covtype md5 confirm (expected: unchanged). 🤖 Review drafted with Claude Code |
3 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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):
fmt's
use_utf8static_assert (fmt/base.h) detects/utf-8by checking how the compiler encodes a "§" literal. Passing-Xcompiler=/utf-8reachescl.exefor nvcc's host-side compile, but not nvcc's own front-end parse of.cufiles, so any.cufile pulling in fmt headers failed to compile. DefinesFMT_UNICODE=0for CUDA translation units to skip the assertion - it only gates an unused Windows-console UTF-8 print codepath, not core formatting.cuda_single_gpu_tree_learner.cppcalledstd::max(static_cast<size_t>(...), 1UL).unsigned longis 32-bit on LLP64 platforms (MSVC/Windows) butsize_tis 64-bit there, sostd::max's matching-type template deduction failed to compile (works on Linux, whereunsigned longis 64-bit). Usesstatic_cast<size_t>(1)instead.kMinScore/kMaxScore/kEpsilon/kZeroThreshold(meta.h) are used directly inside__global__CUDA kernels (e.g.cuda_best_split_finder.cu), but were declared as plainconstrather thanconstexpr. Whether aconstglobal 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.constexprmakes this unambiguous on every toolchain.Test plan
examples/binary_classification,device_type=cuda) - logloss decreased and AUC improved over 10 iterations as expected