[ExecuTorch][WebGPU] Add relu op (shared unary handler; sigmoid adopts make_compute_pipeline)

[JCNTH]

Stack from ghstack (oldest at bottom):

• [ExecuTorch][WebGPU] apply_rotary_emb_hf op tests #20876
• [ExecuTorch][WebGPU] et_vk.apply_rotary_emb_hf: HuggingFace rotate-half RoPE runtime op (unblocks Qwen3) #20875
• [ExecuTorch][WebGPU] et_vk conv2d im2col-GEMM tests #20874
• [ExecuTorch][WebGPU] et_vk conv2d: route standard (groups==1) conv through an im2col tiled GEMM (1.1-2.4x) #20873
• [ExecuTorch][WebGPU] et_vk.sdpa channel-attention tests #20872
• [ExecuTorch][WebGPU] et_vk.sdpa: shape-route QK to a per-entry kernel for channel attention (15-30x faster) #20871
• [ExecuTorch][WebGPU] Migrate quantized_linear to make_compute_pipeline #20870
• [ExecuTorch][WebGPU] Migrate multi-dispatch ops to make_compute_pipeline #20869
• [ExecuTorch][WebGPU] Migrate 9 ops to make_compute_pipeline (layout:"auto") #20868
• [ExecuTorch][WebGPU] Adopt WebGPU layout:"auto" in make_compute_pipeline #20867
• [ExecuTorch][WebGPU] slice_copy native golden test #20866
• [ExecuTorch][WebGPU] slice_copy Double-arg scalar handling + view_copy/alias glue #20865
• [ExecuTorch][WebGPU] relu op tests #20864
• -> [ExecuTorch][WebGPU] Add relu op (shared unary handler; sigmoid adopts make_compute_pipeline) #20863
• [ExecuTorch][WebGPU] max_pool2d op tests #20862
• [ExecuTorch][WebGPU] Add max_pool2d op #20861
• [ExecuTorch][WebGPU] upsample_nearest2d op tests #20860
• [ExecuTorch][WebGPU] Add upsample_nearest2d op #20859
• [ExecuTorch][WebGPU] constant_pad_nd op tests #20858
• [ExecuTorch][WebGPU] Add constant_pad_nd op #20857
• [ExecuTorch][WebGPU] addmm op tests #20856
• [ExecuTorch][WebGPU] Add optimized addmm op #20855
• [ExecuTorch][WebGPU] embedding op tests #20854
• [ExecuTorch][WebGPU] Add embedding op #20853
• [ExecuTorch][WebGPU] SDPA op tests #20852
• [ExecuTorch][WebGPU] Add optimized SDPA op #20851
• [ExecuTorch][WebGPU] conv2d op tests #20850
• [ExecuTorch][WebGPU] Add optimized conv2d op #20849
• [ExecuTorch][WebGPU] linear_fp32 op tests #20848
• [ExecuTorch][WebGPU] Add optimized linear_fp32 op #20847
• [ExecuTorch][WebGPU] layer_norm op tests #20846
• [ExecuTorch][WebGPU] Add optimized layer_norm op #20845
• [ExecuTorch][WebGPU] gelu op tests #20844
• [ExecuTorch][WebGPU] Add optimized gelu op #20843
• [ExecuTorch][WebGPU] Shared pipeline + dispatch-grid + vec4-align helpers (op-suite foundation) #20842

Adds relu to the WebGPU backend via a shared elementwise-unary handler. ReLU is on the SAM2/SAM3 mask-decoder MLP path, so it is needed to delegate those decoders.

Problem — The backend had no aten.relu.default kernel, and sigmoid (the only prior unary op) built its compute pipeline inline rather than through the shared helper — duplicating the bind-group/dispatch boilerplate that a second unary op would repeat.

Solution

• Before: sigmoid was implemented with a bespoke inline pipeline; there was no relu.
• After: a single generic add_unary_op helper (in runtime/ops/sigmoid/UnaryOp.cpp) builds the input/output/params binding and dispatch for any elementwise-unary WGSL; sigmoid_impl and the new relu_impl are thin wrappers over it, so sigmoid now goes through the same utils::make_compute_pipeline path as relu. relu.wgsl is a one-element-per-thread output[idx] = max(input[idx], 0.0).

Implementation

• add_unary_op(graph, in, out, wgsl_source, wg_size_x, op_name) centralizes: the fp32/4-byte-alignment and same-size guards, utils::clamp_workgroup_size + utils::compute_1d_workgroup_count for the 1D dispatch, the wg_size override constant, the uniform (num_elements) via utils::make_uniform, and the three-binding pipeline via utils::make_compute_pipeline.
• Dynamic shapes are supported: a graph.add_tensor_resize_hook recomputes num_elements, rewrites the uniform via wgpuQueueWriteBuffer, and updates the dispatch's workgroup count for the live shape; the graph owns the uniform buffer so the hook can rewrite it.
• Both ops self-register: aten.sigmoid.default -> sigmoid_impl and aten.relu.default -> relu_impl.
• Mirrors Vulkan backends/vulkan/runtime/graph/ops/impl/UnaryOp.cpp (add_unary_op_node); Vulkan expresses relu as clamp(0, inf), whereas this kernel uses a direct max(x, 0.0).

Constraints — fp32 only (both operands 4-byte aligned); input and output must have identical byte size (same-shape elementwise); 1D dispatch only (throws past the 65535 workgroup cap).

Co-authored-with: Claude Code.

Differential Revision: D110836664

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/20863

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

✅ You can merge normally! (1 Unrelated Failure)

As of commit a219bb5 with merge base aceeb40 ():

BROKEN TRUNK - The following job failed but were present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

• pull / unittest-nxp-neutron / linux-job (gh) (trunk failure)
  ##[error]The operation was canceled.

This comment was automatically generated by Dr. CI and updates every 15 minutes.

[github-actions]

This PR needs a release notes: label

If your change should be included in the release notes (i.e. would users of this library care about this change?), please use a label starting with release notes:. This helps us keep track and include your important work in the next release notes.

To add a label, you can comment to pytorchbot, for example
@pytorchbot label "release notes: none"

For more information, see
https://github.com/pytorch/pytorch/wiki/PyTorch-AutoLabel-Bot#why-categorize-for-release-notes-and-how-does-it-work.