diff --git a/backends/webgpu/CMakeLists.txt b/backends/webgpu/CMakeLists.txt
index adbf4301413..93e81192e80 100644
--- a/backends/webgpu/CMakeLists.txt
+++ b/backends/webgpu/CMakeLists.txt
@@ -52,6 +52,9 @@ set(WEBGPU_SRCS
     runtime/ops/cat/Cat.cpp
     runtime/ops/index/Index.cpp
     runtime/ops/sdpa_fd_decode/SdpaFdDecode.cpp
+    runtime/ops/mm/Mm.cpp
+    runtime/ops/bmm/Bmm.cpp
+    runtime/ops/linear/Linear.cpp
 )
 
 add_library(webgpu_backend ${WEBGPU_SRCS})
diff --git a/backends/webgpu/runtime/ops/bmm/Bmm.cpp b/backends/webgpu/runtime/ops/bmm/Bmm.cpp
new file mode 100644
index 00000000000..04c41a51c2e
--- /dev/null
+++ b/backends/webgpu/runtime/ops/bmm/Bmm.cpp
@@ -0,0 +1,224 @@
+/*
+ * 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.
+ */
+
+#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/bmm/bmm_tiled_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/bmm/bmm_vec4_wgsl.h>
+
+#include <webgpu/webgpu.h>
+
+#include <cstdint>
+#include <stdexcept>
+#include <vector>
+
+namespace executorch::backends::webgpu {
+
+namespace {
+
+struct BmmParams {
+  uint32_t B;
+  uint32_t M;
+  uint32_t N;
+  uint32_t K;
+};
+static_assert(sizeof(BmmParams) == 16, "BmmParams must be 16 bytes");
+
+constexpr uint32_t kTile = 32u;
+
+// Batched shared-memory tiled GEMM.
+void bmm_impl(WebGPUGraph& graph, const std::vector<int>& args) {
+  const int a_id = args.at(0);
+  const int b_id = args.at(1);
+  const int out_id = args.at(2);
+
+  WGPUDevice device = graph.device();
+
+  const auto& a = graph.get_tensor(a_id);
+  const auto& b = graph.get_tensor(b_id);
+  const auto& out = graph.get_tensor(out_id);
+
+  if (a.dims.size() != 3 || b.dims.size() != 3) {
+    throw std::runtime_error("WebGPU bmm: inputs must be 3D");
+  }
+  const uint32_t B = static_cast<uint32_t>(a.dims[0]);
+  const uint32_t M = static_cast<uint32_t>(a.dims[1]);
+  const uint32_t K = static_cast<uint32_t>(a.dims[2]);
+  const uint32_t N = static_cast<uint32_t>(b.dims[2]);
+  if (static_cast<uint32_t>(b.dims[0]) != B ||
+      static_cast<uint32_t>(b.dims[1]) != K) {
+    throw std::runtime_error("WebGPU bmm: batch/K mismatch between a and b");
+  }
+  if (B == 0 || M == 0 || N == 0 || K == 0) {
+    throw std::runtime_error("WebGPU bmm: B, M, N, or K == 0");
+  }
+
+  const uint64_t outputs =
+      static_cast<uint64_t>(B) * static_cast<uint64_t>(M) * N;
+  if (a.nbytes != static_cast<uint64_t>(B) * M * K * sizeof(float) ||
+      b.nbytes != static_cast<uint64_t>(B) * K * N * sizeof(float) ||
+      out.nbytes != outputs * sizeof(float)) {
+    throw std::runtime_error("WebGPU bmm: fp32-only (byte-size mismatch)");
+  }
+
+  const uint32_t max_wg = utils::queried_max_workgroups(device);
+  const uint32_t dispatch_x = (N + kTile - 1u) / kTile;
+  // uint64 so B * ceil(M/32) can't wrap before the limit check.
+  const uint64_t dispatch_y64 =
+      static_cast<uint64_t>(B) * ((M + kTile - 1u) / kTile);
+  if (dispatch_x > max_wg || dispatch_y64 > max_wg) {
+    throw std::runtime_error("WebGPU bmm: tile grid exceeds dispatch limit");
+  }
+  const uint32_t dispatch_y = static_cast<uint32_t>(dispatch_y64);
+
+  BmmParams params = {};
+  params.B = B;
+  params.M = M;
+  params.N = N;
+  params.K = K;
+
+  WGPUBuffer uniform_buffer =
+      utils::make_uniform(device, &params, sizeof(BmmParams));
+  graph.add_uniform_buffer_bytes(sizeof(BmmParams));
+
+  // vec4 path when K and N are multiples of 4 (wider 16B loads); else scalar.
+  const bool use_vec4 = (K % 4u == 0u) && (N % 4u == 0u);
+  WGPUShaderSourceWGSL wgsl_desc = {};
+  wgsl_desc.chain.sType = WGPUSType_ShaderSourceWGSL;
+  wgsl_desc.code = {use_vec4 ? kBmmVec4WGSL : kBmmTiledWGSL, WGPU_STRLEN};
+  WGPUShaderModuleDescriptor shader_desc = {};
+  shader_desc.nextInChain = &wgsl_desc.chain;
+  WGPUShaderModule shader = wgpuDeviceCreateShaderModule(device, &shader_desc);
+
+  WGPUBindGroupLayoutEntry entries[4] = {};
+  entries[0].binding = 0;
+  entries[0].visibility = WGPUShaderStage_Compute;
+  entries[0].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[1].binding = 1;
+  entries[1].visibility = WGPUShaderStage_Compute;
+  entries[1].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[2].binding = 2;
+  entries[2].visibility = WGPUShaderStage_Compute;
+  entries[2].buffer.type = WGPUBufferBindingType_Storage;
+  entries[3].binding = 3;
+  entries[3].visibility = WGPUShaderStage_Compute;
+  entries[3].buffer.type = WGPUBufferBindingType_Uniform;
+
+  WGPUBindGroupLayoutDescriptor bgl_desc = {};
+  bgl_desc.entryCount = 4;
+  bgl_desc.entries = entries;
+  WGPUBindGroupLayout bgl = wgpuDeviceCreateBindGroupLayout(device, &bgl_desc);
+
+  WGPUPipelineLayoutDescriptor pl_desc = {};
+  pl_desc.bindGroupLayoutCount = 1;
+  pl_desc.bindGroupLayouts = &bgl;
+  WGPUPipelineLayout pipeline_layout =
+      wgpuDeviceCreatePipelineLayout(device, &pl_desc);
+
+  // Tiled kernel has a fixed @workgroup_size(8, 8, 1) — no override constant.
+  WGPUComputePipelineDescriptor pipeline_desc = {};
+  pipeline_desc.layout = pipeline_layout;
+  pipeline_desc.compute.module = shader;
+  pipeline_desc.compute.entryPoint = {"main", WGPU_STRLEN};
+  WGPUComputePipeline pipeline =
+      wgpuDeviceCreateComputePipeline(device, &pipeline_desc);
+
+  WGPUBindGroupEntry bg_entries[4] = {};
+  bg_entries[0].binding = 0;
+  bg_entries[0].buffer = a.buffer;
+  bg_entries[0].size = a.nbytes;
+  bg_entries[1].binding = 1;
+  bg_entries[1].buffer = b.buffer;
+  bg_entries[1].size = b.nbytes;
+  bg_entries[2].binding = 2;
+  bg_entries[2].buffer = out.buffer;
+  bg_entries[2].size = out.nbytes;
+  bg_entries[3].binding = 3;
+  bg_entries[3].buffer = uniform_buffer;
+  bg_entries[3].size = sizeof(BmmParams);
+
+  WGPUBindGroupDescriptor bg_desc = {};
+  bg_desc.layout = bgl;
+  bg_desc.entryCount = 4;
+  bg_desc.entries = bg_entries;
+  WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup(device, &bg_desc);
+
+  WebGPUDispatch dispatch;
+  dispatch.pipeline = pipeline;
+  dispatch.bind_group = bind_group;
+  dispatch.workgroup_count_x = dispatch_x;
+  dispatch.workgroup_count_y = dispatch_y;
+  dispatch.kernel_name = "bmm";
+  const size_t dispatch_idx = graph.add_dispatch(dispatch);
+
+  // Dynamic shapes: re-derive B, M (from a) and N (from b) + the tile grid.
+  WGPUBuffer params_buf = uniform_buffer;
+  graph.add_tensor_resize_hook(
+      a_id,
+      [a_id, b_id, out_id, B, M, N, K, use_vec4, dispatch_idx, params_buf](
+          WebGPUGraph& g) {
+        const auto& da = g.cur_dims(a_id);
+        const auto& db = g.cur_dims(b_id);
+        if (da.size() != 3 || db.size() != 3 ||
+            static_cast<uint32_t>(da[2]) != K ||
+            static_cast<uint32_t>(db[1]) != K ||
+            static_cast<uint32_t>(db[0]) != static_cast<uint32_t>(da[0])) {
+          throw std::runtime_error("WebGPU bmm: live input dims invalid");
+        }
+        const uint32_t live_b = static_cast<uint32_t>(da[0]);
+        const uint32_t live_m = static_cast<uint32_t>(da[1]);
+        const uint32_t live_n = static_cast<uint32_t>(db[2]);
+        if (live_b == 0u || live_m == 0u || live_n == 0u) {
+          throw std::runtime_error("WebGPU bmm: live B, M, or N == 0");
+        }
+        if (live_b > B || live_m > M || live_n > N) {
+          throw std::runtime_error(
+              "WebGPU bmm: live dims exceed build-time max");
+        }
+        // vec4 baked for N%4==0; a live N breaking it reads past the row.
+        if (use_vec4 && live_n % 4u != 0u) {
+          throw std::runtime_error("WebGPU bmm: live N breaks vec4 alignment");
+        }
+        BmmParams p = {};
+        p.B = live_b;
+        p.M = live_m;
+        p.N = live_n;
+        p.K = K;
+        wgpuQueueWriteBuffer(g.queue(), params_buf, 0, &p, sizeof(p));
+        const uint32_t max_wg = utils::queried_max_workgroups(g.device());
+        const uint32_t dx = (live_n + kTile - 1u) / kTile;
+        const uint64_t dy64 =
+            static_cast<uint64_t>(live_b) * ((live_m + kTile - 1u) / kTile);
+        if (dx > max_wg || dy64 > max_wg) {
+          throw std::runtime_error(
+              "WebGPU bmm(resize): tile grid exceeds dispatch limit");
+        }
+        g.dispatch_at(dispatch_idx).workgroup_count_x = dx;
+        g.dispatch_at(dispatch_idx).workgroup_count_y =
+            static_cast<uint32_t>(dy64);
+        g.set_cur_dims(
+            out_id,
+            {static_cast<int64_t>(live_b),
+             static_cast<int64_t>(live_m),
+             static_cast<int64_t>(live_n)});
+      });
+
+  wgpuShaderModuleRelease(shader);
+  wgpuBindGroupLayoutRelease(bgl);
+  wgpuPipelineLayoutRelease(pipeline_layout);
+  graph.own_uniform_buffer(uniform_buffer);
+}
+
+} // namespace
+
+WEBGPU_REGISTER_OPERATORS {
+  WEBGPU_REGISTER_OP(aten.bmm.default, bmm_impl);
+}
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/bmm/bmm_tiled.wgsl b/backends/webgpu/runtime/ops/bmm/bmm_tiled.wgsl
new file mode 100644
index 00000000000..e7f02225a6d
--- /dev/null
+++ b/backends/webgpu/runtime/ops/bmm/bmm_tiled.wgsl
@@ -0,0 +1,89 @@
+struct Params {
+  B: u32,
+  M: u32,
+  N: u32,
+  K: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(base: u32, row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return a[base + row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(base: u32, krow: u32, col: u32) -> f32 {
+  if (krow < params.K && col < params.N) {
+    return b[base + krow * params.N + col];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tiles_per_m = (params.M + TILE - 1u) / TILE;
+  let bn = wg_id.y / tiles_per_m;
+  let row_tile = wg_id.y % tiles_per_m;
+  let tile_row0 = row_tile * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  let a_base = bn * params.M * params.K;
+  let b_base = bn * params.K * params.N;
+  let out_base = bn * params.M * params.N;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(a_base, tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(b_base, k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[out_base + r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/bmm/bmm_tiled_wgsl.h b/backends/webgpu/runtime/ops/bmm/bmm_tiled_wgsl.h
new file mode 100644
index 00000000000..944a8fc3d9c
--- /dev/null
+++ b/backends/webgpu/runtime/ops/bmm/bmm_tiled_wgsl.h
@@ -0,0 +1,113 @@
+/*
+ * 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 bmm_tiled.wgsl - DO NOT EDIT.
+// wgsl-sha256: cb1dad9102604a87e4b4525e4452bc74dc2d9d258ded282734f2307fed38047b
+inline constexpr const char* kBmmTiledWGSL = R"(
+struct Params {
+  B: u32,
+  M: u32,
+  N: u32,
+  K: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(base: u32, row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return a[base + row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(base: u32, krow: u32, col: u32) -> f32 {
+  if (krow < params.K && col < params.N) {
+    return b[base + krow * params.N + col];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tiles_per_m = (params.M + TILE - 1u) / TILE;
+  let bn = wg_id.y / tiles_per_m;
+  let row_tile = wg_id.y % tiles_per_m;
+  let tile_row0 = row_tile * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  let a_base = bn * params.M * params.K;
+  let b_base = bn * params.K * params.N;
+  let out_base = bn * params.M * params.N;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(a_base, tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(b_base, k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[out_base + r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kBmmTiledWorkgroupSizeX = 8;
+inline constexpr uint32_t kBmmTiledWorkgroupSizeY = 8;
+inline constexpr uint32_t kBmmTiledWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/bmm/bmm_vec4.wgsl b/backends/webgpu/runtime/ops/bmm/bmm_vec4.wgsl
new file mode 100644
index 00000000000..b559d8aa1db
--- /dev/null
+++ b/backends/webgpu/runtime/ops/bmm/bmm_vec4.wgsl
@@ -0,0 +1,84 @@
+struct Params {
+  B: u32,
+  M: u32,
+  N: u32,
+  K: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> b: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<vec4<f32>>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const TILE4: u32 = 8u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tiles_per_m = (params.M + TILE - 1u) / TILE;
+  let bn = wg_id.y / tiles_per_m;
+  let row_tile = wg_id.y % tiles_per_m;
+  let tile_row0 = row_tile * TILE;
+  let tile_col0_4 = wg_id.x * TILE4;
+  let tile_row = local_id.y * RPT;
+  let tile_col4 = local_id.x;
+
+  let k4 = params.K / 4u;
+  let n4 = params.N / 4u;
+  let a_base = bn * params.M * k4;
+  let b_base = bn * params.K * n4;
+  let out_base = bn * params.M * n4;
+
+  var acc: array<vec4<f32>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    acc[ir] = vec4<f32>(0.0);
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4_start = t * TILE4;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      let akcol = k4_start + local_id.x;
+      if (tile_row0 + arow < params.M && akcol < k4) {
+        a_sub[arow][local_id.x] = a[a_base + (tile_row0 + arow) * k4 + akcol];
+      } else {
+        a_sub[arow][local_id.x] = vec4<f32>(0.0);
+      }
+      let brow = t * TILE + arow;
+      let bncol = tile_col0_4 + local_id.x;
+      if (brow < params.K && bncol < n4) {
+        b_sub[arow][local_id.x] = b[b_base + brow * n4 + bncol];
+      } else {
+        b_sub[arow][local_id.x] = vec4<f32>(0.0);
+      }
+    }
+    workgroupBarrier();
+
+    for (var kk: u32 = 0u; kk < TILE4; kk = kk + 1u) {
+      let b0 = b_sub[kk * 4u + 0u][tile_col4];
+      let b1 = b_sub[kk * 4u + 1u][tile_col4];
+      let b2 = b_sub[kk * 4u + 2u][tile_col4];
+      let b3 = b_sub[kk * 4u + 3u][tile_col4];
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let ac = a_sub[tile_row + ir][kk];
+        acc[ir] = acc[ir] + b0 * ac.x + b1 * ac.y + b2 * ac.z + b3 * ac.w;
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    let r = tile_row0 + tile_row + ir;
+    let c4 = tile_col0_4 + tile_col4;
+    if (r < params.M && c4 < n4) {
+      out[out_base + r * n4 + c4] = acc[ir];
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/bmm/bmm_vec4_wgsl.h b/backends/webgpu/runtime/ops/bmm/bmm_vec4_wgsl.h
new file mode 100644
index 00000000000..7616f8e2a8a
--- /dev/null
+++ b/backends/webgpu/runtime/ops/bmm/bmm_vec4_wgsl.h
@@ -0,0 +1,108 @@
+/*
+ * 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 bmm_vec4.wgsl - DO NOT EDIT.
+// wgsl-sha256: 32bd94bcb0c0758b60267224f737bac142aec0e6789f7f2d6433cd4250cfa7b5
+inline constexpr const char* kBmmVec4WGSL = R"(
+struct Params {
+  B: u32,
+  M: u32,
+  N: u32,
+  K: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> b: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<vec4<f32>>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const TILE4: u32 = 8u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tiles_per_m = (params.M + TILE - 1u) / TILE;
+  let bn = wg_id.y / tiles_per_m;
+  let row_tile = wg_id.y % tiles_per_m;
+  let tile_row0 = row_tile * TILE;
+  let tile_col0_4 = wg_id.x * TILE4;
+  let tile_row = local_id.y * RPT;
+  let tile_col4 = local_id.x;
+
+  let k4 = params.K / 4u;
+  let n4 = params.N / 4u;
+  let a_base = bn * params.M * k4;
+  let b_base = bn * params.K * n4;
+  let out_base = bn * params.M * n4;
+
+  var acc: array<vec4<f32>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    acc[ir] = vec4<f32>(0.0);
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4_start = t * TILE4;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      let akcol = k4_start + local_id.x;
+      if (tile_row0 + arow < params.M && akcol < k4) {
+        a_sub[arow][local_id.x] = a[a_base + (tile_row0 + arow) * k4 + akcol];
+      } else {
+        a_sub[arow][local_id.x] = vec4<f32>(0.0);
+      }
+      let brow = t * TILE + arow;
+      let bncol = tile_col0_4 + local_id.x;
+      if (brow < params.K && bncol < n4) {
+        b_sub[arow][local_id.x] = b[b_base + brow * n4 + bncol];
+      } else {
+        b_sub[arow][local_id.x] = vec4<f32>(0.0);
+      }
+    }
+    workgroupBarrier();
+
+    for (var kk: u32 = 0u; kk < TILE4; kk = kk + 1u) {
+      let b0 = b_sub[kk * 4u + 0u][tile_col4];
+      let b1 = b_sub[kk * 4u + 1u][tile_col4];
+      let b2 = b_sub[kk * 4u + 2u][tile_col4];
+      let b3 = b_sub[kk * 4u + 3u][tile_col4];
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let ac = a_sub[tile_row + ir][kk];
+        acc[ir] = acc[ir] + b0 * ac.x + b1 * ac.y + b2 * ac.z + b3 * ac.w;
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    let r = tile_row0 + tile_row + ir;
+    let c4 = tile_col0_4 + tile_col4;
+    if (r < params.M && c4 < n4) {
+      out[out_base + r * n4 + c4] = acc[ir];
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kBmmVec4WorkgroupSizeX = 8;
+inline constexpr uint32_t kBmmVec4WorkgroupSizeY = 8;
+inline constexpr uint32_t kBmmVec4WorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/linear/Linear.cpp b/backends/webgpu/runtime/ops/linear/Linear.cpp
new file mode 100644
index 00000000000..7b0de24918d
--- /dev/null
+++ b/backends/webgpu/runtime/ops/linear/Linear.cpp
@@ -0,0 +1,202 @@
+/*
+ * 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.
+ */
+
+#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/linear/linear_tiled_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/linear/linear_vec4_wgsl.h>
+
+#include <webgpu/webgpu.h>
+
+#include <cstdint>
+#include <cstring>
+#include <stdexcept>
+#include <vector>
+
+namespace executorch::backends::webgpu {
+
+namespace {
+
+struct LinearParams {
+  uint32_t M;
+  uint32_t N;
+  uint32_t K;
+  uint32_t pad_;
+};
+static_assert(sizeof(LinearParams) == 16, "LinearParams must be 16 bytes");
+
+constexpr uint32_t kTile = 32u;
+
+// aten.linear (no bias); shared-memory tiled GEMM.
+void linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
+  // out is the last arg; bias (if present) is unused on this path.
+  const int in_id = args.at(0);
+  const int w_id = args.at(1);
+  const int out_id = args.at(args.size() - 1);
+
+  WGPUDevice device = graph.device();
+  const auto& in = graph.get_tensor(in_id);
+  const auto& w = graph.get_tensor(w_id);
+  const auto& out = graph.get_tensor(out_id);
+
+  if (in.dims.size() != 2 || w.dims.size() != 2) {
+    throw std::runtime_error("WebGPU linear: input/weight must be 2D");
+  }
+  const uint32_t M = static_cast<uint32_t>(in.dims[0]);
+  const uint32_t K = static_cast<uint32_t>(in.dims[1]);
+  const uint32_t N = static_cast<uint32_t>(w.dims[0]);
+  if (static_cast<uint32_t>(w.dims[1]) != K) {
+    throw std::runtime_error("WebGPU linear: weight.dims[1] != input.dims[1]");
+  }
+  if (M == 0 || N == 0 || K == 0) {
+    throw std::runtime_error("WebGPU linear: M, N, or K == 0");
+  }
+
+  const uint64_t outputs = static_cast<uint64_t>(M) * static_cast<uint64_t>(N);
+  if (in.nbytes != static_cast<uint64_t>(M) * K * sizeof(float) ||
+      w.nbytes != static_cast<uint64_t>(N) * K * sizeof(float) ||
+      out.nbytes != outputs * sizeof(float)) {
+    throw std::runtime_error("WebGPU linear: fp32-only (byte-size mismatch)");
+  }
+
+  const uint32_t dispatch_x = (N + kTile - 1u) / kTile;
+  const uint32_t dispatch_y = (M + kTile - 1u) / kTile;
+  const uint32_t max_wg = utils::queried_max_workgroups(device);
+  if (dispatch_x > max_wg || dispatch_y > max_wg) {
+    throw std::runtime_error("WebGPU linear: tile grid exceeds dispatch limit");
+  }
+
+  LinearParams params = {};
+  params.M = M;
+  params.N = N;
+  params.K = K;
+
+  WGPUBufferDescriptor uniform_desc = {};
+  uniform_desc.size = sizeof(LinearParams);
+  uniform_desc.usage = WGPUBufferUsage_Uniform | WGPUBufferUsage_CopyDst;
+  uniform_desc.mappedAtCreation = true;
+  WGPUBuffer uniform_buffer = wgpuDeviceCreateBuffer(device, &uniform_desc);
+  void* mapped =
+      wgpuBufferGetMappedRange(uniform_buffer, 0, sizeof(LinearParams));
+  std::memcpy(mapped, &params, sizeof(LinearParams));
+  wgpuBufferUnmap(uniform_buffer);
+  graph.add_uniform_buffer_bytes(sizeof(LinearParams));
+
+  // vec4-over-K path when K%4==0 (scalar out, only K%4, not N%4 like mm).
+  const bool use_vec4 = (K % 4u == 0u);
+  WGPUShaderSourceWGSL wgsl_desc = {};
+  wgsl_desc.chain.sType = WGPUSType_ShaderSourceWGSL;
+  wgsl_desc.code = {use_vec4 ? kLinearVec4WGSL : kLinearTiledWGSL, WGPU_STRLEN};
+  WGPUShaderModuleDescriptor shader_desc = {};
+  shader_desc.nextInChain = &wgsl_desc.chain;
+  WGPUShaderModule shader = wgpuDeviceCreateShaderModule(device, &shader_desc);
+
+  WGPUBindGroupLayoutEntry entries[4] = {};
+  entries[0].binding = 0;
+  entries[0].visibility = WGPUShaderStage_Compute;
+  entries[0].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[1].binding = 1;
+  entries[1].visibility = WGPUShaderStage_Compute;
+  entries[1].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[2].binding = 2;
+  entries[2].visibility = WGPUShaderStage_Compute;
+  entries[2].buffer.type = WGPUBufferBindingType_Storage;
+  entries[3].binding = 3;
+  entries[3].visibility = WGPUShaderStage_Compute;
+  entries[3].buffer.type = WGPUBufferBindingType_Uniform;
+
+  WGPUBindGroupLayoutDescriptor bgl_desc = {};
+  bgl_desc.entryCount = 4;
+  bgl_desc.entries = entries;
+  WGPUBindGroupLayout bgl = wgpuDeviceCreateBindGroupLayout(device, &bgl_desc);
+
+  WGPUPipelineLayoutDescriptor pl_desc = {};
+  pl_desc.bindGroupLayoutCount = 1;
+  pl_desc.bindGroupLayouts = &bgl;
+  WGPUPipelineLayout pipeline_layout =
+      wgpuDeviceCreatePipelineLayout(device, &pl_desc);
+
+  // Tiled kernel has a fixed @workgroup_size(8, 8, 1) — no override constant.
+  WGPUComputePipelineDescriptor pipeline_desc = {};
+  pipeline_desc.layout = pipeline_layout;
+  pipeline_desc.compute.module = shader;
+  pipeline_desc.compute.entryPoint = {"main", WGPU_STRLEN};
+  WGPUComputePipeline pipeline =
+      wgpuDeviceCreateComputePipeline(device, &pipeline_desc);
+
+  WGPUBindGroupEntry bg_entries[4] = {};
+  bg_entries[0].binding = 0;
+  bg_entries[0].buffer = in.buffer;
+  bg_entries[0].size = in.nbytes;
+  bg_entries[1].binding = 1;
+  bg_entries[1].buffer = w.buffer;
+  bg_entries[1].size = w.nbytes;
+  bg_entries[2].binding = 2;
+  bg_entries[2].buffer = out.buffer;
+  bg_entries[2].size = out.nbytes;
+  bg_entries[3].binding = 3;
+  bg_entries[3].buffer = uniform_buffer;
+  bg_entries[3].size = sizeof(LinearParams);
+
+  WGPUBindGroupDescriptor bg_desc = {};
+  bg_desc.layout = bgl;
+  bg_desc.entryCount = 4;
+  bg_desc.entries = bg_entries;
+  WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup(device, &bg_desc);
+
+  WebGPUDispatch dispatch;
+  dispatch.pipeline = pipeline;
+  dispatch.bind_group = bind_group;
+  dispatch.workgroup_count_x = dispatch_x;
+  dispatch.workgroup_count_y = dispatch_y;
+  dispatch.kernel_name = "linear";
+  const size_t dispatch_idx = graph.add_dispatch(dispatch);
+
+  WGPUBuffer params_buf = uniform_buffer;
+  graph.add_tensor_resize_hook(
+      in_id,
+      [in_id, out_id, M, N, K, dispatch_x, dispatch_idx, params_buf](
+          WebGPUGraph& g) {
+        const auto& d = g.cur_dims(in_id);
+        const uint64_t numel = utils::numel_of(d);
+        if (numel % static_cast<uint64_t>(K) != 0u) {
+          throw std::runtime_error(
+              "WebGPU linear: live input numel not a multiple of K");
+        }
+        const uint32_t m =
+            static_cast<uint32_t>(numel / static_cast<uint64_t>(K));
+        if (m == 0u || m > M) {
+          throw std::runtime_error(
+              "WebGPU linear: live M is 0 or exceeds the build-time max");
+        }
+        LinearParams p = {};
+        p.M = m;
+        p.N = N;
+        p.K = K;
+        wgpuQueueWriteBuffer(g.queue(), params_buf, 0, &p, sizeof(p));
+        g.dispatch_at(dispatch_idx).workgroup_count_x = dispatch_x;
+        g.dispatch_at(dispatch_idx).workgroup_count_y =
+            (m + kTile - 1u) / kTile;
+        g.set_cur_dims(
+            out_id, {static_cast<int64_t>(m), static_cast<int64_t>(N)});
+      });
+
+  wgpuShaderModuleRelease(shader);
+  wgpuBindGroupLayoutRelease(bgl);
+  wgpuPipelineLayoutRelease(pipeline_layout);
+  graph.own_uniform_buffer(uniform_buffer);
+}
+
+} // namespace
+
+WEBGPU_REGISTER_OPERATORS {
+  WEBGPU_REGISTER_OP(aten.linear.default, linear_impl);
+}
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/linear/linear_tiled.wgsl b/backends/webgpu/runtime/ops/linear/linear_tiled.wgsl
new file mode 100644
index 00000000000..2b94814b75e
--- /dev/null
+++ b/backends/webgpu/runtime/ops/linear/linear_tiled.wgsl
@@ -0,0 +1,82 @@
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> input: array<f32>;
+@group(0) @binding(1) var<storage, read> weight: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return input[row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(krow: u32, col: u32) -> f32 {
+  if (col < params.N && krow < params.K) {
+    return weight[col * params.K + krow];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/linear/linear_tiled_wgsl.h b/backends/webgpu/runtime/ops/linear/linear_tiled_wgsl.h
new file mode 100644
index 00000000000..880479b2301
--- /dev/null
+++ b/backends/webgpu/runtime/ops/linear/linear_tiled_wgsl.h
@@ -0,0 +1,106 @@
+/*
+ * 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 linear_tiled.wgsl - DO NOT EDIT.
+// wgsl-sha256: bb054c7ab0937a24089df4944001666c902c3c4b31e9555ec916bf357c2104f7
+inline constexpr const char* kLinearTiledWGSL = R"(
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> input: array<f32>;
+@group(0) @binding(1) var<storage, read> weight: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return input[row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(krow: u32, col: u32) -> f32 {
+  if (col < params.N && krow < params.K) {
+    return weight[col * params.K + krow];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kLinearTiledWorkgroupSizeX = 8;
+inline constexpr uint32_t kLinearTiledWorkgroupSizeY = 8;
+inline constexpr uint32_t kLinearTiledWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/linear/linear_vec4.wgsl b/backends/webgpu/runtime/ops/linear/linear_vec4.wgsl
new file mode 100644
index 00000000000..d3cfaa97d24
--- /dev/null
+++ b/backends/webgpu/runtime/ops/linear/linear_vec4.wgsl
@@ -0,0 +1,83 @@
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> input: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> weight: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const TILE4: u32 = 8u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+fn in4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.M && k4 * 4u < params.K) {
+    return input[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+fn w4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.N && k4 * 4u < params.K) {
+    return weight[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let row0 = wg_id.y * TILE;
+  let col0 = wg_id.x * TILE;
+  let tr = local_id.y * RPT;
+  let tc = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  let flat0 = local_id.y * 8u + local_id.x;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4s = t * TILE4;
+    for (var i: u32 = 0u; i < 4u; i = i + 1u) {
+      let e = flat0 + i * 64u;
+      let rr = e / TILE4;
+      let kk = e % TILE4;
+      a_sub[rr][kk] = in4(row0 + rr, k4s + kk);
+      b_sub[rr][kk] = w4(col0 + rr, k4s + kk);
+    }
+    workgroupBarrier();
+
+    for (var k4: u32 = 0u; k4 < TILE4; k4 = k4 + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let av = a_sub[tr + ir][k4];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + dot(av, b_sub[tc + ic][k4]);
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = row0 + tr + ir;
+      let c = col0 + tc + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/linear/linear_vec4_wgsl.h b/backends/webgpu/runtime/ops/linear/linear_vec4_wgsl.h
new file mode 100644
index 00000000000..361735ea63a
--- /dev/null
+++ b/backends/webgpu/runtime/ops/linear/linear_vec4_wgsl.h
@@ -0,0 +1,107 @@
+/*
+ * 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 linear_vec4.wgsl - DO NOT EDIT.
+// wgsl-sha256: 2d456961a1145cf9858237deb59fd0b37f390dfc253894b240914d2cbe056b66
+inline constexpr const char* kLinearVec4WGSL = R"(
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> input: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> weight: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const TILE4: u32 = 8u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+fn in4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.M && k4 * 4u < params.K) {
+    return input[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+fn w4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.N && k4 * 4u < params.K) {
+    return weight[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let row0 = wg_id.y * TILE;
+  let col0 = wg_id.x * TILE;
+  let tr = local_id.y * RPT;
+  let tc = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  let flat0 = local_id.y * 8u + local_id.x;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4s = t * TILE4;
+    for (var i: u32 = 0u; i < 4u; i = i + 1u) {
+      let e = flat0 + i * 64u;
+      let rr = e / TILE4;
+      let kk = e % TILE4;
+      a_sub[rr][kk] = in4(row0 + rr, k4s + kk);
+      b_sub[rr][kk] = w4(col0 + rr, k4s + kk);
+    }
+    workgroupBarrier();
+
+    for (var k4: u32 = 0u; k4 < TILE4; k4 = k4 + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let av = a_sub[tr + ir][k4];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + dot(av, b_sub[tc + ic][k4]);
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = row0 + tr + ir;
+      let c = col0 + tc + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kLinearVec4WorkgroupSizeX = 8;
+inline constexpr uint32_t kLinearVec4WorkgroupSizeY = 8;
+inline constexpr uint32_t kLinearVec4WorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/mm/Mm.cpp b/backends/webgpu/runtime/ops/mm/Mm.cpp
new file mode 100644
index 00000000000..b93c2aff247
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/Mm.cpp
@@ -0,0 +1,202 @@
+/*
+ * 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.
+ */
+
+#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/mm/mm_tiled_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/mm/mm_vec4_wgsl.h>
+
+#include <webgpu/webgpu.h>
+
+#include <cstdint>
+#include <cstring>
+#include <stdexcept>
+#include <vector>
+
+namespace executorch::backends::webgpu {
+
+namespace {
+
+struct MmParams {
+  uint32_t M;
+  uint32_t N;
+  uint32_t K;
+  uint32_t pad_;
+};
+static_assert(sizeof(MmParams) == 16, "MmParams must be 16 bytes");
+
+// 32x32 output tile per workgroup; shared-memory tiled GEMM.
+constexpr uint32_t kTile = 32u;
+
+void mm_impl(WebGPUGraph& graph, const std::vector<int>& args) {
+  const int a_id = args.at(0);
+  const int b_id = args.at(1);
+  const int out_id = args.at(2);
+
+  WGPUDevice device = graph.device();
+
+  const auto& a = graph.get_tensor(a_id);
+  const auto& b = graph.get_tensor(b_id);
+  const auto& out = graph.get_tensor(out_id);
+
+  if (a.dims.size() != 2 || b.dims.size() != 2) {
+    throw std::runtime_error("WebGPU mm: inputs must be 2D");
+  }
+  const uint32_t M = static_cast<uint32_t>(a.dims[0]);
+  const uint32_t K = static_cast<uint32_t>(a.dims[1]);
+  const uint32_t N = static_cast<uint32_t>(b.dims[1]);
+  if (static_cast<uint32_t>(b.dims[0]) != K) {
+    throw std::runtime_error("WebGPU mm: a.dims[1] != b.dims[0]");
+  }
+  if (M == 0 || N == 0 || K == 0) {
+    throw std::runtime_error("WebGPU mm: M, N, or K == 0");
+  }
+
+  const uint64_t outputs = static_cast<uint64_t>(M) * static_cast<uint64_t>(N);
+  if (a.nbytes != static_cast<uint64_t>(M) * K * sizeof(float) ||
+      b.nbytes != static_cast<uint64_t>(K) * N * sizeof(float) ||
+      out.nbytes != outputs * sizeof(float)) {
+    throw std::runtime_error("WebGPU mm: fp32-only (byte-size mismatch)");
+  }
+
+  const uint32_t dispatch_x = (N + kTile - 1u) / kTile;
+  const uint32_t dispatch_y = (M + kTile - 1u) / kTile;
+  const uint32_t max_wg = utils::queried_max_workgroups(device);
+  if (dispatch_x > max_wg || dispatch_y > max_wg) {
+    throw std::runtime_error("WebGPU mm: tile grid exceeds dispatch limit");
+  }
+
+  MmParams params = {};
+  params.M = M;
+  params.N = N;
+  params.K = K;
+
+  WGPUBufferDescriptor uniform_desc = {};
+  uniform_desc.size = sizeof(MmParams);
+  uniform_desc.usage = WGPUBufferUsage_Uniform | WGPUBufferUsage_CopyDst;
+  uniform_desc.mappedAtCreation = true;
+  WGPUBuffer uniform_buffer = wgpuDeviceCreateBuffer(device, &uniform_desc);
+  void* mapped = wgpuBufferGetMappedRange(uniform_buffer, 0, sizeof(MmParams));
+  std::memcpy(mapped, &params, sizeof(MmParams));
+  wgpuBufferUnmap(uniform_buffer);
+  graph.add_uniform_buffer_bytes(sizeof(MmParams));
+
+  // vec4 path when K and N are multiples of 4 (wider 16B loads); else scalar.
+  const bool use_vec4 = (K % 4u == 0u) && (N % 4u == 0u);
+  WGPUShaderSourceWGSL wgsl_desc = {};
+  wgsl_desc.chain.sType = WGPUSType_ShaderSourceWGSL;
+  wgsl_desc.code = {use_vec4 ? kMmVec4WGSL : kMmTiledWGSL, WGPU_STRLEN};
+  WGPUShaderModuleDescriptor shader_desc = {};
+  shader_desc.nextInChain = &wgsl_desc.chain;
+  WGPUShaderModule shader = wgpuDeviceCreateShaderModule(device, &shader_desc);
+
+  WGPUBindGroupLayoutEntry entries[4] = {};
+  entries[0].binding = 0;
+  entries[0].visibility = WGPUShaderStage_Compute;
+  entries[0].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[1].binding = 1;
+  entries[1].visibility = WGPUShaderStage_Compute;
+  entries[1].buffer.type = WGPUBufferBindingType_ReadOnlyStorage;
+  entries[2].binding = 2;
+  entries[2].visibility = WGPUShaderStage_Compute;
+  entries[2].buffer.type = WGPUBufferBindingType_Storage;
+  entries[3].binding = 3;
+  entries[3].visibility = WGPUShaderStage_Compute;
+  entries[3].buffer.type = WGPUBufferBindingType_Uniform;
+
+  WGPUBindGroupLayoutDescriptor bgl_desc = {};
+  bgl_desc.entryCount = 4;
+  bgl_desc.entries = entries;
+  WGPUBindGroupLayout bgl = wgpuDeviceCreateBindGroupLayout(device, &bgl_desc);
+
+  WGPUPipelineLayoutDescriptor pl_desc = {};
+  pl_desc.bindGroupLayoutCount = 1;
+  pl_desc.bindGroupLayouts = &bgl;
+  WGPUPipelineLayout pipeline_layout =
+      wgpuDeviceCreatePipelineLayout(device, &pl_desc);
+
+  // Tiled kernel has a fixed @workgroup_size(8, 8, 1) — no override constant.
+  WGPUComputePipelineDescriptor pipeline_desc = {};
+  pipeline_desc.layout = pipeline_layout;
+  pipeline_desc.compute.module = shader;
+  pipeline_desc.compute.entryPoint = {"main", WGPU_STRLEN};
+  WGPUComputePipeline pipeline =
+      wgpuDeviceCreateComputePipeline(device, &pipeline_desc);
+
+  WGPUBindGroupEntry bg_entries[4] = {};
+  bg_entries[0].binding = 0;
+  bg_entries[0].buffer = a.buffer;
+  bg_entries[0].size = a.nbytes;
+  bg_entries[1].binding = 1;
+  bg_entries[1].buffer = b.buffer;
+  bg_entries[1].size = b.nbytes;
+  bg_entries[2].binding = 2;
+  bg_entries[2].buffer = out.buffer;
+  bg_entries[2].size = out.nbytes;
+  bg_entries[3].binding = 3;
+  bg_entries[3].buffer = uniform_buffer;
+  bg_entries[3].size = sizeof(MmParams);
+
+  WGPUBindGroupDescriptor bg_desc = {};
+  bg_desc.layout = bgl;
+  bg_desc.entryCount = 4;
+  bg_desc.entries = bg_entries;
+  WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup(device, &bg_desc);
+
+  WebGPUDispatch dispatch;
+  dispatch.pipeline = pipeline;
+  dispatch.bind_group = bind_group;
+  dispatch.workgroup_count_x = dispatch_x;
+  dispatch.workgroup_count_y = dispatch_y;
+  dispatch.kernel_name = "mm";
+  const size_t dispatch_idx = graph.add_dispatch(dispatch);
+
+  // Dynamic shapes: recompute the live M (leading dim) + the y tile count.
+  WGPUBuffer params_buf = uniform_buffer;
+  graph.add_tensor_resize_hook(
+      a_id,
+      [a_id, out_id, M, N, K, dispatch_x, dispatch_idx, params_buf](
+          WebGPUGraph& g) {
+        const auto& d = g.cur_dims(a_id);
+        const uint64_t numel = utils::numel_of(d);
+        if (numel % static_cast<uint64_t>(K) != 0u) {
+          throw std::runtime_error(
+              "WebGPU mm: live input numel not a multiple of K");
+        }
+        const uint32_t m =
+            static_cast<uint32_t>(numel / static_cast<uint64_t>(K));
+        if (m == 0u || m > M) {
+          throw std::runtime_error(
+              "WebGPU mm: live M is 0 or exceeds the build-time max");
+        }
+        MmParams p = {};
+        p.M = m;
+        p.N = N;
+        p.K = K;
+        wgpuQueueWriteBuffer(g.queue(), params_buf, 0, &p, sizeof(p));
+        g.dispatch_at(dispatch_idx).workgroup_count_x = dispatch_x;
+        g.dispatch_at(dispatch_idx).workgroup_count_y =
+            (m + kTile - 1u) / kTile;
+        g.set_cur_dims(
+            out_id, {static_cast<int64_t>(m), static_cast<int64_t>(N)});
+      });
+
+  wgpuShaderModuleRelease(shader);
+  wgpuBindGroupLayoutRelease(bgl);
+  wgpuPipelineLayoutRelease(pipeline_layout);
+  graph.own_uniform_buffer(uniform_buffer);
+}
+
+} // namespace
+
+WEBGPU_REGISTER_OPERATORS {
+  WEBGPU_REGISTER_OP(aten.mm.default, mm_impl);
+}
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/mm/mm.wgsl b/backends/webgpu/runtime/ops/mm/mm.wgsl
new file mode 100644
index 00000000000..a1b7c5d0c7d
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm.wgsl
@@ -0,0 +1,29 @@
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+override wg_size: u32 = 256;
+
+@compute @workgroup_size(wg_size)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let idx = gid.x;
+  let total = params.M * params.N;
+  if (idx >= total) {
+    return;
+  }
+  let m = idx / params.N;
+  let n = idx % params.N;
+  var acc: f32 = 0.0;
+  for (var k: u32 = 0u; k < params.K; k = k + 1u) {
+    acc = acc + a[m * params.K + k] * b[k * params.N + n];
+  }
+  out[idx] = acc;
+}
diff --git a/backends/webgpu/runtime/ops/mm/mm_tiled.wgsl b/backends/webgpu/runtime/ops/mm/mm_tiled.wgsl
new file mode 100644
index 00000000000..0a91b23157d
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm_tiled.wgsl
@@ -0,0 +1,82 @@
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return a[row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(row: u32, col: u32) -> f32 {
+  if (row < params.K && col < params.N) {
+    return b[row * params.N + col];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/mm/mm_tiled_wgsl.h b/backends/webgpu/runtime/ops/mm/mm_tiled_wgsl.h
new file mode 100644
index 00000000000..3728422b367
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm_tiled_wgsl.h
@@ -0,0 +1,106 @@
+/*
+ * 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 mm_tiled.wgsl - DO NOT EDIT.
+// wgsl-sha256: 0cae16d6afbb610124d2a3f91add4a58c23f57e3d8338e7d6e9ce351227efa3e
+inline constexpr const char* kMmTiledWGSL = R"(
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+
+var<workgroup> a_sub: array<array<f32, 32>, 32>;
+var<workgroup> b_sub: array<array<f32, 32>, 32>;
+
+fn read_a(row: u32, col: u32) -> f32 {
+  if (row < params.M && col < params.K) {
+    return a[row * params.K + col];
+  }
+  return 0.0;
+}
+
+fn read_b(row: u32, col: u32) -> f32 {
+  if (row < params.K && col < params.N) {
+    return b[row * params.N + col];
+  }
+  return 0.0;
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0 = wg_id.x * TILE;
+  let tile_row = local_id.y * RPT;
+  let tile_col = local_id.x * RPT;
+
+  var acc: array<array<f32, 4>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      acc[ir][ic] = 0.0;
+    }
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k_start = t * TILE;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      for (var kk: u32 = 0u; kk < RPT; kk = kk + 1u) {
+        let col = local_id.x * RPT + kk;
+        a_sub[arow][col] = read_a(tile_row0 + arow, k_start + col);
+        b_sub[arow][col] = read_b(k_start + arow, tile_col0 + col);
+      }
+    }
+    workgroupBarrier();
+
+    for (var k: u32 = 0u; k < TILE; k = k + 1u) {
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let aval = a_sub[tile_row + ir][k];
+        for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+          acc[ir][ic] = acc[ir][ic] + aval * b_sub[k][tile_col + ic];
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    for (var ic: u32 = 0u; ic < RPT; ic = ic + 1u) {
+      let r = tile_row0 + tile_row + ir;
+      let c = tile_col0 + tile_col + ic;
+      if (r < params.M && c < params.N) {
+        out[r * params.N + c] = acc[ir][ic];
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kMmTiledWorkgroupSizeX = 8;
+inline constexpr uint32_t kMmTiledWorkgroupSizeY = 8;
+inline constexpr uint32_t kMmTiledWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/mm/mm_vec4.wgsl b/backends/webgpu/runtime/ops/mm/mm_vec4.wgsl
new file mode 100644
index 00000000000..7d198fcc8cd
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm_vec4.wgsl
@@ -0,0 +1,79 @@
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> b: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<vec4<f32>>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+const TILE4: u32 = 8u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+fn read_a4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.M && k4 * 4u < params.K) {
+    return a[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+fn read_b4(krow: u32, n4: u32) -> vec4<f32> {
+  if (krow < params.K && n4 * 4u < params.N) {
+    return b[krow * (params.N / 4u) + n4];
+  }
+  return vec4<f32>(0.0);
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0_4 = wg_id.x * TILE4;
+  let tile_row = local_id.y * RPT;
+  let tile_col4 = local_id.x;
+
+  var acc: array<vec4<f32>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    acc[ir] = vec4<f32>(0.0);
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4_start = t * TILE4;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      a_sub[arow][local_id.x] = read_a4(tile_row0 + arow, k4_start + local_id.x);
+      b_sub[arow][local_id.x] =
+          read_b4(t * TILE + arow, tile_col0_4 + local_id.x);
+    }
+    workgroupBarrier();
+
+    for (var k4: u32 = 0u; k4 < TILE4; k4 = k4 + 1u) {
+      let b0 = b_sub[k4 * 4u + 0u][tile_col4];
+      let b1 = b_sub[k4 * 4u + 1u][tile_col4];
+      let b2 = b_sub[k4 * 4u + 2u][tile_col4];
+      let b3 = b_sub[k4 * 4u + 3u][tile_col4];
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let ac = a_sub[tile_row + ir][k4];
+        acc[ir] = acc[ir] + b0 * ac.x + b1 * ac.y + b2 * ac.z + b3 * ac.w;
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    let r = tile_row0 + tile_row + ir;
+    let c4 = tile_col0_4 + tile_col4;
+    if (r < params.M && c4 * 4u < params.N) {
+      out[r * (params.N / 4u) + c4] = acc[ir];
+    }
+  }
+}
diff --git a/backends/webgpu/runtime/ops/mm/mm_vec4_wgsl.h b/backends/webgpu/runtime/ops/mm/mm_vec4_wgsl.h
new file mode 100644
index 00000000000..a0b37cafb57
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm_vec4_wgsl.h
@@ -0,0 +1,103 @@
+/*
+ * 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 mm_vec4.wgsl - DO NOT EDIT.
+// wgsl-sha256: 208e78ed44618a73f5401172e6fc042bff799489cf4a18fbb4616a43619ef668
+inline constexpr const char* kMmVec4WGSL = R"(
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> b: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> out: array<vec4<f32>>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+const TILE: u32 = 32u;
+const RPT: u32 = 4u;
+const TILE4: u32 = 8u;
+
+var<workgroup> a_sub: array<array<vec4<f32>, 8>, 32>;
+var<workgroup> b_sub: array<array<vec4<f32>, 8>, 32>;
+
+fn read_a4(row: u32, k4: u32) -> vec4<f32> {
+  if (row < params.M && k4 * 4u < params.K) {
+    return a[row * (params.K / 4u) + k4];
+  }
+  return vec4<f32>(0.0);
+}
+
+fn read_b4(krow: u32, n4: u32) -> vec4<f32> {
+  if (krow < params.K && n4 * 4u < params.N) {
+    return b[krow * (params.N / 4u) + n4];
+  }
+  return vec4<f32>(0.0);
+}
+
+@compute @workgroup_size(8, 8, 1)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) local_id: vec3<u32>) {
+  let tile_row0 = wg_id.y * TILE;
+  let tile_col0_4 = wg_id.x * TILE4;
+  let tile_row = local_id.y * RPT;
+  let tile_col4 = local_id.x;
+
+  var acc: array<vec4<f32>, 4>;
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    acc[ir] = vec4<f32>(0.0);
+  }
+
+  let num_tiles = (params.K + TILE - 1u) / TILE;
+  for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
+    let k4_start = t * TILE4;
+    for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+      let arow = local_id.y * RPT + ir;
+      a_sub[arow][local_id.x] = read_a4(tile_row0 + arow, k4_start + local_id.x);
+      b_sub[arow][local_id.x] =
+          read_b4(t * TILE + arow, tile_col0_4 + local_id.x);
+    }
+    workgroupBarrier();
+
+    for (var k4: u32 = 0u; k4 < TILE4; k4 = k4 + 1u) {
+      let b0 = b_sub[k4 * 4u + 0u][tile_col4];
+      let b1 = b_sub[k4 * 4u + 1u][tile_col4];
+      let b2 = b_sub[k4 * 4u + 2u][tile_col4];
+      let b3 = b_sub[k4 * 4u + 3u][tile_col4];
+      for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+        let ac = a_sub[tile_row + ir][k4];
+        acc[ir] = acc[ir] + b0 * ac.x + b1 * ac.y + b2 * ac.z + b3 * ac.w;
+      }
+    }
+    workgroupBarrier();
+  }
+
+  for (var ir: u32 = 0u; ir < RPT; ir = ir + 1u) {
+    let r = tile_row0 + tile_row + ir;
+    let c4 = tile_col0_4 + tile_col4;
+    if (r < params.M && c4 * 4u < params.N) {
+      out[r * (params.N / 4u) + c4] = acc[ir];
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kMmVec4WorkgroupSizeX = 8;
+inline constexpr uint32_t kMmVec4WorkgroupSizeY = 8;
+inline constexpr uint32_t kMmVec4WorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/mm/mm_wgsl.h b/backends/webgpu/runtime/ops/mm/mm_wgsl.h
new file mode 100644
index 00000000000..552074dc353
--- /dev/null
+++ b/backends/webgpu/runtime/ops/mm/mm_wgsl.h
@@ -0,0 +1,53 @@
+/*
+ * 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 mm.wgsl - DO NOT EDIT.
+// wgsl-sha256: bbd9aef967c5085c52eb88b8abcd68202cefe7fc4df04a25c4b00781c5b7ca5a
+inline constexpr const char* kMmWGSL = R"(
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  pad_: u32,
+};
+
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+override wg_size: u32 = 256;
+
+@compute @workgroup_size(wg_size)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let idx = gid.x;
+  let total = params.M * params.N;
+  if (idx >= total) {
+    return;
+  }
+  let m = idx / params.N;
+  let n = idx % params.N;
+  var acc: f32 = 0.0;
+  for (var k: u32 = 0u; k < params.K; k = k + 1u) {
+    acc = acc + a[m * params.K + k] * b[k * params.N + n];
+  }
+  out[idx] = acc;
+}
+)";
+
+inline constexpr uint32_t kMmWorkgroupSizeX = 256;
+inline constexpr uint32_t kMmWorkgroupSizeY = 1;
+inline constexpr uint32_t kMmWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu