[Feat] Add HC ops for Qwen3.8-flash-next by Xia-Weiwen · Pull Request #578 · sgl-project/sgl-kernel-xpu

Xia-Weiwen · 2026-10-26T06:26:15Z

Add the following ops for HC in Qwen3.8-flash-next. The implementations of these ops are ported from DeepKLOX.

qwen4_hc_silu
qwen4_hc_gate_mix
qwen4_hc_combine
qwen4_hc_combine_norm

We also modified the kernel codes to improve performance.

Test

pytest tests/test_qwen4_hyperconnection.py

Benchmarks

python benchmark/bench_qwen4_hyperconnection.py

Copilot

🔵 Needs a closer look

The new low-level SYCL paths need XPU validation and human review, particularly because intended misalignment coverage is incomplete.

2 open issues

What changed in this PR

Adds BF16 HyperConnection operations for Qwen3.8-Flash-Next to the Intel XPU kernel library, with Python APIs, correctness tests, and benchmarks.

Changes:

Implements SYCL SiLU, gate mixing, residual combination, and fused combination with grouped Gemma RMSNorm.
Exposes the operations through C++ declarations, Torch registration, and Python wrappers.
Adds tests and benchmarks covering scalar/vector paths and decode/prefill workloads.

File	Description
`tests/test_qwen4_hyperconnection.py`	Tests correctness, validation, aliasing, and kernel-path boundaries.
`tests/run_suite.py`	Registers the HyperConnection tests.
`src/torch_extension_sycl.cc`	Registers four XPU operators with Torch.
`src/sycl/Qwen4HyperConnection.cpp`	Implements kernels, dispatch, and input validation.
`python/sgl_kernel/elementwise.py`	Adds four public Python wrappers.
`python/sgl_kernel/__init__.py`	Exports the new Python APIs.
`include/sgl_kernel_ops.h`	Declares the C++ operator interfaces.
`benchmark/run_suite.py`	Registers the HyperConnection benchmark.
`benchmark/bench_qwen4_hyperconnection.py`	Compares fused kernels with eager and compiled operations.

_{Review effort: Balanced}

Copilot

🔵 Needs a closer look

The new SYCL kernels and fused reductions require Intel XPU build, correctness, and performance validation before human approval.

No open issues 🎉

1 previously missed issue

In code that hasn't changed since last review

Use median return mode for advertised timing statistics

benchmark/bench_qwen4_hyperconnection.py:36

The output labels these timings as medians, but triton.testing.do_bench defaults to the mean when neither quantiles nor return_mode is specified. Consequently, every timing column reports a different statistic than advertised. Pass return_mode="median" explicitly so these results use the stated statistic.

2 resolved since last review

_{Review effort: Balanced}

airMeng · 2026-10-26T22:58:41Z

any effective tflops/BW sharable here?

mingfeima

simplify the code.

ai puts too many defensive checks.

mingfeima · 2026-10-27T01:07:36Z

+struct GateMixVecKernel {
+  const bf16x8* input;
+  const bf16x8* gate;
+  bf16x8* output;
+  int vecs_per_row;
+
+  [[sycl::reqd_sub_group_size(SG_SIZE)]] void operator()(sycl::nd_item<2> item) const {
+    const int64_t token = item.get_global_id(0);
+    const int first = item.get_global_id(1);
+    const int step = item.get_global_range(1);
+    const int64_t base = token * HC * vecs_per_row;
+    bf16x8 xs[ITEMS][HC];
+    bf16x8 gs[ITEMS][HC];
+#pragma unroll
+    for (int p = 0; p < ITEMS; ++p) {
+      const int c = first + p * step;
+      if (c < vecs_per_row) {
+#pragma unroll
+        for (int s = 0; s < HC; ++s) {
+          xs[p][s] = input[base + s * vecs_per_row + c];
+          gs[p][s] = gate[base + s * vecs_per_row + c];
+        }
+      }
+    }
+#pragma unroll
+    for (int p = 0; p < ITEMS; ++p) {
+      const int c = first + p * step;
+      if (c < vecs_per_row) {
+        float acc[VEC_SIZE] = {};
+#pragma unroll
+        for (int s = 0; s < HC; ++s) {
+          float xv[VEC_SIZE], gv[VEC_SIZE];
+          to_float(xs[p][s], xv);
+          to_float(gs[p][s], gv);
+#pragma unroll
+          for (int k = 0; k < VEC_SIZE; ++k) {
+            acc[k] += sigmoid(gv[k]) * xv[k];
+          }
+        }
+#pragma unroll
+        for (int k = 0; k < VEC_SIZE; ++k) {
+          acc[k] /= static_cast<float>(HC);
+        }
+        output[token * vecs_per_row + c] = to_bf16(acc);
+      }
+    }
+  }
+};
+
+struct GateMixScalarKernel {
+  const bf16* input;
+  const bf16* gate;
+  bf16* output;
+  int64_t elements;
+  int64_t hidden;
+  int64_t hc;
+
+  void operator()(sycl::nd_item<1> item) const {
+    const int64_t i = item.get_global_id(0);
+    if (i >= elements) return;
+    const int64_t row_base = (i / hidden) * hidden * hc;
+    const int64_t d = i % hidden;
+    float acc = 0.0f;
+    for (int64_t s = 0; s < hc; ++s) {
+      const int64_t o = row_base + s * hidden + d;
+      acc += sigmoid(static_cast<float>(gate[o])) * static_cast<float>(input[o]);
+    }
+    output[i] = static_cast<bf16>(acc / static_cast<float>(hc));
+  }
+};


here in this file, each kernel has vec and scalar path,

you may consider sharing one vector-width-parameterized implementation between the vec/scalar paths for SiLU, GateMix and Combine, following SigmoidMul.cpp.

this is going to simplify a lot of the code.

Ok. Thanks. This is a port from DeepKLOX. We didn't make many changes. We are simplifying the code.

Code is simplified.

Xia-Weiwen · 2026-10-27T02:02:29Z

any effective tflops/BW sharable here?

hc_silu (low-rank mix activation, 320 wide)
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            10.7            10.3             5.2      0.2      0    0.000      2.06        1.97
       4            10.2             4.7             5.1      1.0      0    0.001      2.02        0.93
      16             9.5             5.0             5.1      4.1      0    0.005      1.87        1.00
      64             7.7             4.8             5.1     16.1      2    0.020      1.50        0.94
     256             8.7             5.8             5.4     61.0      7    0.076      1.63        1.08
    1024            10.3            10.0             7.6    172.1     20    0.215      1.35        1.31
    8192            40.4            50.9            29.3    358.5     41    0.448      1.38        1.74
   32768           191.8           185.9           114.4    366.6     42    0.458      1.68        1.62

hc_gate_mix
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            10.9             4.7             7.0      6.5      1    0.008      1.55        0.67
       4            12.0             6.4             7.3     25.2      3    0.029      1.64        0.87
      16            13.0             6.1             9.4     78.7      9    0.092      1.38        0.65
      64            21.4            11.1            13.3    221.8     26    0.259      1.61        0.84
     256            60.5            31.0            32.3    365.2     42    0.426      1.87        0.96
    1024           356.9           122.0           115.4    409.0     47    0.477      3.09        1.06
    8192          2764.9           985.7           911.6    414.1     48    0.483      3.03        1.08
   32768         10920.4          3942.3          3639.6    414.9     48    0.484      3.00        1.08

hc_combine (precomputed injection)
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            11.7             5.9             5.0      9.3      1    0.004      2.34        1.18
       4            13.2             6.4             6.4     29.0      3    0.013      2.08        1.00
      16            14.7             6.1             6.4    115.7     13    0.051      2.30        0.95
      64            26.4             9.9            12.1    244.3     28    0.109      2.18        0.82
     256            83.0            25.4            25.3    465.5     54    0.207      3.27        1.00
    1024           367.5           110.6           107.6    438.7     51    0.195      3.42        1.03
    8192          2726.2           965.1           949.1    397.8     46    0.177      2.87        1.02
   32768         10775.7          3892.2          3826.9    394.6     45    0.175      2.82        1.02

