diff --git a/js/web/lib/wasm/jsep/webgpu/ops/3rd-party/matmul_packed_webgpu.ts b/js/web/lib/wasm/jsep/webgpu/ops/3rd-party/matmul_packed_webgpu.ts
index c6341f94cf191..d42b5685f3cc3 100644
--- a/js/web/lib/wasm/jsep/webgpu/ops/3rd-party/matmul_packed_webgpu.ts
+++ b/js/web/lib/wasm/jsep/webgpu/ops/3rd-party/matmul_packed_webgpu.ts
@@ -67,19 +67,23 @@ const calculateResultSnippet = (transposeA: boolean, innerElementSize: number) =
         let ACached2 = mm_Asub[k * innerElementSize + 2][localRow];
         ${innerElementSize === 3 ? '' : 'let ACached3 = mm_Asub[k * innerElementSize + 3][localRow];'}
         for (var i = 0; i < rowPerThread; i = i + 1) {
-          acc[i] = BCached0 * ACached0[i] + acc[i];
-          acc[i] = BCached1 * ACached1[i] + acc[i];
-          acc[i] = BCached2 * ACached2[i] + acc[i];
-          ${innerElementSize === 3 ? '' : 'acc[i] = BCached3 * ACached3[i] + acc[i];'}
+          // Explicit f32 casts on each operand are required: Dawn/D3D12 re-demotes
+          // temporaries to f16 when 'enable f16;' is active (issue #26732).
+          acc[i] = vec4<f32>(BCached0) * f32(ACached0[i]) + acc[i];
+          acc[i] = vec4<f32>(BCached1) * f32(ACached1[i]) + acc[i];
+          acc[i] = vec4<f32>(BCached2) * f32(ACached2[i]) + acc[i];
+          ${innerElementSize === 3 ? '' : 'acc[i] = vec4<f32>(BCached3) * f32(ACached3[i]) + acc[i];'}
         }`;
   } else {
     return `
         for (var i = 0; i < rowPerThread; i = i + 1) {
           let ACached = mm_Asub[tileRow + i][k];
-          acc[i] = BCached0 * ACached.x + acc[i];
-          acc[i] = BCached1 * ACached.y + acc[i];
-          acc[i] = BCached2 * ACached.z + acc[i];
-          ${innerElementSize === 3 ? '' : 'acc[i] = BCached3 * ACached.w + acc[i];'}
+          // Explicit f32 casts on each operand are required: Dawn/D3D12 re-demotes
+          // temporaries to f16 when 'enable f16;' is active (issue #26732).
+          acc[i] = vec4<f32>(BCached0) * f32(ACached.x) + acc[i];
+          acc[i] = vec4<f32>(BCached1) * f32(ACached.y) + acc[i];
+          acc[i] = vec4<f32>(BCached2) * f32(ACached.z) + acc[i];
+          ${innerElementSize === 3 ? '' : 'acc[i] = vec4<f32>(BCached3) * f32(ACached.w) + acc[i];'}
         }`;
   }
 };
@@ -140,7 +144,9 @@ fn main(@builtin(local_invocation_id) localId : vec3<u32>,
   let num_tiles = ${splitK ? `${Math.ceil(splitedDimInner / tileInner)}` : '(uniforms.dim_inner - 1) / tileInner + 1'};
   var kStart = ${splitK ? `i32(globalId.z) * ${splitedDimInner}` : '0'};
 
-  var acc: array<vec4<${type}>, rowPerThread>;
+  // Accumulate in f32 to prevent fp16 overflow in long dot products (issue #26732).
+  // Tiles (mm_Asub/mm_Bsub) stay in ${type}: no shared-memory or bandwidth regression.
+  var acc: array<vec4<f32>, rowPerThread>;
 
   // Loop over shared dimension.
   let tileRowB = localRow * ${rowPerThreadB};
@@ -177,7 +183,7 @@ fn main(@builtin(local_invocation_id) localId : vec3<u32>,
   }
 
   for (var innerRow = 0; innerRow < rowPerThread; innerRow = innerRow + 1) {
-      mm_write(batch, globalRow + innerRow, globalCol, acc[innerRow]);
+      mm_write(batch, globalRow + innerRow, globalCol, vec4<${type}>(acc[innerRow]));
   }
 }`;
 };
@@ -268,8 +274,10 @@ export const makeMatMulPackedSource = (
               : `mm_Asub[localRow + innerRow * ${workgroupSize[1]}][k];`
           }
           for (var innerCol = 0; innerCol < colPerThread; innerCol = innerCol + 1) {
+            // Explicit f32 casts required: Dawn/D3D12 re-demotes temporaries to f16
+            // when 'enable f16;' is active (issue #26732).
             acc[innerRow][innerCol] = acc[innerRow][innerCol] +
-                ACached * BCached[innerCol];
+                f32(ACached) * f32(BCached[innerCol]);
           }
         }
       }
@@ -279,7 +287,7 @@ export const makeMatMulPackedSource = (
       let gRow = globalRowStart + localRow + innerRow * ${workgroupSize[1]};
       for (var innerCol = 0; innerCol < colPerThread; innerCol = innerCol + 1) {
         let gCol = globalColStart + localCol + innerCol * ${workgroupSize[0]};
-        mm_write(batch, gRow, gCol, acc[innerRow][innerCol]);
+        mm_write(batch, gRow, gCol, ${type}(acc[innerRow][innerCol]));
       }
     }
     `
@@ -328,7 +336,9 @@ for (var t = 0; t < num_tiles; t = t + 1) {
     for (var innerRow = 0; innerRow < rowPerThread; innerRow = innerRow + 1) {
       ${readDataFromSubASnippet(transposeA)}
       for (var innerCol = 0; innerCol < colPerThread; innerCol = innerCol + 1) {
-        acc[innerRow][innerCol] = acc[innerRow][innerCol] + ACached * BCached[innerCol];
+        // Explicit f32 casts required: Dawn/D3D12 re-demotes temporaries to f16
+        // when 'enable f16;' is active (issue #26732).
+        acc[innerRow][innerCol] = acc[innerRow][innerCol] + f32(ACached) * f32(BCached[innerCol]);
       }
     }
   }
@@ -339,7 +349,7 @@ for (var t = 0; t < num_tiles; t = t + 1) {
 for (var innerRow = 0; innerRow < rowPerThread; innerRow = innerRow + 1) {
   for (var innerCol = 0; innerCol < colPerThread; innerCol = innerCol + 1) {
     mm_write(batch, globalRow + innerRow, globalCol + innerCol,
-        acc[innerRow][innerCol]);
+        ${type}(acc[innerRow][innerCol]));
   }
 }
 `;
@@ -362,7 +372,9 @@ fn main(@builtin(local_invocation_id) localId : vec3<u32>,
     };
     var kStart = ${splitK ? `i32(globalId.z) * ${splitedDimInner}` : '0'};
 
-    var acc : array<array<${type}, colPerThread>, rowPerThread>;
+    // Accumulate in f32 to prevent fp16 overflow in long dot products (issue #26732).
+    // Tiles (mm_Asub/mm_Bsub) stay in ${type}: no shared-memory or bandwidth regression.
+    var acc : array<array<f32, colPerThread>, rowPerThread>;
     ${matmulSnippet}
   }
 `;
diff --git a/js/web/lib/wasm/jsep/webgpu/ops/matmul-shaders.ts b/js/web/lib/wasm/jsep/webgpu/ops/matmul-shaders.ts
index e1f73f137e43e..806ac7dfb2dbc 100644
--- a/js/web/lib/wasm/jsep/webgpu/ops/matmul-shaders.ts
+++ b/js/web/lib/wasm/jsep/webgpu/ops/matmul-shaders.ts
@@ -112,6 +112,10 @@ export const createNaiveMatmulProgramInfo = (
     ];
     appendActivationUniforms(activationAttributes, uniforms);
 
+    // Accumulate in f32 to prevent fp16 overflow in long dot products (issue #26732).
+    // Explicit f32 casts on each operand are required: Dawn/D3D12 re-demotes
+    // temporaries to f16 when 'enable f16;' is active.
+    const accType = components === 1 ? 'f32' : `vec${components}<f32>`;
     const calcResult = (): string => {
       let calcStr = `var a_data: ${a.type.value};`;
       for (let i = 0; i < aComponents; i++) {
@@ -123,7 +127,7 @@ export const createNaiveMatmulProgramInfo = (
 
         for (let j = 0; j < aComponents; j++) {
           calcStr += `
