feat(torchwave): Standalone Triton port of masked_select_jagged and its test (#18860)

[oerling]

Summary:

Explore Triton as a torchwave codegen target by porting the
tw.masked_select_jagged op to standalone Triton kernels. This does NOT
touch torchwave proper (no WaveGraph, no op registration); it adds a new
self-contained package under torchwave_meta/tests/triton/.

Two variants parallel the two CUDA forms in
minimal_viable_ai/torchwave_meta/Select.cuh:

• multikernel: staged launches (per-chunk count -> scan of partials ->
  write exclusive prefix -> scatter -> new segment lengths), mirroring
  torchwave's multi-kernel execution mode.
• cg: a single persistent kernel carrying the same three stages across
  device-wide atomic barriers -- the Triton analogue of the CUDA
  opBarrier / grid.sync (num_programs kept co-resident so the
  spin-barrier cannot deadlock; programs grid-stride the chunks).

The standalone test parallels the CUDA gtest
minimal_viable_ai/torchwave_meta/tests/KernelTest.cu (SelectKernelTest):
same input generation, the same per-segment CPU reference, and the same
sizes; it validates the Triton kernels against the CPU reference and also
covers the grouped mixed-dtype path.

Authored with Claude.

Reviewed By: Yuhta

Differential Revision: D114379433

[netlify]

✅ Deploy Preview for meta-velox canceled.

Name	Link
🔨 Latest commit	05c4aa0
🔍 Latest deploy log	https://app.netlify.com/projects/meta-velox/deploys/6a9c808d5741d50008dda530

[meta-codesync]

@oerling has exported this pull request. If you are a Meta employee, you can view the originating Diff in D114379433.

[github-actions]

Selective Build Plan

Linux release with adapters is running a full build (PR has a standing approval). See the CI workflows README for what this means.

---

Selective build plan

[github-actions]

CI Failure Analysis

Auto-generated by the CI Failure Analysis workflow. This comment is updated in place each time CI fails on a new commit, so it always reflects the latest run — re-pushing or re-running CI will refresh the analysis below. Last updated 2026-09-05 18:26:49 UTC from workflow run 33982956541.

🟡 Presto Fuzzer — FUZZER Failure View logs

Fuzzer crash: Instance 1 (seed 409347951) aborted with a Folly F14 hash table assertion failure during rehash.

Assertion failure: hp.second == srcChunk->tag(srcI)
Message:
File: /deps/include/folly/container/detail/F14Table.h
Line: 2487
Function: rehashImpl
*** Aborted at 1788632247 (Unix time, try 'date -d @1788632247') ***
*** Signal 6 (SIGABRT) received by PID 124 ***

Instances 2, 3, and 4 (seeds 1001860713, 564408579, 783496755) all passed.

---

Correlation with PR changes:

• Not related. This PR modifies only files under velox/experimental/torchwave/ (single-pass scan variants, dynamic shared memory plumbing, __launch_bounds__, frame reuse). The fuzzer crash is in Folly's F14Table::rehashImpl, an internal hash table operation entirely outside the scope of these changes.

Known issues:

• This is a known flaky failure tracked in #18793 — Presto Fuzzer failing in folly rehashImpl, which reports the identical assertion (hp.second == srcChunk->tag(srcI) at F14Table.h:2487) occurring intermittently on main. The issue notes it is not reproducible on devservers.
• The most recent Fuzzer Jobs run on main had the Presto Fuzzer pass, confirming the intermittent nature.

Reproduce locally:

./velox_expression_fuzzer_test \
    --seed 409347951 \
    --enable_variadic_signatures \
    --velox_fuzzer_enable_complex_types \
    --velox_fuzzer_enable_decimal_type \
    --lazy_vector_generation_ratio 0.2 \
    --common_dictionary_wraps_generation_ratio=0.3 \
    --velox_fuzzer_enable_column_reuse \
    --velox_fuzzer_enable_expression_reuse \
    --max_expression_trees_per_step 2 \
    --retry_with_try \
    --special_forms="and,or,cast,coalesce" \
    --enable_dereference \
    --duration_sec 300

Note: Per #18793, this crash has not been reproducible locally so far.

Recommended fix: No action needed on this PR. This is a pre-existing flaky failure unrelated to the PR changes. Consider re-running the workflow.