hc_combine_norm (combine + grouped Gemma RMSNorm; sgl_fused is the kernel under test)
  tokens        eager_us      compile_us  sgl_2kernel_us    sgl_fused_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            48.1             5.9             7.7             6.1     14.4      2    0.012      7.93        0.97
       4            47.1             7.9             7.2             5.9     48.4      6    0.048      7.95        1.33
      16            58.1            10.2             9.4             7.9    137.0     16    0.145      7.33        1.29
      64            85.5            20.7            14.6            11.1    384.2     44    0.412      7.67        1.86
     256           310.3            62.9            37.3            32.3    528.4     61    0.568      9.61        1.95
    1024          1558.1           260.3           211.7           156.7    435.2     50    0.468      9.94        1.66
    8192         11771.6          2302.0          1783.7          1377.2    396.0     46    0.426      8.55        1.67
   32768         46640.5          9265.7          7147.8          5558.3    392.4     45    0.423      8.39        1.67

The kernels are slower than torch.compile in many cases. However, since it's a port, we didn't check performance. We are now tuning the kernel to see if we can close the perf gap.

Xia-Weiwen · 2026-10-27T03:14:31Z

any effective tflops/BW sharable here?

After some optimizations:

hc_silu (low-rank mix activation, 320 wide)
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1             8.7             8.2             4.1      0.3      0    0.000      2.11        1.99
       4            10.9             5.1             4.2      1.2      0    0.002      2.61        1.22
      16             9.0             4.9             4.2      4.9      1    0.006      2.14        1.16
      64             8.3             5.9             5.1     16.1      2    0.020      1.63        1.15
     256             7.9             6.0             4.6     71.2      8    0.089      1.72        1.31
    1024             9.9             9.8             6.5    201.7     23    0.252      1.53        1.50
    8192            40.4            49.5            21.0    499.8     58    0.625      1.92        2.36
   32768           191.8           185.9            90.1    465.6     54    0.582      2.13        2.06

hc_gate_mix
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            11.2             5.0             4.6     10.0      1    0.012      2.43        1.10
       4            11.2             5.2             4.7     39.6      5    0.046      2.41        1.11
      16            13.4             6.1             7.9     93.8     11    0.109      1.70        0.77
      64            22.5            11.0            10.7    275.2     32    0.321      2.10        1.03
     256            61.3            33.4            29.0    406.9     47    0.475      2.11        1.15
    1024           357.0           121.7           113.0    417.6     48    0.487      3.16        1.08
    8192          2765.0           985.7           908.3    415.6     48    0.485      3.04        1.09
   32768         10920.3          3942.4          3631.0    415.9     48    0.485      3.01        1.09

hc_combine (precomputed injection)
  tokens        eager_us      compile_us          sgl_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            11.9             4.6             4.3     10.8      1    0.005      2.78        1.08
       4            12.1             6.0             4.4     41.6      5    0.018      2.72        1.36
      16            14.8             5.5             5.3    140.4     16    0.062      2.81        1.05
      64            26.7             9.9             8.5    348.7     40    0.155      3.16        1.17
     256            76.7            28.9            21.0    560.6     65    0.249      3.65        1.38
    1024           367.4           110.6           102.5    460.6     53    0.205      3.59        1.08
    8192          2726.8           964.9           938.7    402.2     46    0.179      2.90        1.03
   32768         10775.0          3891.4          3791.2    398.4     46    0.177      2.84        1.03

hc_combine_norm (combine + grouped Gemma RMSNorm; sgl_fused is the kernel under test)
  tokens        eager_us      compile_us  sgl_2kernel_us    sgl_fused_us     GB/s  %copy   TFLOPS  vs_eager  vs_compile
       1            42.2             6.3             6.7             4.9     17.8      2    0.015      8.62        1.30
       4            47.2             7.4             7.2             5.1     56.5      7    0.056      9.30        1.46
      16            54.4            11.6             7.9             6.2    174.1     20    0.184      8.73        1.86
      64            85.2            20.9            16.9            10.0    430.1     50    0.461      8.56        2.10
     256           319.1            63.0            33.3            28.8    591.9     68    0.637     11.07        2.19
    1024          1558.3           260.3           206.4           152.9    446.0     51    0.480     10.19        1.70
    8192         11770.4          2301.9          1773.9          1358.0    401.6     46    0.432      8.67        1.70
   32768         46630.7          9266.4          7113.3          5469.9    398.8     46    0.429      8.52        1.69

airMeng · 2026-10-27T17:42:42Z

+    const int local_row = sg_id / SPLIT;
+    const int part = sg_id - local_row * SPLIT;
+    const uint32_t group_row = item.get_group(0) * ROWS_PER_GROUP + local_row;
+    const bool active = group_row < rows;
+    const uint32_t row = active ? group_row : 0;