-            values[${i}] = fma(${b.type.value}(a_data${aComponents === 1 ? '' : `[${j}]`}), b_data${j}, values[${i}]);\n`;
+            values[${i}] = fma(${accType}(a_data${aComponents === 1 ? '' : `[${j}]`}), ${accType}(b_data${j}), values[${i}]);\n`;
         }
       }
       return calcStr;
@@ -155,12 +159,13 @@ export const createNaiveMatmulProgramInfo = (
     ${b.indicesSet('b_indices', b.rank - 2, 0)}
     ${b.indicesSet('b_indices', b.rank - 1, 0)}
     let b_offset = ${b.indicesToOffset('b_indices')};
-    var values: array<${output.type.value}, ${outputNumber}>;
+    var values: array<${accType}, ${outputNumber}>;
     for (var k: u32 = 0u; k < uniforms.K; k = k + ${aComponents}) {
       ${calcResult()}
     }
     for (var i = 0u; i < ${outputNumber}u; i++) {
-      var value = values[i];
+      // Downcast to the output type only at the final write.
+      var value = ${output.type.value}(values[i]);
       ${processBias}
       ${applyActivation}
       let cur_indices = ${output.type.indices}(batch, row + i, col);
diff --git a/js/web/lib/wasm/jsep/webgpu/ops/matmulnbits.ts b/js/web/lib/wasm/jsep/webgpu/ops/matmulnbits.ts
index d22acdbe0e782..46b23d8f05516 100644
--- a/js/web/lib/wasm/jsep/webgpu/ops/matmulnbits.ts
+++ b/js/web/lib/wasm/jsep/webgpu/ops/matmulnbits.ts
@@ -180,8 +180,8 @@ export const createMatMulNBitsProgramInfo = (
               (_, i) =>
                 `${
                   aComponents === 1
-                    ? `a_data${pass > 0 ? pass : ''}[${i}] * b_dequantized_values[${i}]`
-                    : `dot(a_data${pass > 0 ? pass : ''}[${i}], b_dequantized_values[${i}])`
+                    ? `f32(a_data${pass > 0 ? pass : ''}[${i}]) * f32(b_dequantized_values[${i}])`
+                    : `dot(vec${aComponents}<f32>(a_data${pass > 0 ? pass : ''}[${i}]), vec${aComponents}<f32>(b_dequantized_values[${i}]))`
                 }`,
             ).join(' + ')};
           `;
@@ -242,8 +242,13 @@ export const createMatMulNBitsProgramInfo = (
             var b_dequantized_values: ${qDqDataType};`;
       return calcStr;
     };
+    // Accumulate in f32 to prevent fp16 overflow in long dot products (issue #26732).
+    // Weights and activations stay in their original dtype; explicit f32 casts on each
+    // operand are required because Dawn/D3D12 re-demotes temporaries to f16 when
+    // 'enable f16;' is active. The result is downcast to the output type only at the write.
+    const accType = components === 1 ? 'f32' : `vec${components}<f32>`;
     return `
-        var<workgroup> workgroup_shared: array<${output.type.value}, ${outputNumber * workgroupSize}>;
+        var<workgroup> workgroup_shared: array<${accType}, ${outputNumber * workgroupSize}>;
         ${shaderHelper.declareVariables(...inputVariables, output)}
         ${shaderHelper.mainStart([workgroupSize, 1, 1])}
           let output_indices = ${output.offsetToIndices(`(global_idx / ${workgroupSize}) * ${outputNumber}`)};
@@ -267,13 +272,13 @@ export const createMatMulNBitsProgramInfo = (
           workgroupBarrier();
 
           if (local_id.x < ${outputNumber}) {
-            var output_value: ${output.type.value} = ${output.type.value}(0);
+            var output_value: ${accType} = ${accType}(0);
             var workgroup_shared_offset: u32 = local_id.x;
             for (var b: u32 = 0u; b < ${workgroupSize}u; b++) {
               output_value += workgroup_shared[workgroup_shared_offset];
               workgroup_shared_offset += ${outputNumber};
             }
-            ${output.setByIndices(`${output.type.indices}(batch, row, col + local_id.x)`, 'output_value')};
+            ${output.setByIndices(`${output.type.indices}(batch, row, col + local_id.x)`, `${output.type.value}(output_value)`)};
           }
         }`;
   };
