Add compile-time 256-bit vector guard for pre-Blackwell - #19794
Conversation
|
Warning You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again! |
cac753f to
6b4e4de
Compare
| #define SGL_ARCH_IS_HOPPER_PLUS (__CUDA_ARCH__ >= 900) | ||
| #define SGL_ARCH_IS_BLACKWELL_PLUS ((__CUDA_ARCH__ >= 1000) && (CUDA_VERSION >= 12090)) |
There was a problem hiding this comment.
SGL_ARCH_IS_HOPPER_PLUS / SGL_ARCH_IS_BLACKWELL_PLUS . Can we improve the naming? Maybe we can refer to other open-source libraries as example.
There was a problem hiding this comment.
What about SGL_ARCH_HOPPER_OR_GREATER / SGL_ARCH_BLACKWELL_OR_GREATER, inspired by .NET macro naming. See https://learn.microsoft.com/en-us/dotnet/standard/frameworks#preprocessor-symbols ,it looks pretty clear
| @@ -58,6 +58,7 @@ __global__ void qknorm_across_heads_reg_kernel( | |||
| const T* __restrict__ k_weight, | |||
| int vec_hidden_size, | |||
| float eps) { | |||
| if constexpr (VEC_SIZE_IN_BYTE > device::kMaxVecBytes) return; | |||
There was a problem hiding this comment.
Given that we are testing the architecture from compile-time macros, we should remove this line and change the kernel dispatch logic based on the compile-time macro values. We should never instantiate invalid kernel (e.g. VEC_SIZE_IN_BYTE = 32 on hopper), i.e. the constexpr guard in device side should be moved to host side.
There was a problem hiding this comment.
found a blocker for such change
-
SGL_ARCH_IS_BLACKWELL_PLUSrelies on__CUDA_ARCH__, which is only defined on the device side. In host-siderun()(qknorm_across_heads.cuh:L168), it always evaluates to 0, even on Blackwell. -
cc_majoris a runtime value. It gives the correct answer but can't prevent template instantiation. Bothkernel<16>andkernel<32>are always instantiated (qknorm_across_heads.cuh:L216). -
To achieve zero-instantiation on the host side, we'd need the JIT compilation pipeline to pass a compiler flag (e.g.
-DSGL_CC_MAJOR=10) so that host code can use#ifto select the kernel template. This requires changes toload_jit().
The current device-side if constexpr approach is correct, though the invalid kernel gets instantiated with an empty body. Considering the additional complexity involved (also somewhat out of scope for this PR), I'd suggest keeping it as-is and maybe addressing it in a follow-up PR if needed. What do you think?
There was a problem hiding this comment.
Actually, in Python-side, we can get the architecture information from torch. We may add some macros which indicate the target architecture (which works for both host/device code) in default compile flags.
There was a problem hiding this comment.
Makes sense. I'll add the default architecture macro in load_jit() and update the host-side dispatch in a separate PR, since it will have a broader impact. This PR will keep the current approach as-is.
There was a problem hiding this comment.
I think that a value that differs in host/device code is highly error prune. BTW this PR seems to break some JIT kernel unit-tests https://github.com/sgl-project/sglang/actions/runs/22722498873/job/65904747412?pr=19794
PTAL
There was a problem hiding this comment.
Why we not do like this:
void run_kernel(...) {
// ...
#if defined(SGL_ARCH_IS_BLACKWELL_PLUS)
launch_kernel<32>(...);
#else
launch_kernel<16>(...);
#endif
}There was a problem hiding this comment.
Why we not do like this:
void run_kernel(...) { // ... #if defined(SGL_ARCH_IS_BLACKWELL_PLUS) launch_kernel<32>(...); #else launch_kernel<16>(...); #endif }
That's because "SGL_ARCH_IS_BLACKWELL_PLUS relies on CUDA_ARCH, which is only defined on the device side. In host-side run() (qknorm_across_heads.cuh:L168), it always evaluates to 0, even on Blackwell."
That's why I suggest we introduce new compile flags to specify the target architecture in both host/device code.
There was a problem hiding this comment.
I opened a follow-up PR #20103 to inject target arch flag into JIT compilation so that arch-dependent macros work in both host and device passes, as @DarkSharpness suggested. cc @BBuf
After that follow-up PR merged, this PR will be much simpler.
|
/tag-and-rerun-ci |
69ae625 to
bab3637
Compare
bab3637 to
7b24cdb
Compare
|
Hi @DarkSharpness, I rewrote this PR, rebased on #20103. Added |
Motivation
Follow-up from #19770 review discussion. On pre-Blackwell GPUs (SM < 100), CUDA only supports 128-bit vector load/store. The 32-byte
AlignedVectorpath compiles but should never run. This PR adds a compile-time guard so invalid 256-bit instantiations are caught at compile time rather than relying solely on runtime dispatch.Modifications
utils.cuh: AddSGL_ARCH_IS_HOPPER_PLUSandSGL_ARCH_IS_BLACKWELL_PLUSmacros. Refactor existing PDL__CUDA_ARCH__checks to use them.vec.cuh: SplitAlignedVectorstatic_assertby arch — 16 bytes max on pre-Blackwell, 32 bytes on Blackwell+.fused_add_rmsnorm.cuh/qknorm_across_heads.cuh: Addif constexprearly return in 32B kernel template on pre-Blackwell, preventing the body (which instantiates 32-byteAlignedVector) from being compiled.Accuracy Tests
No model output changes. This is a compile-time guard only — kernel logic is unchanged.
Benchmarking and Profiling
No performance impact. The guard is resolved at compile time (
if constexpr/#ifpreprocessor), producing identical codegen for all existing paths.Checklist