[Portability][MSVC] GNU always_inline attributes on SM90 device lambdas fail to compile

Author: rogerobergCreated Sep 6, 2026Updated Sep 6, 2026

Background

Hi!

I encountered this issue while building FlashMLA through SystemPanic/vllm-windows for a deployment serving DeepSeek-V4-Flash-0731 on H200 NVL GPUs.

The Windows build uses MSVC as the NVCC host compiler and targets SM90a.

I understand that native Windows may not be part of the supported build matrix. This report concerns a localized C++/CUDA source-portability issue in:

csrc/sm90/prefill/sparse/phase1.cuh

Problem

Inside:

cpp

KernelTemplate<D_QK, HAVE_TOPK_LENGTH>::devfunc

three device lambdas place the GNU-style attribute __attribute__((always_inline)) after the lambda declarator:

cpp

auto pipelined_wait_and_qkt_gemm_l =
    [&]() __attribute__((always_inline)) {
      // ...
    };

The same form is used for:

pipelined_wait_and_qkt_gemm_l
pipelined_wait_and_qkt_gemm_r
pipelined_wait_and_qkt_gemm

MSVC does not accept the GNU attribute at this position in a lambda expression. NVCC's device compilation pass also rejects the syntax before the lambda body.

Reproduction

The issue was reproduced with a compile-only source containing three lambdas inside a __device__ function, matching the relevant syntax used in phase1.cuh.

Original form:

cpp

auto pipelined_wait_and_qkt_gemm_l =
    [&]() __attribute__((always_inline)) {
      // ...
    };

auto pipelined_wait_and_qkt_gemm_r =
    [&]() __attribute__((always_inline)) {
      // ...
    };

auto pipelined_wait_and_qkt_gemm =
    [&]() __attribute__((always_inline)) {
      // ...
    };

The device function was instantiated from a __global__ launcher to ensure that NVCC compiled the relevant device path.

This was a compile-only reproduction. No executable was run and no GPU kernel was executed.

The original form failed with MSVC:

return code: 2

error C3260: 'identifier': skipping unexpected token(s) before lambda body
error C3260: 'identifier': skipping unexpected token(s) before lambda body
error C3260: 'identifier': skipping unexpected token(s) before lambda body

error C2064: term does not evaluate to a function taking 0 arguments
error C2064: term does not evaluate to a function taking 0 arguments
error C2064: term does not evaluate to a function taking 0 arguments

The original form also failed with NVCC:

return code: 2

error: expected a "{" introducing a lambda body
error: expected a "{" introducing a lambda body
error: expected a "{" introducing a lambda body

The corrected form, with only the three attributes removed, compiled successfully:

cl:
  return code 0

nvcc:
  return code 0

Tested fix

Remove __attribute__((always_inline)) from the three device lambdas while preserving each lambda's capture list, parameter list, and body:

diff

diff --git a/csrc/sm90/prefill/sparse/phase1.cuh b/csrc/sm90/prefill/sparse/phase1.cuh
index bf2fff8..f409a0d 100644
--- a/csrc/sm90/prefill/sparse/phase1.cuh
+++ b/csrc/sm90/prefill/sparse/phase1.cuh
@@ -259,7 +259,7 @@ __device__ void KernelTemplate<D_QK, HAVE_TOPK_LENGTH>::devfunc(const SparseAttn
         if (warpgroup_idx == 0) {
             // Warpgroup 0

-            auto pipelined_wait_and_qkt_gemm_l = [&]() __attribute__((always_inline)) {
+            auto pipelined_wait_and_qkt_gemm_l = [&]() {
                 plan.bar_k0_ready[0].wait(cur_bar_wait_phase);
                 qkt_gemm_one_tile(Warpgroup0{}, 0, true);
                 qkt_gemm_one_tile(Warpgroup0{}, 1, false);
@@ -268,7 +268,7 @@ __device__ void KernelTemplate<D_QK, HAVE_TOPK_LENGTH>::devfunc(const SparseAttn
                 warpgroup_commit_batch();
             };

-            auto pipelined_wait_and_qkt_gemm_r = [&]() __attribute__((always_inline)) {
+            auto pipelined_wait_and_qkt_gemm_r = [&]() {
                 plan.bar_k0_ready[1].wait(cur_bar_wait_phase);
                 qkt_gemm_one_tile(Warpgroup0{}, 4, false);
                 qkt_gemm_one_tile(Warpgroup0{}, 5, false);
@@ -377,7 +377,7 @@ __device__ void KernelTemplate<D_QK, HAVE_TOPK_LENGTH>::devfunc(const SparseAttn
         } else {
             // Warpgroup 1

-            auto pipelined_wait_and_qkt_gemm = [&]() __attribute__((always_inline)) {
+            auto pipelined_wait_and_qkt_gemm = [&]() {
                 plan.bar_k1_ready[1].wait(cur_bar_wait_phase);
                 qkt_gemm_one_tile(Warpgroup1{}, 4, true);
                 qkt_gemm_one_tile(Warpgroup1{}, 5, false);

Controlled compile-only results:

Original:
  cl:   return code 2
  nvcc: return code 2

Corrected:
  cl:   return code 0
  nvcc: return code 0

Behavior

The lambda capture lists remain unchanged:

cpp

[&]

The empty parameter lists remain unchanged:

cpp

()

The lambda bodies, calls, synchronization, barrier waits, warp-group operations, and QK-tile operations also remain unchanged.

Only the unsupported __attribute__((always_inline)) syntax is removed.

The corrected source allows both the MSVC host parsing pass and the NVCC device compilation pass to complete successfully.

The compiler may still choose to inline these lambdas through its normal optimization process. Removing the forced-inline attribute could theoretically affect generated code, register pressure, code size, or performance, but those effects were not measured as part of this compile-only reproduction.

Environment

Windows Server 2025
Visual Studio Build Tools with MSVC 19.51.36248
MSVC toolset 14.51.36231
CUDA/NVCC 13.2, V13.2.51
C++20
--gpu-architecture=sm_90a
--expt-relaxed-constexpr
--expt-extended-lambda
Python 3.12
PyTorch 2.11.0+cu130
Tested for H200 NVL GPUs
SystemPanic/vllm-windows 0.25-based deployment
DeepSeek-V4-Flash-0731

Source: deepseek-ai/FlashMLA

View original on GitHub View discussion on GitHub