@@ -368,7 +373,8 @@ export const createMatMulNBitsBlockSize32ProgramInfo = (
 
     return `
         var<workgroup> sub_a: array<${a.type.value}, ${aLengthPerTile}>;
-        var<workgroup> inter_results: array<array<${output.type.value}, ${workgroupX}>, ${workgroupY}>;
+        // Accumulate in f32 to prevent fp16 overflow in long dot products (issue #26732).
+        var<workgroup> inter_results: array<array<f32, ${workgroupX}>, ${workgroupY}>;
         ${shaderHelper.declareVariables(...inputVariables, output)}
         ${shaderHelper.mainStart([workgroupX, workgroupY, 1])}
           let output_indices = ${output.offsetToIndices(`workgroup_index * ${workgroupY}`)};
@@ -444,9 +450,11 @@ export const createMatMulNBitsBlockSize32ProgramInfo = (
                   (_, i) => `${dataType}(b_value_lower[${i}]), ${dataType}(b_value_upper[${i}])`,
                 ).join(', ')});
                 let b_dequantized_values = (b_quantized_values - mat2x4<${dataType}>(${Array(8).fill('zero_point').join(',')})) * scale;
+                // Explicit f32 casts on each operand are required: Dawn/D3D12 re-demotes
+                // temporaries to f16 when 'enable f16;' is active (issue #26732).
                 inter_results[local_id.y][local_id.x] += ${Array.from(
                   { length: 2 },
-                  (_, i) => `${`dot(a_data${i}, b_dequantized_values[${i}])`}`,
+                  (_, i) => `${`dot(vec4<f32>(a_data${i}), vec4<f32>(b_dequantized_values[${i}]))`}`,
                 ).join(' + ')};
               }
               word_offset += ${8 / aComponents};`;
@@ -458,13 +466,13 @@ export const createMatMulNBitsBlockSize32ProgramInfo = (
           }
 
           if (local_idx < ${workgroupY}) {
-            var output_value: ${output.type.value} = ${output.type.value}(0);
+            var output_value: f32 = f32(0);
             for (var b = 0u; b < ${workgroupX}; b++) {
               output_value += inter_results[local_idx][b];
             }
             if (col + local_idx < uniforms.output_shape[2])
             {
-              ${output.setByIndices(`${output.type.indices}(batch, row, col + local_idx)`, 'output_value')}
+              ${output.setByIndices(`${output.type.indices}(batch, row, col + local_idx)`, `${output.type.value}(output_value)`)}
             }
           }
         }`;
diff --git a/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template b/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template
index 342d73a7941ea..5a059f408ed0d 100644
--- a/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template
@@ -217,12 +217,14 @@ $MAIN {
         base_B = subtile_idy * 16;
         a_idx = sg_id;
     }
-    var lane_outputs: array<output_element_t, 16>;
+    // Accumulate in f32 to avoid f16 overflow (max ~65504) when summing partial dot products over K.
+    var lane_outputs: array<f32, 16>;
 #else
-    var lane_output1: vec4<output_element_t>;
-    var lane_output2: vec4<output_element_t>;
-    var lane_output3: vec4<output_element_t>;
-    var lane_output4: vec4<output_element_t>;
+    // Accumulate in f32 to avoid f16 overflow (max ~65504) when summing partial dot products over K.
+    var lane_output1: vec4<f32>;
+    var lane_output2: vec4<f32>;
+    var lane_output3: vec4<f32>;
+    var lane_output4: vec4<f32>;
 #endif
     // K's vectorization is 16 items per index. See input_a/input_b.
     // tile_size_k_vec - is the k tile size in vectorized space (1/16). That is
@@ -264,14 +266,14 @@ $MAIN {
             // Step 2: Access registers across the subgroup using subgroupShuffle and perform the matmul.
             for (var i = 0u; i < 16u; i++)
             {
-                lane_outputs[i] += SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a, subgroupShuffle(zero, i));
+                lane_outputs[i] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a, subgroupShuffle(zero, i)));
             }
         }
         else
         {
             for (var i = 0u; i < 16u; i++)
             {
-                lane_outputs[i] += SDP8AI(own_a0, tile_B[0][base_B + i], own_a1, tile_B[1][base_B + i],  own_scale_a * scale_B[base_B + i], zeroes[base_B + i]);
+                lane_outputs[i] += f32(SDP8AI(own_a0, tile_B[0][base_B + i], own_a1, tile_B[1][base_B + i],  own_scale_a * scale_B[base_B + i], zeroes[base_B + i]));
             }
         }
  #else
@@ -282,50 +284,50 @@ $MAIN {
             var own_scale_b: output_element_t  = scale_B[base_B + sg_id];
             var zero = zeroes[base_B + sg_id];
             // Step 2: Access registers across the subgroup using subgroupShuffle and perform the matmul.
-            lane_output1[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 0), own_a1, subgroupShuffle(own_b1, 0), subgroupShuffle(own_scale_b, 0) * own_scale_a, subgroupShuffle(zero, 0));
-            lane_output1[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 1), own_a1, subgroupShuffle(own_b1, 1), subgroupShuffle(own_scale_b, 1) * own_scale_a, subgroupShuffle(zero, 1));
-            lane_output1[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 2), own_a1, subgroupShuffle(own_b1, 2), subgroupShuffle(own_scale_b, 2) * own_scale_a, subgroupShuffle(zero, 2));
-            lane_output1[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 3), own_a1, subgroupShuffle(own_b1, 3), subgroupShuffle(own_scale_b, 3) * own_scale_a, subgroupShuffle(zero, 3));
+            lane_output1[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 0), own_a1, subgroupShuffle(own_b1, 0), subgroupShuffle(own_scale_b, 0) * own_scale_a, subgroupShuffle(zero, 0)));
+            lane_output1[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 1), own_a1, subgroupShuffle(own_b1, 1), subgroupShuffle(own_scale_b, 1) * own_scale_a, subgroupShuffle(zero, 1)));
+            lane_output1[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 2), own_a1, subgroupShuffle(own_b1, 2), subgroupShuffle(own_scale_b, 2) * own_scale_a, subgroupShuffle(zero, 2)));
+            lane_output1[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 3), own_a1, subgroupShuffle(own_b1, 3), subgroupShuffle(own_scale_b, 3) * own_scale_a, subgroupShuffle(zero, 3)));
 
-            lane_output2[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 4), own_a1, subgroupShuffle(own_b1, 4), subgroupShuffle(own_scale_b, 4) * own_scale_a, subgroupShuffle(zero, 4));
-            lane_output2[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 5), own_a1, subgroupShuffle(own_b1, 5), subgroupShuffle(own_scale_b, 5) * own_scale_a, subgroupShuffle(zero, 5));
-            lane_output2[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 6), own_a1, subgroupShuffle(own_b1, 6), subgroupShuffle(own_scale_b, 6) * own_scale_a, subgroupShuffle(zero, 6));
-            lane_output2[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 7), own_a1, subgroupShuffle(own_b1, 7), subgroupShuffle(own_scale_b, 7) * own_scale_a, subgroupShuffle(zero, 7));
+            lane_output2[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 4), own_a1, subgroupShuffle(own_b1, 4), subgroupShuffle(own_scale_b, 4) * own_scale_a, subgroupShuffle(zero, 4)));
+            lane_output2[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 5), own_a1, subgroupShuffle(own_b1, 5), subgroupShuffle(own_scale_b, 5) * own_scale_a, subgroupShuffle(zero, 5)));
+            lane_output2[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 6), own_a1, subgroupShuffle(own_b1, 6), subgroupShuffle(own_scale_b, 6) * own_scale_a, subgroupShuffle(zero, 6)));
+            lane_output2[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 7), own_a1, subgroupShuffle(own_b1, 7), subgroupShuffle(own_scale_b, 7) * own_scale_a, subgroupShuffle(zero, 7)));
 
-            lane_output3[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 8), own_a1, subgroupShuffle(own_b1, 8), subgroupShuffle(own_scale_b, 8) * own_scale_a, subgroupShuffle(zero, 8));
-            lane_output3[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 9), own_a1, subgroupShuffle(own_b1, 9), subgroupShuffle(own_scale_b, 9) * own_scale_a, subgroupShuffle(zero, 9));
-            lane_output3[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 10), own_a1, subgroupShuffle(own_b1, 10), subgroupShuffle(own_scale_b, 10) * own_scale_a, subgroupShuffle(zero, 10));
-            lane_output3[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 11), own_a1, subgroupShuffle(own_b1, 11), subgroupShuffle(own_scale_b, 11) * own_scale_a, subgroupShuffle(zero, 11));
+            lane_output3[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 8), own_a1, subgroupShuffle(own_b1, 8), subgroupShuffle(own_scale_b, 8) * own_scale_a, subgroupShuffle(zero, 8)));
+            lane_output3[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 9), own_a1, subgroupShuffle(own_b1, 9), subgroupShuffle(own_scale_b, 9) * own_scale_a, subgroupShuffle(zero, 9)));
+            lane_output3[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 10), own_a1, subgroupShuffle(own_b1, 10), subgroupShuffle(own_scale_b, 10) * own_scale_a, subgroupShuffle(zero, 10)));
+            lane_output3[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 11), own_a1, subgroupShuffle(own_b1, 11), subgroupShuffle(own_scale_b, 11) * own_scale_a, subgroupShuffle(zero, 11)));
 
-            lane_output4[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 12), own_a1, subgroupShuffle(own_b1, 12), subgroupShuffle(own_scale_b, 12) * own_scale_a, subgroupShuffle(zero, 12));
-            lane_output4[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 13), own_a1, subgroupShuffle(own_b1, 13), subgroupShuffle(own_scale_b, 13) * own_scale_a, subgroupShuffle(zero, 13));
-            lane_output4[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 14), own_a1, subgroupShuffle(own_b1, 14), subgroupShuffle(own_scale_b, 14) * own_scale_a, subgroupShuffle(zero, 14));
-            lane_output4[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 15), own_a1, subgroupShuffle(own_b1, 15), subgroupShuffle(own_scale_b, 15) * own_scale_a, subgroupShuffle(zero, 15));
+            lane_output4[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 12), own_a1, subgroupShuffle(own_b1, 12), subgroupShuffle(own_scale_b, 12) * own_scale_a, subgroupShuffle(zero, 12)));
+            lane_output4[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 13), own_a1, subgroupShuffle(own_b1, 13), subgroupShuffle(own_scale_b, 13) * own_scale_a, subgroupShuffle(zero, 13)));
+            lane_output4[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 14), own_a1, subgroupShuffle(own_b1, 14), subgroupShuffle(own_scale_b, 14) * own_scale_a, subgroupShuffle(zero, 14)));
+            lane_output4[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 15), own_a1, subgroupShuffle(own_b1, 15), subgroupShuffle(own_scale_b, 15) * own_scale_a, subgroupShuffle(zero, 15)));
         }
         else
         {
             // Code for other subgroup sizes, simply doesnt use subgroups at all.
             // Relies on reads from single location tile_B[][base_B + col] by all
             // being optimized by the hardware.
-            lane_output1[0] += SDP8AI(own_a0, tile_B[0][base_B + 0], own_a1, tile_B[1][base_B + 0],  own_scale_a * scale_B[base_B + 0], zeroes[base_B + 0]);
-            lane_output1[1] += SDP8AI(own_a0, tile_B[0][base_B + 1], own_a1, tile_B[1][base_B + 1],  own_scale_a * scale_B[base_B + 1], zeroes[base_B + 1]);
-            lane_output1[2] += SDP8AI(own_a0, tile_B[0][base_B + 2], own_a1, tile_B[1][base_B + 2],  own_scale_a * scale_B[base_B + 2], zeroes[base_B + 2]);
-            lane_output1[3] += SDP8AI(own_a0, tile_B[0][base_B + 3], own_a1, tile_B[1][base_B + 3],  own_scale_a * scale_B[base_B + 3], zeroes[base_B + 3]);
+            lane_output1[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 0], own_a1, tile_B[1][base_B + 0],  own_scale_a * scale_B[base_B + 0], zeroes[base_B + 0]));
+            lane_output1[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 1], own_a1, tile_B[1][base_B + 1],  own_scale_a * scale_B[base_B + 1], zeroes[base_B + 1]));
+            lane_output1[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 2], own_a1, tile_B[1][base_B + 2],  own_scale_a * scale_B[base_B + 2], zeroes[base_B + 2]));
+            lane_output1[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 3], own_a1, tile_B[1][base_B + 3],  own_scale_a * scale_B[base_B + 3], zeroes[base_B + 3]));
 
-            lane_output2[0] += SDP8AI(own_a0, tile_B[0][base_B + 4], own_a1, tile_B[1][base_B + 4],  own_scale_a * scale_B[base_B + 4], zeroes[base_B + 4]);
-            lane_output2[1] += SDP8AI(own_a0, tile_B[0][base_B + 5], own_a1, tile_B[1][base_B + 5],  own_scale_a * scale_B[base_B + 5], zeroes[base_B + 5]);
-            lane_output2[2] += SDP8AI(own_a0, tile_B[0][base_B + 6], own_a1, tile_B[1][base_B + 6],  own_scale_a * scale_B[base_B + 6], zeroes[base_B + 6]);
-            lane_output2[3] += SDP8AI(own_a0, tile_B[0][base_B + 7], own_a1, tile_B[1][base_B + 7],  own_scale_a * scale_B[base_B + 7], zeroes[base_B + 7]);
+            lane_output2[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 4], own_a1, tile_B[1][base_B + 4],  own_scale_a * scale_B[base_B + 4], zeroes[base_B + 4]));
+            lane_output2[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 5], own_a1, tile_B[1][base_B + 5],  own_scale_a * scale_B[base_B + 5], zeroes[base_B + 5]));
+            lane_output2[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 6], own_a1, tile_B[1][base_B + 6],  own_scale_a * scale_B[base_B + 6], zeroes[base_B + 6]));
+            lane_output2[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 7], own_a1, tile_B[1][base_B + 7],  own_scale_a * scale_B[base_B + 7], zeroes[base_B + 7]));
 
-            lane_output3[0] += SDP8AI(own_a0, tile_B[0][base_B + 8], own_a1, tile_B[1][base_B + 8],  own_scale_a * scale_B[base_B + 8], zeroes[base_B + 8]);
-            lane_output3[1] += SDP8AI(own_a0, tile_B[0][base_B + 9], own_a1, tile_B[1][base_B + 9],  own_scale_a * scale_B[base_B + 9], zeroes[base_B + 9]);
-            lane_output3[2] += SDP8AI(own_a0, tile_B[0][base_B + 10], own_a1, tile_B[1][base_B + 10],  own_scale_a * scale_B[base_B + 10], zeroes[base_B + 10]);
-            lane_output3[3] += SDP8AI(own_a0, tile_B[0][base_B + 11], own_a1, tile_B[1][base_B + 11],  own_scale_a * scale_B[base_B + 11], zeroes[base_B + 11]);
+            lane_output3[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 8], own_a1, tile_B[1][base_B + 8],  own_scale_a * scale_B[base_B + 8], zeroes[base_B + 8]));
+            lane_output3[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 9], own_a1, tile_B[1][base_B + 9],  own_scale_a * scale_B[base_B + 9], zeroes[base_B + 9]));
+            lane_output3[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 10], own_a1, tile_B[1][base_B + 10],  own_scale_a * scale_B[base_B + 10], zeroes[base_B + 10]));
+            lane_output3[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 11], own_a1, tile_B[1][base_B + 11],  own_scale_a * scale_B[base_B + 11], zeroes[base_B + 11]));
 
-            lane_output4[0] += SDP8AI(own_a0, tile_B[0][base_B + 12], own_a1, tile_B[1][base_B + 12],  own_scale_a * scale_B[base_B + 12], zeroes[base_B + 12]);
-            lane_output4[1] += SDP8AI(own_a0, tile_B[0][base_B + 13], own_a1, tile_B[1][base_B + 13],  own_scale_a * scale_B[base_B + 13], zeroes[base_B + 13]);
-            lane_output4[2] += SDP8AI(own_a0, tile_B[0][base_B + 14], own_a1, tile_B[1][base_B + 14],  own_scale_a * scale_B[base_B + 14], zeroes[base_B + 14]);
-            lane_output4[3] += SDP8AI(own_a0, tile_B[0][base_B + 15], own_a1, tile_B[1][base_B + 15],  own_scale_a * scale_B[base_B + 15], zeroes[base_B + 15]);
+            lane_output4[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 12], own_a1, tile_B[1][base_B + 12],  own_scale_a * scale_B[base_B + 12], zeroes[base_B + 12]));
+            lane_output4[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 13], own_a1, tile_B[1][base_B + 13],  own_scale_a * scale_B[base_B + 13], zeroes[base_B + 13]));
+            lane_output4[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 14], own_a1, tile_B[1][base_B + 14],  own_scale_a * scale_B[base_B + 14], zeroes[base_B + 14]));
+            lane_output4[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 15], own_a1, tile_B[1][base_B + 15],  own_scale_a * scale_B[base_B + 15], zeroes[base_B + 15]));
         }
  #endif
 #else
@@ -344,14 +346,14 @@ $MAIN {
             // Step 2: Access registers across the subgroup using subgroupShuffle and perform the matmul.
             for (var i = 0u; i < 16u; i++)
             {
-                lane_outputs[i] += SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a);
+                lane_outputs[i] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a));
             }
         }
         else
         {
             for (var i = 0u; i < 16u; i++)
             {
-                lane_outputs[i] += SDP8AI(own_a0, tile_B[0][base_B + i], own_a1, tile_B[1][base_B + i],  own_scale_a * scale_B[base_B + i]);
+                lane_outputs[i] += f32(SDP8AI(own_a0, tile_B[0][base_B + i], own_a1, tile_B[1][base_B + i],  own_scale_a * scale_B[base_B + i]));
             }
         }
  #else
@@ -361,50 +363,50 @@ $MAIN {
             var own_b1: vec4<u32> = tile_B[1][base_B + sg_id];
             var own_scale_b: output_element_t  = scale_B[base_B + sg_id];
             // Step 2: Access registers across the subgroup using subgroupShuffle and perform the matmul.
-            lane_output1[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 0), own_a1, subgroupShuffle(own_b1, 0), subgroupShuffle(own_scale_b, 0) * own_scale_a);
-            lane_output1[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 1), own_a1, subgroupShuffle(own_b1, 1), subgroupShuffle(own_scale_b, 1) * own_scale_a);
-            lane_output1[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 2), own_a1, subgroupShuffle(own_b1, 2), subgroupShuffle(own_scale_b, 2) * own_scale_a);
-            lane_output1[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 3), own_a1, subgroupShuffle(own_b1, 3), subgroupShuffle(own_scale_b, 3) * own_scale_a);
+            lane_output1[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 0), own_a1, subgroupShuffle(own_b1, 0), subgroupShuffle(own_scale_b, 0) * own_scale_a));
+            lane_output1[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 1), own_a1, subgroupShuffle(own_b1, 1), subgroupShuffle(own_scale_b, 1) * own_scale_a));
+            lane_output1[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 2), own_a1, subgroupShuffle(own_b1, 2), subgroupShuffle(own_scale_b, 2) * own_scale_a));
+            lane_output1[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 3), own_a1, subgroupShuffle(own_b1, 3), subgroupShuffle(own_scale_b, 3) * own_scale_a));
 
-            lane_output2[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 4), own_a1, subgroupShuffle(own_b1, 4), subgroupShuffle(own_scale_b, 4) * own_scale_a);
-            lane_output2[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 5), own_a1, subgroupShuffle(own_b1, 5), subgroupShuffle(own_scale_b, 5) * own_scale_a);
-            lane_output2[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 6), own_a1, subgroupShuffle(own_b1, 6), subgroupShuffle(own_scale_b, 6) * own_scale_a);
-            lane_output2[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 7), own_a1, subgroupShuffle(own_b1, 7), subgroupShuffle(own_scale_b, 7) * own_scale_a);
+            lane_output2[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 4), own_a1, subgroupShuffle(own_b1, 4), subgroupShuffle(own_scale_b, 4) * own_scale_a));
+            lane_output2[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 5), own_a1, subgroupShuffle(own_b1, 5), subgroupShuffle(own_scale_b, 5) * own_scale_a));
+            lane_output2[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 6), own_a1, subgroupShuffle(own_b1, 6), subgroupShuffle(own_scale_b, 6) * own_scale_a));
+            lane_output2[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 7), own_a1, subgroupShuffle(own_b1, 7), subgroupShuffle(own_scale_b, 7) * own_scale_a));
 
-            lane_output3[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 8), own_a1, subgroupShuffle(own_b1, 8), subgroupShuffle(own_scale_b, 8) * own_scale_a);
-            lane_output3[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 9), own_a1, subgroupShuffle(own_b1, 9), subgroupShuffle(own_scale_b, 9) * own_scale_a);
-            lane_output3[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 10), own_a1, subgroupShuffle(own_b1, 10), subgroupShuffle(own_scale_b, 10) * own_scale_a);
-            lane_output3[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 11), own_a1, subgroupShuffle(own_b1, 11), subgroupShuffle(own_scale_b, 11) * own_scale_a);
+            lane_output3[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 8), own_a1, subgroupShuffle(own_b1, 8), subgroupShuffle(own_scale_b, 8) * own_scale_a));
+            lane_output3[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 9), own_a1, subgroupShuffle(own_b1, 9), subgroupShuffle(own_scale_b, 9) * own_scale_a));
+            lane_output3[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 10), own_a1, subgroupShuffle(own_b1, 10), subgroupShuffle(own_scale_b, 10) * own_scale_a));
+            lane_output3[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 11), own_a1, subgroupShuffle(own_b1, 11), subgroupShuffle(own_scale_b, 11) * own_scale_a));
 
-            lane_output4[0] += SDP8AI(own_a0, subgroupShuffle(own_b0, 12), own_a1, subgroupShuffle(own_b1, 12), subgroupShuffle(own_scale_b, 12) * own_scale_a);
-            lane_output4[1] += SDP8AI(own_a0, subgroupShuffle(own_b0, 13), own_a1, subgroupShuffle(own_b1, 13), subgroupShuffle(own_scale_b, 13) * own_scale_a);
-            lane_output4[2] += SDP8AI(own_a0, subgroupShuffle(own_b0, 14), own_a1, subgroupShuffle(own_b1, 14), subgroupShuffle(own_scale_b, 14) * own_scale_a);
-            lane_output4[3] += SDP8AI(own_a0, subgroupShuffle(own_b0, 15), own_a1, subgroupShuffle(own_b1, 15), subgroupShuffle(own_scale_b, 15) * own_scale_a);
+            lane_output4[0] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 12), own_a1, subgroupShuffle(own_b1, 12), subgroupShuffle(own_scale_b, 12) * own_scale_a));
+            lane_output4[1] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 13), own_a1, subgroupShuffle(own_b1, 13), subgroupShuffle(own_scale_b, 13) * own_scale_a));
+            lane_output4[2] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 14), own_a1, subgroupShuffle(own_b1, 14), subgroupShuffle(own_scale_b, 14) * own_scale_a));
+            lane_output4[3] += f32(SDP8AI(own_a0, subgroupShuffle(own_b0, 15), own_a1, subgroupShuffle(own_b1, 15), subgroupShuffle(own_scale_b, 15) * own_scale_a));
         }
         else
         {
             // Code for other subgroup sizes, simply doesnt use subgroups at all.
             // Relies on reads from single location tile_B[][base_B + col] by all
             // being optimized by the hardware.
-            lane_output1[0] += SDP8AI(own_a0, tile_B[0][base_B + 0], own_a1, tile_B[1][base_B + 0],  own_scale_a * scale_B[base_B + 0]);
-            lane_output1[1] += SDP8AI(own_a0, tile_B[0][base_B + 1], own_a1, tile_B[1][base_B + 1],  own_scale_a * scale_B[base_B + 1]);
-            lane_output1[2] += SDP8AI(own_a0, tile_B[0][base_B + 2], own_a1, tile_B[1][base_B + 2],  own_scale_a * scale_B[base_B + 2]);
-            lane_output1[3] += SDP8AI(own_a0, tile_B[0][base_B + 3], own_a1, tile_B[1][base_B + 3],  own_scale_a * scale_B[base_B + 3]);
+            lane_output1[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 0], own_a1, tile_B[1][base_B + 0],  own_scale_a * scale_B[base_B + 0]));
+            lane_output1[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 1], own_a1, tile_B[1][base_B + 1],  own_scale_a * scale_B[base_B + 1]));
+            lane_output1[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 2], own_a1, tile_B[1][base_B + 2],  own_scale_a * scale_B[base_B + 2]));
+            lane_output1[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 3], own_a1, tile_B[1][base_B + 3],  own_scale_a * scale_B[base_B + 3]));
 
-            lane_output2[0] += SDP8AI(own_a0, tile_B[0][base_B + 4], own_a1, tile_B[1][base_B + 4],  own_scale_a * scale_B[base_B + 4]);
-            lane_output2[1] += SDP8AI(own_a0, tile_B[0][base_B + 5], own_a1, tile_B[1][base_B + 5],  own_scale_a * scale_B[base_B + 5]);
-            lane_output2[2] += SDP8AI(own_a0, tile_B[0][base_B + 6], own_a1, tile_B[1][base_B + 6],  own_scale_a * scale_B[base_B + 6]);
-            lane_output2[3] += SDP8AI(own_a0, tile_B[0][base_B + 7], own_a1, tile_B[1][base_B + 7],  own_scale_a * scale_B[base_B + 7]);
+            lane_output2[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 4], own_a1, tile_B[1][base_B + 4],  own_scale_a * scale_B[base_B + 4]));
+            lane_output2[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 5], own_a1, tile_B[1][base_B + 5],  own_scale_a * scale_B[base_B + 5]));
+            lane_output2[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 6], own_a1, tile_B[1][base_B + 6],  own_scale_a * scale_B[base_B + 6]));
+            lane_output2[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 7], own_a1, tile_B[1][base_B + 7],  own_scale_a * scale_B[base_B + 7]));
 
-            lane_output3[0] += SDP8AI(own_a0, tile_B[0][base_B + 8], own_a1, tile_B[1][base_B + 8],  own_scale_a * scale_B[base_B + 8]);
-            lane_output3[1] += SDP8AI(own_a0, tile_B[0][base_B + 9], own_a1, tile_B[1][base_B + 9],  own_scale_a * scale_B[base_B + 9]);
-            lane_output3[2] += SDP8AI(own_a0, tile_B[0][base_B + 10], own_a1, tile_B[1][base_B + 10],  own_scale_a * scale_B[base_B + 10]);
-            lane_output3[3] += SDP8AI(own_a0, tile_B[0][base_B + 11], own_a1, tile_B[1][base_B + 11],  own_scale_a * scale_B[base_B + 11]);
+            lane_output3[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 8], own_a1, tile_B[1][base_B + 8],  own_scale_a * scale_B[base_B + 8]));
+            lane_output3[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 9], own_a1, tile_B[1][base_B + 9],  own_scale_a * scale_B[base_B + 9]));
+            lane_output3[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 10], own_a1, tile_B[1][base_B + 10],  own_scale_a * scale_B[base_B + 10]));
+            lane_output3[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 11], own_a1, tile_B[1][base_B + 11],  own_scale_a * scale_B[base_B + 11]));
 
-            lane_output4[0] += SDP8AI(own_a0, tile_B[0][base_B + 12], own_a1, tile_B[1][base_B + 12],  own_scale_a * scale_B[base_B + 12]);
-            lane_output4[1] += SDP8AI(own_a0, tile_B[0][base_B + 13], own_a1, tile_B[1][base_B + 13],  own_scale_a * scale_B[base_B + 13]);
-            lane_output4[2] += SDP8AI(own_a0, tile_B[0][base_B + 14], own_a1, tile_B[1][base_B + 14],  own_scale_a * scale_B[base_B + 14]);
-            lane_output4[3] += SDP8AI(own_a0, tile_B[0][base_B + 15], own_a1, tile_B[1][base_B + 15],  own_scale_a * scale_B[base_B + 15]);
+            lane_output4[0] += f32(SDP8AI(own_a0, tile_B[0][base_B + 12], own_a1, tile_B[1][base_B + 12],  own_scale_a * scale_B[base_B + 12]));
+            lane_output4[1] += f32(SDP8AI(own_a0, tile_B[0][base_B + 13], own_a1, tile_B[1][base_B + 13],  own_scale_a * scale_B[base_B + 13]));
+            lane_output4[2] += f32(SDP8AI(own_a0, tile_B[0][base_B + 14], own_a1, tile_B[1][base_B + 14],  own_scale_a * scale_B[base_B + 14]));
+            lane_output4[3] += f32(SDP8AI(own_a0, tile_B[0][base_B + 15], own_a1, tile_B[1][base_B + 15],  own_scale_a * scale_B[base_B + 15]));
         }
  #endif
 #endif
@@ -455,15 +457,15 @@ $MAIN {
             bias[b_global + 14 + b_bias_offset],
             bias[b_global + 15 + b_bias_offset]
         );
-        output.setByOffset(output_idx, vec4<output_element_t>(lane_outputs[0], lane_outputs[1], lane_outputs[2], lane_outputs[3]) + bias_vec1);
-        output.setByOffset(output_idx+1, vec4<output_element_t>(lane_outputs[4], lane_outputs[5], lane_outputs[6], lane_outputs[7]) + bias_vec2);
-        output.setByOffset(output_idx+2, vec4<output_element_t>(lane_outputs[8], lane_outputs[9], lane_outputs[10], lane_outputs[11]) + bias_vec3);
-        output.setByOffset(output_idx+3, vec4<output_element_t>(lane_outputs[12], lane_outputs[13], lane_outputs[14], lane_outputs[15]) + bias_vec4);
+        output.setByOffset(output_idx, vec4<output_element_t>(vec4<f32>(lane_outputs[0], lane_outputs[1], lane_outputs[2], lane_outputs[3])) + bias_vec1);
+        output.setByOffset(output_idx+1, vec4<output_element_t>(vec4<f32>(lane_outputs[4], lane_outputs[5], lane_outputs[6], lane_outputs[7])) + bias_vec2);
+        output.setByOffset(output_idx+2, vec4<output_element_t>(vec4<f32>(lane_outputs[8], lane_outputs[9], lane_outputs[10], lane_outputs[11])) + bias_vec3);
+        output.setByOffset(output_idx+3, vec4<output_element_t>(vec4<f32>(lane_outputs[12], lane_outputs[13], lane_outputs[14], lane_outputs[15])) + bias_vec4);
  #else
-        output.setByOffset(output_idx, vec4<output_element_t>(lane_outputs[0], lane_outputs[1], lane_outputs[2], lane_outputs[3]));
-        output.setByOffset(output_idx+1, vec4<output_element_t>(lane_outputs[4], lane_outputs[5], lane_outputs[6], lane_outputs[7]));
-        output.setByOffset(output_idx+2, vec4<output_element_t>(lane_outputs[8], lane_outputs[9], lane_outputs[10], lane_outputs[11]));
-        output.setByOffset(output_idx+3, vec4<output_element_t>(lane_outputs[12], lane_outputs[13], lane_outputs[14], lane_outputs[15]));
+        output.setByOffset(output_idx, vec4<output_element_t>(vec4<f32>(lane_outputs[0], lane_outputs[1], lane_outputs[2], lane_outputs[3])));
+        output.setByOffset(output_idx+1, vec4<output_element_t>(vec4<f32>(lane_outputs[4], lane_outputs[5], lane_outputs[6], lane_outputs[7])));
+        output.setByOffset(output_idx+2, vec4<output_element_t>(vec4<f32>(lane_outputs[8], lane_outputs[9], lane_outputs[10], lane_outputs[11])));
+        output.setByOffset(output_idx+3, vec4<output_element_t>(vec4<f32>(lane_outputs[12], lane_outputs[13], lane_outputs[14], lane_outputs[15])));
  #endif
 #else
  #if has_bias
@@ -492,15 +494,15 @@ $MAIN {
             bias[b_global + 14 + b_bias_offset],
             bias[b_global + 15 + b_bias_offset]
         );
-        output.setByOffset(output_idx, lane_output1 + bias_vec1);
-        output.setByOffset(output_idx+1, lane_output2 + bias_vec2);
-        output.setByOffset(output_idx+2, lane_output3 + bias_vec3);
-        output.setByOffset(output_idx+3, lane_output4 + bias_vec4);
+        output.setByOffset(output_idx, vec4<output_element_t>(lane_output1) + bias_vec1);
+        output.setByOffset(output_idx+1, vec4<output_element_t>(lane_output2) + bias_vec2);
+        output.setByOffset(output_idx+2, vec4<output_element_t>(lane_output3) + bias_vec3);
+        output.setByOffset(output_idx+3, vec4<output_element_t>(lane_output4) + bias_vec4);
  #else
-        output.setByOffset(output_idx, lane_output1);
-        output.setByOffset(output_idx+1, lane_output2);
-        output.setByOffset(output_idx+2, lane_output3);
-        output.setByOffset(output_idx+3, lane_output4);
+        output.setByOffset(output_idx, vec4<output_element_t>(lane_output1));
+        output.setByOffset(output_idx+1, vec4<output_element_t>(lane_output2));
+        output.setByOffset(output_idx+2, vec4<output_element_t>(lane_output3));
+        output.setByOffset(output_idx+3, vec4<output_element_t>(lane_output4));
  #endif
 #endif
     }
diff --git a/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_small_m.wgsl.template b/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_small_m.wgsl.template
index 027fc77043cbc..589bda85af073 100644
--- a/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_small_m.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_small_m.wgsl.template
@@ -40,7 +40,8 @@
 
 const double_tile_size_k_vec = 2 * tile_size_k_vec;
 
-var<workgroup> inter_results: array<array<output_element_t, tile_size_k_vec>, tile_size>;
+// Accumulate in f32 to avoid f16 overflow (max ~65504) when summing partial dot products over K.
+var<workgroup> inter_results: array<array<f32, tile_size_k_vec>, tile_size>;
 // Need 2 * tile_size_k_vec to store a tile_A since b is quantized as 4 bits and a is quantized as 8 bits.
 var<workgroup> tile_A : array<vec4<u32>, double_tile_size_k_vec>;
 // double_tile_size_k_vec * 16 / 128
@@ -164,14 +165,14 @@ $MAIN {
                 let b_value = b.getByOffset(b_offset);
                 let own_b = DequantizedFrom4BitsTo8Bits(b_value.xy, zero);
                 let own_b1 = DequantizedFrom4BitsTo8Bits(b_value.zw, zero);
-                inter_results[row_offset + local_row][local_col] += SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b);
+                inter_results[row_offset + local_row][local_col] += f32(SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b));
 #elif n_bits == 8
                 let own_b = AlignWithZeroPoint(b.getByOffset(b_offset * 2));
                 let own_b1 = AlignWithZeroPoint(b.getByOffset(b_offset * 2 + 1));
 #if has_zero_points
-                inter_results[row_offset + local_row][local_col] += SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b, zero);
+                inter_results[row_offset + local_row][local_col] += f32(SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b, zero));
 #else
-                inter_results[row_offset + local_row][local_col] += SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b);
+                inter_results[row_offset + local_row][local_col] += f32(SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b));
 #endif
 
 #elif n_bits == 2
@@ -183,7 +184,7 @@ $MAIN {
                 let own_b = DequantizedFrom2BitsTo8Bits(b_value.x);
                 let own_b1 = DequantizedFrom2BitsTo8Bits(b_value.y);
 #endif
-                inter_results[row_offset + local_row][local_col] += SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b);
+                inter_results[row_offset + local_row][local_col] += f32(SDP8AI(own_a, own_b, own_a1, own_b1, own_scale_a * own_scale_b));
 #endif
             }
         }
@@ -192,7 +193,7 @@ $MAIN {
 
     if (local_idx < tile_size) {
       // Do reduce sum to get final output.
-      var output_value = output_element_t(0);
+      var output_value = f32(0);
       for (var b = 0u; b < tile_size_k_vec; b++) {
         output_value += inter_results[local_idx][b];
       }
@@ -200,10 +201,11 @@ $MAIN {
       let output_idx = batch * uniforms.dispatch_M * uniforms.N + a_global * uniforms.N + b_global;
       if (b_global < uniforms.N) {
 #if has_bias
-        let bias_value = bias[b_global + b_bias_offset];
-        output.setByOffset(output_idx, output_value + bias_value);
+        let bias_value = f32(bias[b_global + b_bias_offset]);
+        // Downcast to the output element type only at the final store.
+        output.setByOffset(output_idx, output_element_t(output_value + bias_value));
 #else
-        output.setByOffset(output_idx, output_value);
+        output.setByOffset(output_idx, output_element_t(output_value));
 #endif
       }
     }
diff --git a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template
index 58f4baadae99f..3cce5f1376b59 100644
--- a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template
@@ -22,7 +22,10 @@
 
 // Shared memory
 var<workgroup> tile_A : array<input_a_value_t, a_length_per_tile>;
-var<workgroup> inter_results: array<array<output_element_t, tile_size_k_vec>, tile_size>;
+// Accumulate partial sums along K in f32: with fp16 outputs, summing 2048+ f16
+// products overflows the f16 max (65504) and poisons the output with +Inf/NaN
+// (issue #26732). Only the block-local `sum` stays in output_element_t.
+var<workgroup> inter_results: array<array<f32, tile_size_k_vec>, tile_size>;
 
 fn loadSHMA(batch: u32, a_global: u32, kidx: u32, col: u32)
 {
@@ -222,7 +225,9 @@ $MAIN {
 #endif
 #endif
 
-        inter_results[local_row_offset + idy][idx] += sum;
+        // Explicit f32 cast: Dawn/D3D12 re-demotes temporaries to f16 when
+        // 'enable f16;' is active (issue #26732).
+        inter_results[local_row_offset + idy][idx] += f32(sum);
       }
     }
     workgroupBarrier();
@@ -233,7 +238,7 @@ $MAIN {
   }
 
   if (local_idx < tile_size) {
-    var output_value = output_element_t(0);
+    var output_value = f32(0);
     for (var b = 0u; b < tile_size_k_vec; b++) {
       output_value += inter_results[local_idx][b];
     }
@@ -241,9 +246,10 @@ $MAIN {
     let output_idx = batch * uniforms.dispatch_M * uniforms.N + a_global * uniforms.N + b_global;
     if (b_global < uniforms.N) {
 #if has_bias
-      output_value += bias[b_global + b_bias_offset];
+      output_value += f32(bias[b_global + b_bias_offset]);
 #endif
-      output.setByOffset(output_idx, output_value);
+      // Downcast to the output type only at the final write.
+      output.setByOffset(output_idx, output_element_t(output_value));
     }
   }
 }  // MAIN
diff --git a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.wgsl.template b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.wgsl.template
index f64f0d38f24e2..04fe7a2cb8616 100644
--- a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.wgsl.template
@@ -29,8 +29,9 @@
 var<workgroup> sum_squared_shared : array<f32, workgroup_size_x>;
 #endif
 var<workgroup> tile_A : array<input_a_value_t, a_length_per_tile>;
-var<workgroup> gate_inter_results : array<array<output_element_t, tile_size_k_vec + 1>, tile_size>;
-var<workgroup> up_inter_results : array<array<output_element_t, tile_size_k_vec + 1>, tile_size>;
+// Accumulate in f32 to avoid f16 overflow (max ~65504) when summing long K dot products.
+var<workgroup> gate_inter_results : array<array<f32, tile_size_k_vec + 1>, tile_size>;
+var<workgroup> up_inter_results : array<array<f32, tile_size_k_vec + 1>, tile_size>;
 
 const default_zero_point = output_element_t(8);
 
@@ -64,7 +65,7 @@ fn loadSHMA(batch: u32, b_global_base: u32, kidx: u32, col: u32, inv_std: f32)
     }
 }
 
-fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> vec2<output_element_t> {
+fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> vec2<f32> {
 #if single_scale_weights
   let gate_scale_b = gate_scales_b.getByOffset(0);
   let up_scale_b = up_scales_b.getByOffset(0);
@@ -76,8 +77,8 @@ fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> ve
   let gate_b_value = gate_b.getByOffset(b_global * uniforms.K_of_b + k_offset);
   let up_b_value = up_b.getByOffset(b_global * uniforms.K_of_b + k_offset);
 
-  var gate_sum = output_element_t(0);
-  var up_sum = output_element_t(0);
+  var gate_sum = f32(0);
+  var up_sum = f32(0);
   var a_offset = idx * (8 / component_a) * component_b;
 #if component_b == 1
   let gate_b_value_lower = vec4<output_element_t>(unpack4xU8(gate_b_value & 0x0F0F0F0Fu)) - vec4<output_element_t>(default_zero_point);
@@ -91,18 +92,18 @@ fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> ve
 #if component_a == 1
   let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1], tile_A[a_offset + 2], tile_A[a_offset + 3]);
   let a1 = vec4<output_element_t>(tile_A[a_offset + 4], tile_A[a_offset + 5], tile_A[a_offset + 6], tile_A[a_offset + 7]);
-  gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-  up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+  gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+  up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
 #elif component_a == 2
   let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1]);
   let a1 = vec4<output_element_t>(tile_A[a_offset + 2], tile_A[a_offset + 3]);
-  gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-  up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+  gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+  up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
 #elif component_a == 4
   let a0 = tile_A[a_offset];
   let a1 = tile_A[a_offset + 1];
-  gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-  up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+  gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+  up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
 #endif
 #else
   for (var i = 0u; i < component_b; i++) {
@@ -117,26 +118,26 @@ fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> ve
 #if component_a == 1
       let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1], tile_A[a_offset + 2], tile_A[a_offset + 3]);
       let a1 = vec4<output_element_t>(tile_A[a_offset + 4], tile_A[a_offset + 5], tile_A[a_offset + 6], tile_A[a_offset + 7]);
-      gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-      up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+      gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+      up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
       a_offset += 8;
 #elif component_a == 2
       let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1]);
       let a1 = vec4<output_element_t>(tile_A[a_offset + 2], tile_A[a_offset + 3]);
-      gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-      up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+      gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+      up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
       a_offset += 4;
 #elif component_a == 4
       let a0 = tile_A[a_offset];
       let a1 = tile_A[a_offset + 1];
-      gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
-      up_sum += dot(a0, up_b0) + dot(a1, up_b1);
+      gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
+      up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
       a_offset += 2;
 #endif
   }
 #endif
 
-  return vec2<output_element_t>(gate_sum, up_sum);
+  return vec2<f32>(gate_sum, up_sum);
 }
 
 fn process_k_tile(batch: u32, b_global_base: u32, thread_idx: u32, idx: u32, idy: u32, kidx: u32, inv_std: f32) {
@@ -168,8 +169,8 @@ $MAIN {
 
   if (local_idx < tile_size) {
     for (var b = 0u; b < tile_size_k_vec; b++) {
-      gate_inter_results[local_idx][b] = output_element_t(0);
-      up_inter_results[local_idx][b] = output_element_t(0);
+      gate_inter_results[local_idx][b] = f32(0);
+      up_inter_results[local_idx][b] = f32(0);
     }
   }
   workgroupBarrier();
@@ -239,8 +240,8 @@ $MAIN {
   }
 
   if (local_idx < tile_size) {
-    var gate_output_value = output_element_t(0);
-    var up_output_value = output_element_t(0);
+    var gate_output_value = f32(0);
+    var up_output_value = f32(0);
     for (var b = 0u; b < tile_size_k_vec; b++) {
       gate_output_value += gate_inter_results[local_idx][b];
       up_output_value += up_inter_results[local_idx][b];
@@ -249,18 +250,19 @@ $MAIN {
     let output_idx = batch * uniforms.N + b_global;
     if (b_global < uniforms.N) {
 #if has_gate_bias
-      gate_output_value += gate_bias[b_global];
+      gate_output_value += f32(gate_bias[b_global]);
 #endif
 #if has_up_bias
-      up_output_value += up_bias[b_global];
+      up_output_value += f32(up_bias[b_global]);
 #endif
-      let one = output_element_t(1.0);
+      let one = f32(1.0);
 #if activation_kind == 0
       // SiLU(x) = x * sigmoid(x). New activations are added with additional
       // `#elif activation_kind == N` blocks (must match MlpActivationKind).
       let activated_value = gate_output_value * (one / (one + exp(-gate_output_value)));
 #endif