local == 0
groups == rows
then row always equals to group_row

there is only 1 template activated

Fixed. thanks.

airMeng · 2026-10-27T17:43:36Z

+}
+
+// An output may alias an input exactly (in place) but must not partially overlap it.
+void check_no_partial_overlap(const at::Tensor& output, std::initializer_list<const at::Tensor*> inputs) {


over design, all the callers supply only one input tensor

Fixed. thanks.

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

…alignment Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…nd TFLOPS per op Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…tead of IEEE divides, 2D launches without integer division, preloaded combine_norm loads with row-count-tuned tiles, compile-time injection dispatch in combine, and 2-element vectors for small silu/gate_mix inputs

Xia-Weiwen added this pull request to stack #576 October 26, 2026 06:26

Xia-Weiwen requested a balanced review from Copilot October 26, 2026 07:43

Copilot started reviewing on behalf of Xia-Weiwen October 26, 2026 07:43

View session

Copilot AI reviewed Oct 26, 2026

View reviewed changes

Comment thread tests/test_qwen4_hyperconnection.py Resolved

Comment thread tests/test_qwen4_hyperconnection.py Resolved

Xia-Weiwen requested a balanced review from Copilot October 26, 2026 07:53

Copilot started reviewing on behalf of Xia-Weiwen October 26, 2026 07:53

View session

Copilot AI reviewed Oct 26, 2026

View reviewed changes

Xia-Weiwen marked this pull request as ready for review October 26, 2026 08:14

Xia-Weiwen added the run-ci label Oct 26, 2026

Xia-Weiwen requested review from airMeng and mingfeima October 26, 2026 08:21

mingfeima reviewed Oct 26, 2026

View reviewed changes

Comment thread benchmark/bench_qwen4_hyperconnection.py Resolved

Comment thread src/sycl/Qwen4HyperConnection.cpp Resolved

Comment thread src/sycl/Qwen4HyperConnection.cpp Resolved

Xia-Weiwen requested a review from mingfeima October 26, 2026 09:47

mingfeima requested changes Oct 27, 2026

View reviewed changes

Xia-Weiwen requested a review from mingfeima October 27, 2026 03:45

airMeng reviewed Oct 27, 2026

View reviewed changes

Xia-Weiwen requested a review from airMeng October 28, 2026 02:58

airMeng approved these changes Oct 28, 2026

View reviewed changes

mingfeima approved these changes Oct 29, 2026

View reviewed changes

Base automatically changed from feat/qwen4-short-conv to main October 29, 2026 07:54

Xia-Weiwen and others added 7 commits October 29, 2026 00:54

[Feat] Add HC ops for Qwen3.8-flash-next

a754b95

Add qwen4_hc_combine_norm

5e6f5f6

Refine code

9b1cbfa

Apply batched suggestions from code review

389caf7

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Address HC review: vector naming, element-wise bf16 helpers, test mis…

ecba799

…alignment Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Simplify HC kernels with shared vector-width templates; report GB/s a…

6da9eeb

…nd TFLOPS per op Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Xia-Weiwen added 5 commits October 29, 2026 00:54

Use sycl_kernel_submit

5b9997e

Reuse some kernel code

0382bc8

Simplify code

04b1727

Move python wrappers to hc.py

6dabd5b

Refine code

8684311

Xia-Weiwen force-pushed the feat/qwen4-hc-ops branch from d2540a3 to 8684311 Compare October 29, 2026 07:54

Xia-Weiwen merged commit 81bbce4 into main Oct 30, 2026
6 checks passed

Conversation

Xia-Weiwen commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Uh oh!

Uh oh!

Uh oh!

airMeng commented Oct 26, 2026

Uh oh!

mingfeima left a comment

Choose a reason for hiding this comment

Uh oh!

mingfeima Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

Xia-Weiwen Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

Xia-Weiwen Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

Xia-Weiwen commented Oct 27, 2026

Uh oh!

Xia-Weiwen commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

airMeng Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

Xia-Weiwen Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

airMeng Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

Xia-Weiwen Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Xia-Weiwen commented Oct 26, 2026 •

edited

Loading

Xia-Weiwen commented Oct 27, 2026 •

edited

Loading