[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
meta-codesync merged 9 commits into from
Jul 9, 2026
Merged

[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM #20752

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
161baf5
Update
JCNTH Jul 6, 2026
d1ad151
Update
JCNTH Jul 6, 2026
4c20d00
Update
JCNTH Jul 7, 2026
642b82a
Update
JCNTH Jul 7, 2026
b8be2d1
Update
JCNTH Jul 9, 2026
523e082
Update
JCNTH Jul 9, 2026
c522131
Update
JCNTH Jul 9, 2026
9a0c964
Update
JCNTH Jul 9, 2026
061b444
Update
JCNTH Jul 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
* LICENSE file in the root directory of this source tree.
*/

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
: use_steel ? kQ4gswLinearGemmSteelWGSL
: use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
: kQ4gswLinearWGSL;
// f16-multiply steel: only when the device negotiated shader-f16; else the
// f32 steel kernel runs (fail-closed). Same bindings and tile.
if (use_steel) {
const WebGPUContext* ctx = get_default_webgpu_context();
if (ctx != nullptr && ctx->shader_f16_supported) {
shader_src = kQ4gswLinearGemmSteelHalfWGSL;
}
}
const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
device,
use_gemv,
Expand Down
12 changes: 10 additions & 2 deletions backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl
View file
Open in desktop
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
var nib: u32;
if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
$if DTYPE == "half":
dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
$else:
dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
}
Bs[br * BN + bc + j] = dqv;
}
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
for (var m: u32 = 0u; m < 4u; m = m + 1u) {
for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
$if DTYPE == "half":
for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
$else:
for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
}
}
workgroupBarrier();
Expand Down
2 changes: 2 additions & 0 deletions backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
DTYPE:
- VALUE: float
SUFFIX: ""
- VALUE: half
SUFFIX: half
shader_variants:
- NAME: q4gsw_linear_gemm_steel
126 changes: 126 additions & 0 deletions backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
* Copyright (c) Meta Platforms, Inc. and affiliates.
* All rights reserved.
*
* This source code is licensed under the BSD-style license found in the
* LICENSE file in the root directory of this source tree.
*/

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
M: u32,
N: u32,
K: u32,
K_packed: u32,
group_size: u32,
padded_N: u32,
has_bias: u32,
_pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
@builtin(local_invocation_id) lid: vec3<u32>) {
let nbN = (params.N + BN - 1u) / BN;
let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
let by = wid.x / nbN;
let row0 = by * BM;
let col0 = bx * BN;
let tid = lid.y * 16u + lid.x;
var acc: array<array<f32, 4>, 4>;
for (var m: u32 = 0u; m < 4u; m = m + 1u) {
for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
}
// A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
let ar = tid / 4u; // 0..63 (row in tile)
let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
// B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
let br = tid / 16u; // 0..15 (K within BK)
let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

var k0: u32 = 0u;
loop {
if (k0 >= params.K) { break; }
// stage activations (edge-masked on M; K is a multiple of BK for our shapes)
let arow = row0 + ar;
if (arow < params.M) {
let base = arow * params.K + k0 + ac;
As[ar * BK + ac + 0u] = f16(t_input[base]);
As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
} else {
As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
}
// stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
let kk = k0 + br; // K index for this shmem row
let scale_row = (kk / params.group_size) * params.padded_N;
for (var j: u32 = 0u; j < 4u; j = j + 1u) {
let n = col0 + bc + j;
var dqv: f16 = 0.0;
if (n < params.N) {
let byte_idx = n * params.K_packed + (kk >> 1u);
let word = t_weight[byte_idx >> 2u];
let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
var nib: u32;
if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
}
Bs[br * BN + bc + j] = dqv;
}
workgroupBarrier();
for (var k: u32 = 0u; k < BK; k = k + 1u) {
var a: array<f16, 4>;
var bvec: array<f16, 4>;
for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
for (var m: u32 = 0u; m < 4u; m = m + 1u) {
for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
}
}
workgroupBarrier();
k0 = k0 + BK;
}
for (var m: u32 = 0u; m < 4u; m = m + 1u) {
for (var n: u32 = 0u; n < 4u; n = n + 1u) {
let r = row0 + lid.y * 4u + m;
let c = col0 + lid.x * 4u + n;
if (r < params.M && c < params.N) {
var v = acc[m][n];
if (params.has_bias != 0u) { v = v + t_bias[c]; }
t_out[r * params.N + c] = v;
}
}
}
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading
, 'i'); if (__m === '*' || __re.test(location.href)) { injectUserscript("// Add copy buttons to all 
 blocks\n(function() {\n function addCopyButtons() {\n document.querySelectorAll('pre code').forEach(function(codeBlock) {\n if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;\n codeBlock.parentElement.setAttribute('data-copy-added', 'true');\n \n var btn = document.createElement('button');\n btn.textContent = 'Copy';\n btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';\n btn.onmouseover = function() { this.style.opacity = '1'; };\n btn.onmouseout = function() { this.style.opacity = '0.7'; };\n btn.onclick = function() {\n navigator.clipboard.writeText(codeBlock.textContent).then(function() {\n btn.textContent = 'Copied!';\n setTimeout(function() { btn.textContent = 'Copy'; }, 1500);\n });\n };\n codeBlock.parentElement.style.position = 'relative';\n codeBlock.parentElement.appendChild(btn);\n });\n }\n \n addCopyButtons();\n \n // Re-run on dynamic content\n var observer = new MutationObserver(addCopyButtons);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Add Copy Buttons to Code Blocks");
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Force GitHub README to respect dark mode\n(function() {\n var style = document.createElement('style');\n style.textContent = '\n .markdown-body {\n color-scheme: dark light;\n }\n .markdown-body pre { background: #161b22 !important; }\n .markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }\n .markdown-body table th, .markdown-body table td { border-color: #30363d !important; }\n .markdown-body img { background: #0d1117; }\n .markdown-body blockquote { border-left-color: #8b949e; }\n .markdown-body hr { border-color: #30363d; }\n ';\n document.head.appendChild(style);\n})();", "GitHub Dark Mode README Fix");
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Highlight search terms from Google/DuckDuckGo/Bing referrer\n(function() {\n var ref = document.referrer;\n var terms = [];\n \n if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {\n var url = new URL(ref);\n var q = url.searchParams.get('q') || url.searchParams.get('p');\n if (q) {\n terms = q.split(/\\s+/).filter(function(t) { return t.length > 2; });\n }\n }\n \n if (terms.length === 0) return;\n \n var style = document.createElement('style');\n style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';\n document.head.appendChild(style);\n \n function highlight(node) {\n if (node.nodeType === 3) { // text node\n var text = node.textContent;\n var found = false;\n terms.forEach(function(term) {\n var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\\]\\\\]/g, '\\\\') + ')', 'gi');\n if (regex.test(text)) {\n found = true;\n var frag = document.createDocumentFragment();\n var parts = text.split(regex);\n parts.forEach(function(part, i) {\n if (i % 2 === 0) {\n frag.appendChild(document.createTextNode(part));\n } else {\n var span = document.createElement('span');\n span.className = 'userscript-highlight';\n span.textContent = part;\n frag.appendChild(span);\n }\n });\n node.parentNode.replaceChild(frag, node);\n }\n });\n } else if (node.nodeType === 1 && node.childNodes) { // element\n var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];\n if (!skipTags.includes(node.tagName)) {\n Array.from(node.childNodes).forEach(highlight);\n }\n }\n }\n \n highlight(document.body);\n \n // Re-highlight on dynamic content\n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1 || node.nodeType === 3) highlight(node);\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Highlight Search Terms");
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Strip utm_, fbclid, gclid, etc. from all links on page\n(function() {\n var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',\n 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',\n 'ref', 'ref_src', 'source', 'medium', 'campaign'];\n \n function cleanUrl(url) {\n try {\n var u = new URL(url, window.location.origin);\n var changed = false;\n trackingParams.forEach(function(p) {\n if (u.searchParams.has(p)) {\n u.searchParams.delete(p);\n changed = true;\n }\n });\n return changed ? u.toString() : url;\n } catch (e) {\n return url;\n }\n }\n \n function cleanLinks() {\n document.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n \n cleanLinks();\n \n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1) {\n if (node.tagName === 'A') cleanLinks();\n node.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Remove Tracking Parameters from Links");
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Auto-enable theater mode on YouTube\n(function() {\n function tryTheater() {\n var btn = document.querySelector('button[aria-label=\"Theater mode\"], ytd-player #player button[title=\"Theater mode\"]');\n if (btn && !btn.classList.contains('activated')) {\n btn.click();\n }\n }\n \n // Try immediately\n tryTheater();\n \n // Try after navigation (SPA)\n var lastUrl = location.href;\n setInterval(function() {\n if (location.href !== lastUrl) {\n lastUrl = location.href;\n setTimeout(tryTheater, 500);\n }\n }, 1000);\n \n // Also try on player load\n var observer = new MutationObserver(tryTheater);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "YouTube Theater Mode Default");
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Remove or un-stick sticky/fixed headers that block content\n(function() {\n function unstick() {\n document.querySelectorAll('header, nav, [role=\"banner\"], .header, .navbar, .sticky, .fixed-top, [style*=\"position: fixed\"], [style*=\"position:sticky\"]').forEach(function(el) {\n if (el.style.position === 'fixed' || el.style.position === 'sticky' || \n getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {\n el.style.position = 'static';\n el.style.top = 'auto';\n el.style.zIndex = 'auto';\n }\n });\n }\n \n unstick();\n \n var observer = new MutationObserver(unstick);\n observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });\n})();", "Kill Sticky Headers");
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Universal Dark Mode - works on any site\n(function() {\n var enabled = true;\n \n function applyDarkMode() {\n if (!enabled) return;\n \n // Create style element if it doesn't exist\n var style = document.getElementById('universal-dark-mode-style');\n if (!style) {\n style = document.createElement('style');\n style.id = 'universal-dark-mode-style';\n document.head.appendChild(style);\n }\n \n // Dark mode CSS - inverts colors but preserves images/video\n style.textContent = '\n /* Invert everything except media */\n html {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #1a1a2e !important;\n }\n \n /* Restore images, videos, iframes, canvas */\n img, video, iframe, canvas, svg, picture, [style*=\"background-image\"] {\n filter: invert(1) hue-rotate(180deg) !important;\n }\n \n /* Preserve specific elements that should not be inverted */\n .no-dark-mode, .no-dark-mode *,\n [data-theme=\"light\"], [data-theme=\"light\"],\n .ace_editor, .ace_editor *,\n .CodeMirror, .CodeMirror *,\n .monaco-editor, .monaco-editor *,\n .markdown-body pre, .markdown-body pre *,\n .highlight, .highlight *,\n pre code, pre code * {\n filter: none !important;\n }\n \n /* Fix common UI elements */\n .modal, .popup, .dropdown-menu, .tooltip, .popover {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #2d2d44 !important;\n border-color: #444 !important;\n }\n \n /* Scrollbars */\n ::-webkit-scrollbar { background: #1a1a2e !important; }\n ::-webkit-scrollbar-thumb { background: #444 !important; }\n ::-webkit-scrollbar-thumb:hover { background: #555 !important; }\n \n /* Selection */\n ::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ';\n }\n \n function removeDarkMode() {\n var style = document.getElementById('universal-dark-mode-style');\n if (style) style.remove();\n }\n \n // Toggle with Alt+Shift+D\n document.addEventListener('keydown', function(e) {\n if (e.altKey && e.shiftKey && e.key === 'D') {\n e.preventDefault();\n enabled = !enabled;\n if (enabled) {\n applyDarkMode();\n console.log('[Universal Dark Mode] Enabled');\n } else {\n removeDarkMode();\n console.log('[Universal Dark Mode] Disabled');\n }\n }\n });\n \n // Apply on load\n applyDarkMode();\n \n // Re-apply on dynamic content\n var observer = new MutationObserver(function(mutations) {\n if (enabled && !document.getElementById('universal-dark-mode-style')) {\n applyDarkMode();\n }\n });\n observer.observe(document.head, { childList: true });\n \n console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');\n})();", "Universal Dark Mode");
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();

Skip to content
[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM#20752

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

meta-codesync
merged 9 commits into

from

Jul 9, 2026
 Merged


[ExecuTorch][WebGPU] Add f16-multiply variant of the steel q4gsw prefill GEMM
#20752
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


9 commits


Select commit
Hold shift + click to select a range
161baf5

Update


JCNTH Jul 6, 2026
d1ad151

Update


JCNTH Jul 6, 2026
4c20d00

Update


JCNTH Jul 7, 2026
642b82a

Update


JCNTH Jul 7, 2026
b8be2d1

Update


JCNTH Jul 9, 2026
523e082

Update


JCNTH Jul 9, 2026
c522131

Update


JCNTH Jul 9, 2026
9a0c964

Update


JCNTH Jul 9, 2026
061b444

Update


JCNTH Jul 9, 2026

File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


10 changes: 10 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -6,11 +6,13 @@
 * LICENSE file in the root directory of this source tree.
 */

#include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
#include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
#include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
#include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>

Expand DownExpand Up@@ -263,6 +265,14 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
 : use_steel ? kQ4gswLinearGemmSteelWGSL
 : use_shmem_gemm ? kQ4gswLinearGemmShmemWGSL
 : kQ4gswLinearWGSL;
 // f16-multiply steel: only when the device negotiated shader-f16; else the
 // f32 steel kernel runs (fail-closed). Same bindings and tile.
 if (use_steel) {
 const WebGPUContext* ctx = get_default_webgpu_context();
 if (ctx != nullptr && ctx->shader_f16_supported) {
 shader_src = kQ4gswLinearGemmSteelHalfWGSL;
 }
 }
 const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
 device,
 use_gemv,
Expand Down

12 changes: 10 additions & 2 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
$if DTYPE == "half":
 enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
Expand DownExpand Up@@ -70,7 +72,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 $if DTYPE == "half":
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 $else:
 dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
 }
 Bs[br * BN + bc + j] = dqv;
 }
Expand All@@ -81,7 +86,10 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 $if DTYPE == "half":
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 $else:
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
 }
 }
 workgroupBarrier();
Expand Down

2 changes: 2 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -5,5 +5,7 @@ q4gsw_linear_gemm_steel:
 DTYPE:
 - VALUE: float
 SUFFIX: ""
 - VALUE: half
 SUFFIX: half
 shader_variants:
 - NAME: q4gsw_linear_gemm_steel

126 changes: 126 additions & 0 deletions
backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
@@ -0,0 +1,126 @@
/*
 * Copyright (c) Meta Platforms, Inc. and affiliates.
 * All rights reserved.
 *
 * This source code is licensed under the BSD-style license found in the
 * LICENSE file in the root directory of this source tree.
 */

#pragma once

#include <cstdint>

namespace executorch::backends::webgpu {

// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
enable f16;
@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@group(0) @binding(1) var<storage, read> t_input: array<f32>;
@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
@group(0) @binding(4) var<storage, read> t_bias: array<f32>;

struct Params {
 M: u32,
 N: u32,
 K: u32,
 K_packed: u32,
 group_size: u32,
 padded_N: u32,
 has_bias: u32,
 _pad: u32,
}
@group(0) @binding(5) var<uniform> params: Params;

// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
// The "steel" name + register-tiled dequant-to-shared GEMM structure are
// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
// mlx/backend/metal/kernels/steel).
const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
var<workgroup> As: array<f16, 1024>; // BM*BK
var<workgroup> Bs: array<f16, 1024>; // BK*BN
@compute @workgroup_size(16, 16)
fn main(@builtin(workgroup_id) wid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>) {
 let nbN = (params.N + BN - 1u) / BN;
 let bx = wid.x % nbN; // decode 2D tile id from 1D dispatch
 let by = wid.x / nbN;
 let row0 = by * BM;
 let col0 = bx * BN;
 let tid = lid.y * 16u + lid.x;
 var acc: array<array<f32, 4>, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
 }
 // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
 let ar = tid / 4u; // 0..63 (row in tile)
 let ac = (tid % 4u) * 4u; // 0,4,8,12 (K offset, 4 contiguous)
 // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
 let br = tid / 16u; // 0..15 (K within BK)
 let bc = (tid % 16u) * 4u; // 0,4,..60 (N offset, 4 contiguous)

 var k0: u32 = 0u;
 loop {
 if (k0 >= params.K) { break; }
 // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
 let arow = row0 + ar;
 if (arow < params.M) {
 let base = arow * params.K + k0 + ac;
 As[ar * BK + ac + 0u] = f16(t_input[base]);
 As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
 As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
 As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
 } else {
 As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
 As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
 }
 // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
 let kk = k0 + br; // K index for this shmem row
 let scale_row = (kk / params.group_size) * params.padded_N;
 for (var j: u32 = 0u; j < 4u; j = j + 1u) {
 let n = col0 + bc + j;
 var dqv: f16 = 0.0;
 if (n < params.N) {
 let byte_idx = n * params.K_packed + (kk >> 1u);
 let word = t_weight[byte_idx >> 2u];
 let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
 var nib: u32;
 if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
 dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
 }
 Bs[br * BN + bc + j] = dqv;
 }
 workgroupBarrier();
 for (var k: u32 = 0u; k < BK; k = k + 1u) {
 var a: array<f16, 4>;
 var bvec: array<f16, 4>;
 for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
 }
 }
 workgroupBarrier();
 k0 = k0 + BK;
 }
 for (var m: u32 = 0u; m < 4u; m = m + 1u) {
 for (var n: u32 = 0u; n < 4u; n = n + 1u) {
 let r = row0 + lid.y * 4u + m;
 let c = col0 + lid.x * 4u + n;
 if (r < params.M && c < params.N) {
 var v = acc[m][n];
 if (params.has_bias != 0u) { v = v + t_bias[c]; }
 t_out[r * params.N + c] = v;
 }
 }
 }
}
)";

inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeX = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeY = 16;
inline constexpr uint32_t kQ4gswLinearGemmSteelHalfWorkgroupSizeZ = 1;

} // namespace executorch::backends::webgpu
Loading