-      output.setByOffset(output_idx, activated_value * up_output_value);
+      // Downcast to the output element type only at the final store.
+      output.setByOffset(output_idx, output_element_t(activated_value * up_output_value));
     }
   }
 }  // MAIN
diff --git a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.wgsl.template b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.wgsl.template
index 60f34e9ef2530..209ccfcd45153 100644
--- a/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.wgsl.template
@@ -21,9 +21,10 @@
 var<workgroup> sum_squared_shared : array<f32, workgroup_size_x>;
 #endif
 var<workgroup> tile_A : array<input_a_value_t, a_length_per_tile>;
-var<workgroup> q_inter_results : array<array<q_output_element_t, tile_size_k_vec + 1>, tile_size>;
-var<workgroup> k_inter_results : array<array<q_output_element_t, tile_size_k_vec + 1>, tile_size>;
-var<workgroup> v_inter_results : array<array<q_output_element_t, tile_size_k_vec + 1>, tile_size>;
+// Accumulate in f32 to avoid f16 overflow (max ~65504) when summing long K dot products.
+var<workgroup> q_inter_results : array<array<f32, tile_size_k_vec + 1>, tile_size>;
+var<workgroup> k_inter_results : array<array<f32, tile_size_k_vec + 1>, tile_size>;
+var<workgroup> v_inter_results : array<array<f32, tile_size_k_vec + 1>, tile_size>;
 
 const default_zero_point = vec4<q_output_element_t>(q_output_element_t(8));
 
