diff --git a/mojo_opset/backends/ttx/kernels/npu/layernorm.py b/mojo_opset/backends/ttx/kernels/npu/layernorm.py index 7bf0dcbda..f779f803f 100644 --- a/mojo_opset/backends/ttx/kernels/npu/layernorm.py +++ b/mojo_opset/backends/ttx/kernels/npu/layernorm.py @@ -95,6 +95,7 @@ def _layernorm_fwd_kernel( b_chunk = tl.load(B_ptr + cols_off, mask=cols_mask, other=0.0).to(tl.float32) x_centered = x_chunk - mean[:, None] + x_centered = tl.where(block_mask, x_centered, 0.0) # Filter out invalid columns var_acc += tl.sum(x_centered * x_centered, axis=1)