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

### 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:

```text
csrc/sm90/prefill/sparse/phase1.cuh
```

### Problem

Inside:

```cpp
KernelTemplate::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:

```text
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:

```text
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:

```text
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:

```text
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::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::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::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:

```text
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

Start in csrc/sm90/prefill/sparse/phase1.cuh at KernelTemplate::devfunc and inspect the three named device lambdas. Reproduce the compile-only path with MSVC and NVCC for SM90a, then verify both compilers complete successfully while the captures, bodies, and calls remain unchanged.

Written by the indexing model from the issue text.