@@ -95,8 +96,8 @@ fn loadSHMA(batch: u32, b_global_base: u32, kidx: u32, col: u32, inv_std: f32) {
 
 fn compute_projection_sum(weight: q_b_value_t,
                           scale: q_output_element_t,
-                          idx: u32) -> q_output_element_t {
-  var sum = q_output_element_t(0);
+                          idx: u32) -> f32 {
+  var sum = f32(0);
   var a_offset = idx * (8 / component_a) * component_b;
 #if component_b == 1
   let weight_lower = unpack_nibble_values(weight & 0x0F0F0F0Fu) - default_zero_point;
@@ -106,15 +107,15 @@ fn compute_projection_sum(weight: q_b_value_t,
 #if component_a == 1
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 4);
-  sum += dot(a0, w0) + dot(a1, w1);
+  sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
 #elif component_a == 2
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 2);
-  sum += dot(a0, w0) + dot(a1, w1);
+  sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
 #elif component_a == 4
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 1);
-  sum += dot(a0, w0) + dot(a1, w1);
+  sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
 #endif
 #else
   for (var i = 0u; i < component_b; i++) {
@@ -125,17 +126,17 @@ fn compute_projection_sum(weight: q_b_value_t,
 #if component_a == 1
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 4);
-    sum += dot(a0, w0) + dot(a1, w1);
+    sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
     a_offset += 8;
 #elif component_a == 2
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 2);
-    sum += dot(a0, w0) + dot(a1, w1);
+    sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
     a_offset += 4;
 #elif component_a == 4
   let a0 = load_a_vec4(a_offset);
   let a1 = load_a_vec4(a_offset + 1);
-    sum += dot(a0, w0) + dot(a1, w1);
+    sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
     a_offset += 2;
 #endif
   }
@@ -193,9 +194,9 @@ $MAIN {
 
   if (local_idx < tile_size) {
     for (var b = 0u; b < tile_size_k_vec; b++) {
-      q_inter_results[local_idx][b] = q_output_element_t(0);
-      k_inter_results[local_idx][b] = q_output_element_t(0);
-      v_inter_results[local_idx][b] = q_output_element_t(0);
+      q_inter_results[local_idx][b] = f32(0);
+      k_inter_results[local_idx][b] = f32(0);
+      v_inter_results[local_idx][b] = f32(0);
     }
   }
 
@@ -262,9 +263,9 @@ $MAIN {
 
   if (local_idx < tile_size) {
     let b_global = b_global_base + local_idx;
-    var q_output_value = q_output_element_t(0);
-    var k_output_value = q_output_element_t(0);
-    var v_output_value = q_output_element_t(0);
+    var q_output_value = f32(0);
+    var k_output_value = f32(0);
+    var v_output_value = f32(0);
     for (var b = 0u; b < tile_size_k_vec; b++) {
       q_output_value += q_inter_results[local_idx][b];
       k_output_value += k_inter_results[local_idx][b];