Summary:


# Problem
Diff 1/3 implementing [#203385](#203385). We want to generalize tiling from [X, Y, Z] to [X, Y, Z0, Z1, ..., ZN]. 

# Solution
This prepares for that by renaming the existing `z` axis to `z0_` and `ZBLOCK` to `Z0_BLOCK`. Update Inductor and its MTIA consumers together so the intermediate revision stays compatible. 

# Alternatives considered
Naming tradeoff:
- `X`, `Y`, `Z0_`, `Z1_`, ... preserves the heavily used X/Y conventions and makes every further pointwise axis part of one numbered family, like `r0_`, `r1_`, ... . Ordering, detection, and max-block lookup can treat that family uniformly. The name itself has no performance advantage.
- Keeping `X`, `Y`, `Z` and adding `W0_`, `W1_`, ... would avoid this rename. The measured unformatted rename changes 156 lines across 24 files (156 additions plus 156 deletions, 312 lines of diff churn, zero net code growth). The separate Z axis would need an estimated 0-10 extra helper lines.
- Renaming every pointwise axis to `X0_`, `X1_`, ... would make naming uniform but require an estimated 1,000-1,500 additional line replacements (2,000-3,000 lines of diff churn) within the files already touched by this stack. The scan found 1,163 currently unchanged X/Y reference lines; that is a candidate count, not an exact migration count. Other consumers would increase the scope. Uniform naming could shorten the helpers by roughly 10-25 lines.

# Testing

- Update existing generated-code and indexing assertions for `Z0_BLOCK`/`z0_`, including strided pointwise/reduction kernels, loop ordering, CUDA regressions, and native batched matmul.
- Update existing autotuning configuration tests for the renamed Z axis while preserving block-size and reduction expectations.
- Update existing MTIA heuristic, scheduling metadata, and generated-kernel tests for the renamed axis. This rename adds no new test scenarios.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

Test Plan:
`arc lint --only-changed --never-apply-patches --output json`

Draft submission precedes the remaining validation, as requested. Test results will be added after validation.

```bash
buck2 test mode/opt fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'test_boundary_check_block_multiple_(True|False)_ynumel_exceed_ygrid_size_False_include_z_True_cuda'
```
https://www.internalfb.com/intern/testinfra/testrun/9007199470724220

7 passing results, no failures. The former `ZBLOCK`/`Z0_BLOCK` fixed-config argument mismatch is resolved.

```bash
buck2 test mode/opt fbcode//mtia/compiler/graph_compiler/inductor/test:test_compiler_subprocess -- --regex 'test_timeout_kills_descendants|test_timeout_does_not_wait_for_escaped_pipe_holder' --run-disabled
```
https://www.internalfb.com/intern/testinfra/testrun/32932572330111445

Both previously disabled subprocess tests passed.

Differential Revision: D130812310

Summary:

# Problem

Diff 2/3 implementing [#203385](#203385), which requests pointwise tiling beyond three dimensions. After the Z-axis rename in [D130812310](https://www.internalfb.com/diff/D130812310) ([PyTorch PR #204273](#204273)), Inductor still hard-codes the axis list and the three-axis limit in codegen and autotuning. These assumptions need to be removed before [D130812313](https://www.internalfb.com/diff/D130812313) can enable additional axes.

# Solution

Centralize pointwise axis enumeration, block-name conversion, ordering, and maximum block lookup in shared helpers. Use those helpers in range-tree construction, codegen traversal, and autotuning. Replace the literal three-axis clamp with the size of the supported axis collection, which remains `x`, `y`, and `z0_` in this diff.

This is a behavior-preserving refactor: existing kernels keep their axis order, supported rank, candidate configurations, and tuning constants. Add unit coverage for the generalized axis handling.

# Testing

- Add prefix-helper unit coverage for numbered Z axes, numeric ordering (including `z10_`), separation from reduction axes, rank-based prefix generation, block suffix conversion, and shared Z block limits.
- Generalize the existing autotuning block-limit test to check every emitted block field.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

Test Plan:
`arc lint -a`: no lint issues.
`arc pyrefly check-owning-targets caffe2/torch/_inductor/codegen/triton.py`: the owning target does not enable type checking.

https://www.internalfb.com/intern/testinfra/testrun/9288674450596823
Ran foreach tests on this refactor commit, without ND descendants: 132 passed. All 20 codegen and all 48 Athena hardware cases that failed in v0.4 now pass. Twenty of the corresponding QModel cases pass; 28 fail in setup because the `workloadd` DBus socket is missing, before Inductor runs.

https://www.internalfb.com/intern/testinfra/testrun/33776997260249910
Retried the failing QModel `foreach_erfc` operator in isolation: both dtype cases passed (three passing results including the listing).

The v0.4 foreach failures reached `_trace_symbol_prefix` and raised `NameError` for the missing `re` import. The fix is in shared OSS code; no MTIA extension API change is required for these failures.

Differential Revision: D130812312

Summary:


# Problem
Diff 1/3 implementing [#203385](#203385). We want to generalize tiling from [X, Y, Z] to [X, Y, Z0, Z1, ..., ZN]. 

# Solution
This prepares for that by renaming the existing `z` axis to `z0_` and `ZBLOCK` to `Z0_BLOCK`. Update Inductor and its MTIA consumers together so the intermediate revision stays compatible. 

# Alternatives considered
Naming tradeoff:
- `X`, `Y`, `Z0_`, `Z1_`, ... preserves the heavily used X/Y conventions and makes every further pointwise axis part of one numbered family, like `r0_`, `r1_`, ... . Ordering, detection, and max-block lookup can treat that family uniformly. The name itself has no performance advantage.
- Keeping `X`, `Y`, `Z` and adding `W0_`, `W1_`, ... would avoid this rename. The measured unformatted rename changes 156 lines across 24 files (156 additions plus 156 deletions, 312 lines of diff churn, zero net code growth). The separate Z axis would need an estimated 0-10 extra helper lines.
- Renaming every pointwise axis to `X0_`, `X1_`, ... would make naming uniform but require an estimated 1,000-1,500 additional line replacements (2,000-3,000 lines of diff churn) within the files already touched by this stack. The scan found 1,163 currently unchanged X/Y reference lines; that is a candidate count, not an exact migration count. Other consumers would increase the scope. Uniform naming could shorten the helpers by roughly 10-25 lines.

# Testing

- Update existing generated-code and indexing assertions for `Z0_BLOCK`/`z0_`, including strided pointwise/reduction kernels, loop ordering, CUDA regressions, and native batched matmul.
- Update existing autotuning configuration tests for the renamed Z axis while preserving block-size and reduction expectations.
- Update existing MTIA heuristic, scheduling metadata, and generated-kernel tests for the renamed axis. This rename adds no new test scenarios.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

Test Plan:
`arc lint --only-changed --never-apply-patches --output json`

Draft submission precedes the remaining validation, as requested. Test results will be added after validation.

```bash
buck2 test mode/opt fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'test_boundary_check_block_multiple_(True|False)_ynumel_exceed_ygrid_size_False_include_z_True_cuda'
```
https://www.internalfb.com/intern/testinfra/testrun/9007199470724220

7 passing results, no failures. The former `ZBLOCK`/`Z0_BLOCK` fixed-config argument mismatch is resolved.

```bash
buck2 test mode/opt fbcode//mtia/compiler/graph_compiler/inductor/test:test_compiler_subprocess -- --regex 'test_timeout_kills_descendants|test_timeout_does_not_wait_for_escaped_pipe_holder' --run-disabled
```
https://www.internalfb.com/intern/testinfra/testrun/32932572330111445

Both previously disabled subprocess tests passed.

Differential Revision: D130812310

Summary:

# Problem

Diff 2/3 implementing [#203385](#203385), which requests pointwise tiling beyond three dimensions. After the Z-axis rename in [D130812310](https://www.internalfb.com/diff/D130812310) ([PyTorch PR #204273](#204273)), Inductor still hard-codes the axis list and the three-axis limit in codegen and autotuning. These assumptions need to be removed before [D130812313](https://www.internalfb.com/diff/D130812313) can enable additional axes.

# Solution

Centralize pointwise axis enumeration, block-name conversion, ordering, and maximum block lookup in shared helpers. Use those helpers in range-tree construction, codegen traversal, and autotuning. Replace the literal three-axis clamp with the size of the supported axis collection, which remains `x`, `y`, and `z0_` in this diff.

This is a behavior-preserving refactor: existing kernels keep their axis order, supported rank, candidate configurations, and tuning constants. Add unit coverage for the generalized axis handling.

# Testing

- Add prefix-helper unit coverage for numbered Z axes, numeric ordering (including `z10_`), separation from reduction axes, rank-based prefix generation, block suffix conversion, and shared Z block limits.
- Generalize the existing autotuning block-limit test to check every emitted block field.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

Test Plan:
`arc lint -a`: no lint issues.
`arc pyrefly check-owning-targets caffe2/torch/_inductor/codegen/triton.py`: the owning target does not enable type checking.

https://www.internalfb.com/intern/testinfra/testrun/9288674450596823
Ran foreach tests on this refactor commit, without ND descendants: 132 passed. All 20 codegen and all 48 Athena hardware cases that failed in v0.4 now pass. Twenty of the corresponding QModel cases pass; 28 fail in setup because the `workloadd` DBus socket is missing, before Inductor runs.

https://www.internalfb.com/intern/testinfra/testrun/33776997260249910
Retried the failing QModel `foreach_erfc` operator in isolation: both dtype cases passed (three passing results including the listing).

The v0.4 foreach failures reached `_trace_symbol_prefix` and raised `NameError` for the missing `re` import. The fix is in shared OSS code; no MTIA extension API change is required for these failures.

Differential Revision: D130812312

Summary:
Pull Request resolved: #204273

# Problem
Diff 1/3 implementing [#203385](#203385). We want to generalize tiling from [X, Y, Z] to [X, Y, Z0, Z1, ..., ZN].

# Solution
This prepares for that by renaming the existing `z` axis to `z0_` and `ZBLOCK` to `Z0_BLOCK`. Update Inductor and its MTIA consumers together so the intermediate revision stays compatible.

# Alternatives considered
Naming tradeoff:
- `X`, `Y`, `Z0_`, `Z1_`, ... preserves the heavily used X/Y conventions and makes every further pointwise axis part of one numbered family, like `r0_`, `r1_`, ... . Ordering, detection, and max-block lookup can treat that family uniformly. The name itself has no performance advantage.
- Keeping `X`, `Y`, `Z` and adding `W0_`, `W1_`, ... would avoid this rename. The measured unformatted rename changes 156 lines across 24 files (156 additions plus 156 deletions, 312 lines of diff churn, zero net code growth). The separate Z axis would need an estimated 0-10 extra helper lines.
- Renaming every pointwise axis to `X0_`, `X1_`, ... would make naming uniform but require an estimated 1,000-1,500 additional line replacements (2,000-3,000 lines of diff churn) within the files already touched by this stack. The scan found 1,163 currently unchanged X/Y reference lines; that is a candidate count, not an exact migration count. Other consumers would increase the scope. Uniform naming could shorten the helpers by roughly 10-25 lines.

# Testing

- Update existing generated-code and indexing assertions for `Z0_BLOCK`/`z0_`, including strided pointwise/reduction kernels, loop ordering, CUDA regressions, and native batched matmul.
- Update existing autotuning configuration tests for the renamed Z axis while preserving block-size and reduction expectations.
- Update existing MTIA heuristic, scheduling metadata, and generated-kernel tests for the renamed axis. This rename adds no new test scenarios.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

#build_oncall[mtia_sw_pt2]

Test Plan:
`arc lint --only-changed --never-apply-patches --output json`

Draft submission precedes the remaining validation, as requested. Test results will be added after validation.

```bash
buck2 test mode/opt fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'test_boundary_check_block_multiple_(True|False)_ynumel_exceed_ygrid_size_False_include_z_True_cuda'
```
https://www.internalfb.com/intern/testinfra/testrun/9007199470724220

7 passing results, no failures. The former `ZBLOCK`/`Z0_BLOCK` fixed-config argument mismatch is resolved.

```bash
buck2 test mode/opt fbcode//mtia/compiler/graph_compiler/inductor/test:test_compiler_subprocess -- --regex 'test_timeout_kills_descendants|test_timeout_does_not_wait_for_escaped_pipe_holder' --run-disabled
```
https://www.internalfb.com/intern/testinfra/testrun/32932572330111445

Both previously disabled subprocess tests passed.

Differential Revision: D130812310

Summary:
Pull Request resolved: #204281

# Problem

Diff 2/3 implementing [#203385](#203385), which requests pointwise tiling beyond three dimensions. After the Z-axis rename in [D130812310](https://www.internalfb.com/diff/D130812310) ([PyTorch PR #204273](#204273)), Inductor still hard-codes the axis list and the three-axis limit in codegen and autotuning. These assumptions need to be removed before [D130812313](https://www.internalfb.com/diff/D130812313) can enable additional axes.

# Solution

Centralize pointwise axis enumeration, block-name conversion, ordering, and maximum block lookup in shared helpers. Use those helpers in range-tree construction, codegen traversal, and autotuning. Replace the literal three-axis clamp with the size of the supported axis collection, which remains `x`, `y`, and `z0_` in this diff.

This is a behavior-preserving refactor: existing kernels keep their axis order, supported rank, candidate configurations, and tuning constants. Add unit coverage for the generalized axis handling.

# Testing

- Add prefix-helper unit coverage for numbered Z axes, numeric ordering (including `z10_`), separation from reduction axes, rank-based prefix generation, block suffix conversion, and shared Z block limits.
- Generalize the existing autotuning block-limit test to check every emitted block field.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313) — [PyTorch PR #204324](#204324)

Test Plan:
`arc lint -a`: no lint issues.
`arc pyrefly check-owning-targets caffe2/torch/_inductor/codegen/triton.py`: the owning target does not enable type checking.

https://www.internalfb.com/intern/testinfra/testrun/9288674450596823
Ran foreach tests on this refactor commit, without ND descendants: 132 passed. All 20 codegen and all 48 Athena hardware cases that failed in v0.4 now pass. Twenty of the corresponding QModel cases pass; 28 fail in setup because the `workloadd` DBus socket is missing, before Inductor runs.

https://www.internalfb.com/intern/testinfra/testrun/33776997260249910
Retried the failing QModel `foreach_erfc` operator in isolation: both dtype cases passed (three passing results including the listing).

The v0.4 foreach failures reached `_trace_symbol_prefix` and raised `NameError` for the missing `re` import. The fix is in shared OSS code; no MTIA extension API change is required for these failures.

Differential Revision: D130812312

Summary:
Pull Request resolved: #204324

# Problem

Diff 3/3 implementing [#203385](#203385), which requests pointwise tiling beyond three dimensions. [D130812312](https://www.internalfb.com/diff/D130812312) ([PyTorch PR #204281](#204281)) generalizes axis handling, but Inductor still cannot generate and launch kernels with four or more pointwise tiles. Triton exposes only three program-ID dimensions, and existing autotuning candidates provide at most three pointwise block budgets.

# Solution

Enable numbered Z axes through the existing `triton.max_tiles` setting and `prefer_nd_tiling` selection. For example, a 4D tile uses `z1_`, `z0_`, `y`, and `x`. Pack the numbered Z tile grids into Triton's third launch dimension and reconstruct only their logical coordinates with division and remainder. Reuse the existing X/Y launch dimensions and program IDs. Use `GridND` for ordinary 3D and higher-rank kernels, retaining the native batched-matmul axis mapping. Generate pointwise prefixes with the existing helpers and use `TRITON_LAUNCH_GRID_NDIM` for the legacy launch-rank boundary.

Generalize block-budget expansion by distributing the final supplied budget over the remaining pointwise axes, inner first. Reuse the existing candidate lists, tuning constants, and reduction-budget expansion. Carry numbered Z block sizes through Python and C++ wrappers using the existing `zblocks` field.

Add numerical coverage for 4D broadcasts, permute/reshape refinement, and dynamic shapes, plus unit coverage for block-budget expansion and packed launch coordinates. The default tile count remains unchanged.

# Testing

- Add launch-grid unit tests for 3D/4D/5D kernels, checking that Z packing reconstructs every logical tile exactly once and preserves X/Y program IDs with partial boundary tiles.
- Add block-budget unit tests for ND expansion, unit dimensions, block limits, malformed inputs, and unchanged existing 1D/2D/3D configurations and warp counts.
- Add reduction configuration regression tests for existing reduction-budget expansion and joint pointwise/reduction block budgets.
- Extend prefix-helper unit coverage to verify pointwise block products and axis exclusions.
- Add numerical tests for 4D broadcasts, 5D permute/reshape refinement, and dynamic shapes (including empty dimensions); tensor descriptor variants also assert the requested numbered Z axes and tensor descriptors in generated kernels.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Enable ND pointwise tiling](https://www.internalfb.com/diff/D130812313)

Test Plan:
```lang=bash
buck2 test mode/opt --agent-context direct_call=true --agent-context 'reason=Validate unified GridND for 3D and ND launches, unchanged autotuning budgets, and Hopper pointwise tensor descriptors' fbcode//caffe2/test/inductor:simd_range_trees fbcode//caffe2/test/inductor:triton_heuristics fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'TestSIMDRangeTrees|pointwise_config|pointwise_budget|existing_pointwise|reduction_config_budget|test_nd_pointwise|test_pointwise_|test_broadcast_'
```
https://www.internalfb.com/intern/testinfra/testrun/20547673367950394

123 passed, 0 failed; 32 skipped and one already-disabled test omitted. Covers 3D/4D/5D launch coordinates, existing and expanded block budgets, pointwise broadcasts, dynamic shapes, and 5D Hopper tensor descriptors.

Differential Revision: D130812313

Summary:
# Problem
Diff 1/3 implementing [#203385](#203385). We want to generalize tiling from [X, Y, Z] to [X, Y, Z0, Z1, ..., ZN].

# Solution
This prepares for that by renaming the existing `z` axis to `z0_` and `ZBLOCK` to `Z0_BLOCK`. Update Inductor and its MTIA consumers together so the intermediate revision stays compatible.

# Alternatives considered
Naming tradeoff:
- `X`, `Y`, `Z0_`, `Z1_`, ... preserves the heavily used X/Y conventions and makes every further pointwise axis part of one numbered family, like `r0_`, `r1_`, ... . Ordering, detection, and max-block lookup can treat that family uniformly. The name itself has no performance advantage.
- Keeping `X`, `Y`, `Z` and adding `W0_`, `W1_`, ... would avoid this rename. The measured unformatted rename changes 156 lines across 24 files (156 additions plus 156 deletions, 312 lines of diff churn, zero net code growth). The separate Z axis would need an estimated 0-10 extra helper lines.
- Renaming every pointwise axis to `X0_`, `X1_`, ... would make naming uniform but require an estimated 1,000-1,500 additional line replacements (2,000-3,000 lines of diff churn) within the files already touched by this stack. The scan found 1,163 currently unchanged X/Y reference lines; that is a candidate count, not an exact migration count. Other consumers would increase the scope. Uniform naming could shorten the helpers by roughly 10-25 lines.

# Testing

- Update existing generated-code and indexing assertions for `Z0_BLOCK`/`z0_`, including strided pointwise/reduction kernels, loop ordering, CUDA regressions, and native batched matmul.
- Update existing autotuning configuration tests for the renamed Z axis while preserving block-size and reduction expectations.
- Update existing MTIA heuristic, scheduling metadata, and generated-kernel tests for the renamed axis. This rename adds no new test scenarios.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Support ND pointwise and reduction tiling](https://www.internalfb.com/diff/D130812313) — [PyTorch PR #204324](#204324)

Test Plan:
`arc lint --only-changed --never-apply-patches --output json`

Draft submission precedes the remaining validation, as requested. Test results will be added after validation.

```bash
buck2 test mode/opt fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'test_boundary_check_block_multiple_(True|False)_ynumel_exceed_ygrid_size_False_include_z_True_cuda'
```
https://www.internalfb.com/intern/testinfra/testrun/9007199470724220

7 passing results, no failures. The former `ZBLOCK`/`Z0_BLOCK` fixed-config argument mismatch is resolved.

```bash
buck2 test mode/opt fbcode//mtia/compiler/graph_compiler/inductor/test:test_compiler_subprocess -- --regex 'test_timeout_kills_descendants|test_timeout_does_not_wait_for_escaped_pipe_holder' --run-disabled
```
https://www.internalfb.com/intern/testinfra/testrun/32932572330111445

Both previously disabled subprocess tests passed.

ND stack split validation (2026-11-10):
H100 validation: https://www.internalfb.com/intern/testinfra/testrun/35465847120931160

- 24 independent padded reduction variants passed: four/five axes, persistent/looped, with/without last-axis broadcast, CUDA block pointers and host/device tensor descriptors.
- 11 ND pointwise/joint cases passed, including four pointwise plus four reduction axes.
- 14 existing and extended configuration cases passed.
- Eight numerical examples and generated kernels captured with `TORCH_COMPILE_DEBUG=1` passed, including descriptor generation and the `max_tiles=2` counterexamples.
- Eight Triton CPU descriptor cases skipped because that backend is unavailable on this runner.

`arc lint -a`: passed.
Conda `lintrunner -a --paths-from /tmp/nd_lintrunner_oss_paths.txt --skip CODEOWNERS_TAXONOMY,PYREFLY --output=json`: passed. CODEOWNERS_TAXONOMY needs a Git checkout and PYREFLY needs the matching built OSS environment; these previously agreed environment exclusions remain in place.

Differential Revision: D130812310

Summary:
Pull Request resolved: #204273

# Problem
Diff 1/3 implementing [#203385](#203385). We want to generalize tiling from [X, Y, Z] to [X, Y, Z0, Z1, ..., ZN].

# Solution
This prepares for that by renaming the existing `z` axis to `z0_` and `ZBLOCK` to `Z0_BLOCK`. Update Inductor and its MTIA consumers together so the intermediate revision stays compatible.

# Alternatives considered
Naming tradeoff:
- `X`, `Y`, `Z0_`, `Z1_`, ... preserves the heavily used X/Y conventions and makes every further pointwise axis part of one numbered family, like `r0_`, `r1_`, ... . Ordering, detection, and max-block lookup can treat that family uniformly. The name itself has no performance advantage.
- Keeping `X`, `Y`, `Z` and adding `W0_`, `W1_`, ... would avoid this rename. The measured unformatted rename changes 156 lines across 24 files (156 additions plus 156 deletions, 312 lines of diff churn, zero net code growth). The separate Z axis would need an estimated 0-10 extra helper lines.
- Renaming every pointwise axis to `X0_`, `X1_`, ... would make naming uniform but require an estimated 1,000-1,500 additional line replacements (2,000-3,000 lines of diff churn) within the files already touched by this stack. The scan found 1,163 currently unchanged X/Y reference lines; that is a candidate count, not an exact migration count. Other consumers would increase the scope. Uniform naming could shorten the helpers by roughly 10-25 lines.

# Testing

- Update existing generated-code and indexing assertions for `Z0_BLOCK`/`z0_`, including strided pointwise/reduction kernels, loop ordering, CUDA regressions, and native batched matmul.
- Update existing autotuning configuration tests for the renamed Z axis while preserving block-size and reduction expectations.
- Update existing MTIA heuristic, scheduling metadata, and generated-kernel tests for the renamed axis. This rename adds no new test scenarios.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Support ND pointwise and reduction tiling](https://www.internalfb.com/diff/D130812313) — [PyTorch PR #204324](#204324)

Test Plan:
`arc lint --only-changed --never-apply-patches --output json`

Draft submission precedes the remaining validation, as requested. Test results will be added after validation.

```bash
buck2 test mode/opt fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'test_boundary_check_block_multiple_(True|False)_ynumel_exceed_ygrid_size_False_include_z_True_cuda'
```
https://www.internalfb.com/intern/testinfra/testrun/9007199470724220

7 passing results, no failures. The former `ZBLOCK`/`Z0_BLOCK` fixed-config argument mismatch is resolved.

```bash
buck2 test mode/opt fbcode//mtia/compiler/graph_compiler/inductor/test:test_compiler_subprocess -- --regex 'test_timeout_kills_descendants|test_timeout_does_not_wait_for_escaped_pipe_holder' --run-disabled
```
https://www.internalfb.com/intern/testinfra/testrun/32932572330111445

Both previously disabled subprocess tests passed.

ND stack split validation (2026-11-10):
H100 validation: https://www.internalfb.com/intern/testinfra/testrun/35465847120931160

- 24 independent padded reduction variants passed: four/five axes, persistent/looped, with/without last-axis broadcast, CUDA block pointers and host/device tensor descriptors.
- 11 ND pointwise/joint cases passed, including four pointwise plus four reduction axes.
- 14 existing and extended configuration cases passed.
- Eight numerical examples and generated kernels captured with `TORCH_COMPILE_DEBUG=1` passed, including descriptor generation and the `max_tiles=2` counterexamples.
- Eight Triton CPU descriptor cases skipped because that backend is unavailable on this runner.

`arc lint -a`: passed.
Conda `lintrunner -a --paths-from /tmp/nd_lintrunner_oss_paths.txt --skip CODEOWNERS_TAXONOMY,PYREFLY --output=json`: passed. CODEOWNERS_TAXONOMY needs a Git checkout and PYREFLY needs the matching built OSS environment; these previously agreed environment exclusions remain in place.

Differential Revision: D130812310

Summary:
# Problem

Diff 2/3 implementing [#203385](#203385), which requests pointwise tiling beyond three dimensions. After the Z-axis rename in [D130812310](https://www.internalfb.com/diff/D130812310) ([PyTorch PR #204273](#204273)), Inductor still hard-codes the axis list and the three-axis limit in codegen and autotuning. These assumptions need to be removed before [D130812313](https://www.internalfb.com/diff/D130812313) can enable additional axes.

# Solution

This diff replaces any hard-coded (x, y, z) limits with reusable helper functions which adapt to `triton.max_tiles`. It does not actually support `triton.max_tiles > 3`--that is handled in the next diff. This is a pure refactor, no functional change. Reduction offsets, pointer advancement, descriptor ownership, and configuration expansion use the same rank-based helpers; pointwise and reduction selection retain their existing limits until diff 3.

# Testing

- Add prefix-helper unit coverage for numbered Z axes, numeric ordering (including `z10_`), separation from reduction axes, rank-based prefix generation, block suffix conversion, and shared Z block limits.
- Generalize the existing autotuning block-limit test to check every emitted block field.
- Generalize the shared strided-block test assertions to check the complete numbered pointwise and reduction axis sets; keep the existing numerical cases unchanged.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Support ND pointwise and reduction tiling](https://www.internalfb.com/diff/D130812313) — [PyTorch PR #204324](#204324)

Test Plan:
`arc lint -a`: no lint issues.
`arc pyrefly check-owning-targets caffe2/torch/_inductor/codegen/triton.py`: the owning target does not enable type checking.

https://www.internalfb.com/intern/testinfra/testrun/9288674450596823
Ran foreach tests on this refactor commit, without ND descendants: 132 passed. All 20 codegen and all 48 Athena hardware cases that failed in v0.4 now pass. Twenty of the corresponding QModel cases pass; 28 fail in setup because the `workloadd` DBus socket is missing, before Inductor runs.

https://www.internalfb.com/intern/testinfra/testrun/33776997260249910
Retried the failing QModel `foreach_erfc` operator in isolation: both dtype cases passed (three passing results including the listing).

The v0.4 foreach failures reached `_trace_symbol_prefix` and raised `NameError` for the missing `re` import. The fix is in shared OSS code; no MTIA extension API change is required for these failures.

Generic reduction refactors moved here; the pointwise-only child still rejects four reduction tilings and caps the padded reduction at three axes: https://www.internalfb.com/intern/testinfra/testrun/20829148344842151

ND stack split validation (2026-11-10):
H100 validation: https://www.internalfb.com/intern/testinfra/testrun/35465847120931160

- 24 independent padded reduction variants passed: four/five axes, persistent/looped, with/without last-axis broadcast, CUDA block pointers and host/device tensor descriptors.
- 11 ND pointwise/joint cases passed, including four pointwise plus four reduction axes.
- 14 existing and extended configuration cases passed.
- Eight numerical examples and generated kernels captured with `TORCH_COMPILE_DEBUG=1` passed, including descriptor generation and the `max_tiles=2` counterexamples.
- Eight Triton CPU descriptor cases skipped because that backend is unavailable on this runner.

`arc lint -a`: passed.
Conda `lintrunner -a --paths-from /tmp/nd_lintrunner_oss_paths.txt --skip CODEOWNERS_TAXONOMY,PYREFLY --output=json`: passed. CODEOWNERS_TAXONOMY needs a Git checkout and PYREFLY needs the matching built OSS environment; these previously agreed environment exclusions remain in place.

Differential Revision: D130812312

Summary:
# Problem

Diff 3/3 implementing [#203385](#203385), which requests tiling beyond three dimensions. [D130812312](https://www.internalfb.com/diff/D130812312) ([PyTorch PR #204281](#204281)) generalizes axis handling while preserving the existing selection limits. Inductor still cannot retain four or more independent pointwise or reduction axes in generated kernels.

# Solution

- Expand pointwise tiling to support arbitrary `triton.max_tiles`.
- Extend reduction tiling to additional numbered axes, including joint ND pointwise/reduction kernels.
- Reuse the existing block size heuristics and autotuning candidates across the added dimensions.

The generic codegen refactors are in diff 2. The default tile count and tuning constants remain unchanged.

# Example kernels

These examples were executed on an H100 with `TORCH_COMPILE_DEBUG=1` and checked against eager PyTorch. The generated Triton function bodies below come from `output_code.py`; Python launch wrappers and autotuning decorators are omitted. Tensor descriptors are explicitly enabled.

## 4D permute

Copy a padded, permuted 4D tensor into contiguous storage. The input uses a descriptor and the kernel has `Z1_BLOCK`.

```python
import torch

def fn(x):
    return x.permute(1, 3, 0, 2).contiguous()

x = torch.randn(3, 5, 7, 32, device="cuda")[..., :16]
compiled = torch.compile(fn, options={
    "triton.max_tiles": 4,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(x)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_poi_fused_clone_permute_0(in_ptr0, out_ptr0, z1_numel, z0_numel, ynumel, xnumel, Z1_BLOCK : tl.constexpr, Z0_BLOCK : tl.constexpr, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    z1_numel = 5
    z0_numel = 16
    ynumel = 3
    xnumel = 7
    z1_offset = (tl.program_id(2) // tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z1_BLOCK
    z1_index = z1_offset + tl.arange(0, Z1_BLOCK)[:, None, None, None]
    z1_mask = z1_index < z1_numel
    z0_offset = (tl.program_id(2) % tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z0_BLOCK
    z0_index = z0_offset + tl.arange(0, Z0_BLOCK)[None, :, None, None]
    z0_mask = z0_index < z0_numel
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, None, :, None]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[None, None, None, :]
    xmask = xindex < xnumel
    x3 = xindex
    y2 = yindex
    z0_1 = z0_index
    z1_0 = z1_index
    tmp0 = tl.trans(tl.make_tensor_descriptor(in_ptr0, shape=[3, 5, 7, 16], strides=[1120, 224, 32, 1], block_shape=[YBLOCK, Z1_BLOCK, XBLOCK, Z0_BLOCK]).load([yoffset, z1_offset, xoffset, z0_offset]), [1, 3, 0, 2])
    tl.store(out_ptr0 + (x3 + 7*y2 + 21*z0_1 + 336*z1_0), tmp0, xmask & ymask & z0_mask & z1_mask)
```

## 4D broadcast

Add tensors that broadcast along alternating axes. Both inputs and the output use descriptors.

```python
import torch

def fn(a, b):
    return a + b

a = torch.randn(8, 1, 8, 1, device="cuda")
b = torch.randn(1, 8, 1, 16, device="cuda")
compiled = torch.compile(fn, options={
    "triton.max_tiles": 4,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(a, b)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_poi_fused_add_0(in_ptr0, in_ptr1, out_ptr0, z1_numel, z0_numel, ynumel, xnumel, Z1_BLOCK : tl.constexpr, Z0_BLOCK : tl.constexpr, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    z1_numel = 8
    z0_numel = 8
    ynumel = 8
    xnumel = 16
    z1_offset = (tl.program_id(2) // tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z1_BLOCK
    z1_index = z1_offset + tl.arange(0, Z1_BLOCK)[:, None, None, None]
    z1_mask = z1_index < z1_numel
    z0_offset = (tl.program_id(2) % tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z0_BLOCK
    z0_index = z0_offset + tl.arange(0, Z0_BLOCK)[None, :, None, None]
    z0_mask = z0_index < z0_numel
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, None, :, None]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[None, None, None, :]
    xmask = xindex < xnumel
    y2 = yindex
    z1_0 = z1_index
    x3 = xindex
    z0_1 = z0_index
    tmp0 = tl.trans(tl.make_tensor_descriptor(in_ptr0, shape=[8, 8], strides=[8, 1], block_shape=[Z1_BLOCK, YBLOCK]).load([z1_offset, yoffset])[:, :, None, None], [0, 2, 1, 3])
    tmp1 = tl.trans(tl.make_tensor_descriptor(in_ptr1, shape=[8, 16], strides=[16, 1], block_shape=[Z0_BLOCK, XBLOCK]).load([z0_offset, xoffset])[:, :, None, None], [2, 0, 3, 1])
    tmp2 = tmp0 + tmp1
    tl.make_tensor_descriptor(out_ptr0, shape=[8, 8, 8, 16], strides=[1024, 128, 16, 1], block_shape=[Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).store([z1_offset, z0_offset, yoffset, xoffset], tl.broadcast_to(tmp2, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).to(tl.float32))
```

## Fused broadcast multiply + sum

Retain four pointwise axes while reducing the final axis. The single kernel has `Z1_BLOCK` and `R0_BLOCK`, and uses descriptors for both inputs and the output.

```python
import torch

def fn(a, b):
    return (a * b).sum(dim=-1)

a = torch.randn(8, 1, 8, 1, 32, device="cuda")
b = torch.randn(1, 8, 1, 8, 32, device="cuda")
compiled = torch.compile(fn, options={
    "triton.max_tiles": 4,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(a, b)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_per_fused_mul_sum_0(in_ptr0, in_ptr1, out_ptr0, z1_numel, z0_numel, ynumel, xnumel, r0_numel, Z1_BLOCK : tl.constexpr, Z0_BLOCK : tl.constexpr, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    z1_numel = 8
    z0_numel = 8
    ynumel = 8
    xnumel = 8
    r0_numel = 32
    R0_BLOCK: tl.constexpr = 32
    rnumel = r0_numel
    RBLOCK: tl.constexpr = R0_BLOCK
    z1_offset = (tl.program_id(2) // tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z1_BLOCK
    z1_index = z1_offset + tl.arange(0, Z1_BLOCK)[:, None, None, None, None]
    z1_mask = z1_index < z1_numel
    z0_offset = (tl.program_id(2) % tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z0_BLOCK
    z0_index = z0_offset + tl.arange(0, Z0_BLOCK)[None, :, None, None, None]
    z0_mask = z0_index < z0_numel
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, None, :, None, None]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[None, None, None, :, None]
    xmask = xindex < xnumel
    r0_index = tl.arange(0, R0_BLOCK)[None, None, None, None, :]
    r0_offset = 0
    r0_mask = tl.full([R0_BLOCK], True, tl.int1)[None, None, None, None, :]
    roffset = r0_offset
    rindex = r0_index
    r0_4 = r0_index
    y2 = yindex
    z1_0 = z1_index
    x3 = xindex
    z0_1 = z0_index
    tmp0 = tl.trans(tl.make_tensor_descriptor(in_ptr0, shape=[8, 8, 32], strides=[256, 32, 1], block_shape=[Z1_BLOCK, YBLOCK, R0_BLOCK]).load([z1_offset, yoffset, r0_offset])[:, :, :, None, None], [0, 3, 1, 4, 2])
    tmp1 = tl.trans(tl.make_tensor_descriptor(in_ptr1, shape=[8, 8, 32], strides=[256, 32, 1], block_shape=[Z0_BLOCK, XBLOCK, R0_BLOCK]).load([z0_offset, xoffset, r0_offset])[:, :, :, None, None], [3, 0, 4, 1, 2])
    tmp2 = tmp0 * tmp1
    tmp3 = tl.broadcast_to(tmp2, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK, R0_BLOCK])
    tmp5 = tl.where(xmask & ymask & z0_mask & z1_mask, tmp3, 0)
    tmp6 = tl.sum(tmp5, 4)[:, :, :, :, None].to(tl.float32)
    tl.make_tensor_descriptor(out_ptr0, shape=[8, 8, 8, 8], strides=[512, 64, 8, 1], block_shape=[Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).store([z1_offset, z0_offset, yoffset, xoffset], tl.reshape(tl.broadcast_to(tmp6, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK, 1]), [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).to(tl.float32))
```

## Counterexample: the same permute with the default two tiles

Set `triton.max_tiles=2`, its default, while keeping descriptors enabled and using the same padded input. This kernel has only X/Y tiles and uses `tl.load`/`tl.store`, with no descriptor. The collapsed strides in this example cannot form a valid descriptor; two tiles do not universally disable descriptors.

```python
import torch

def fn(x):
    return x.permute(1, 3, 0, 2).contiguous()

x = torch.randn(3, 5, 7, 32, device="cuda")[..., :16]
compiled = torch.compile(fn, options={
    "triton.max_tiles": 2,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(x)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_poi_fused_clone_permute_0(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    ynumel = 240
    xnumel = 7
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[:, None]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[None, :]
    xmask = xindex < xnumel
    x3 = xindex
    y0 = (yindex % 3)
    y1 = ((yindex // 3) % 16)
    y2 = yindex // 48
    y4 = yindex
    tmp0 = tl.load(in_ptr0 + (y1 + 32*x3 + 224*y2 + 1120*y0), xmask & ymask, eviction_policy='evict_last')
    tl.store(out_ptr0 + (x3 + 7*y4), tmp0, xmask & ymask)
```

## 4 independent reduction axes

Sum a padded tensor with 4 independent strides. The kernel retains `R0_BLOCK` through `R3_BLOCK` and loads the input through a tensor descriptor.

```python
import torch

def fn(x):
    return x.sum()

x = torch.randn(6, 6, 6, 16, device="cuda")[:3, :3, :3, :8]
compiled = torch.compile(fn, options={
    "triton.max_tiles": 4,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(x)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_per_fused_sum_0(in_ptr0, out_ptr0, xnumel, r0_numel, r1_numel, r2_numel, r3_numel, XBLOCK : tl.constexpr):
    xnumel = 1
    r0_numel = 3
    R0_BLOCK: tl.constexpr = 4
    r1_numel = 3
    R1_BLOCK: tl.constexpr = 4
    r2_numel = 3
    R2_BLOCK: tl.constexpr = 4
    r3_numel = 8
    R3_BLOCK: tl.constexpr = 8
    rnumel = r0_numel * r1_numel * r2_numel * r3_numel
    RBLOCK: tl.constexpr = R0_BLOCK*R1_BLOCK*R2_BLOCK*R3_BLOCK
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None, None, None, None]
    xmask = tl.full([XBLOCK], True, tl.int1)[:, None, None, None, None]
    r0_index = tl.arange(0, R0_BLOCK)[None, :, None, None, None]
    r0_offset = 0
    r0_mask = r0_index < r0_numel
    r1_index = tl.arange(0, R1_BLOCK)[None, None, :, None, None]
    r1_offset = 0
    r1_mask = r1_index < r1_numel
    r2_index = tl.arange(0, R2_BLOCK)[None, None, None, :, None]
    r2_offset = 0
    r2_mask = r2_index < r2_numel
    r3_index = tl.arange(0, R3_BLOCK)[None, None, None, None, :]
    r3_offset = 0
    r3_mask = r3_index < r3_numel
    roffset = r3_offset + r2_offset*r3_numel + r1_offset*r2_numel*r3_numel + r0_offset*r1_numel*r2_numel*r3_numel
    rindex = r3_index + r2_index*r3_numel + r1_index*r2_numel*r3_numel + r0_index*r1_numel*r2_numel*r3_numel
    r0_0 = r0_index
    r1_1 = r1_index
    r2_2 = r2_index
    r3_3 = r3_index
    tmp0 = tl.trans(tl.make_tensor_descriptor(in_ptr0, shape=[3, 3, 3, 8], strides=[576, 96, 16, 1], block_shape=[R0_BLOCK, R1_BLOCK, R2_BLOCK, R3_BLOCK]).load([r0_offset, r1_offset, r2_offset, r3_offset])[:, :, :, :, None], [4, 0, 1, 2, 3])
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, R0_BLOCK, R1_BLOCK, R2_BLOCK, R3_BLOCK])
    tmp3 = tl.where(r0_mask & r1_mask & r2_mask & r3_mask, tmp1, 0)
    tmp4 = tl.reshape(tmp3, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp4, 1)[:, None, None, None, None].to(tl.float32)
    tl.store(out_ptr0 + (tl.full([1, 1, 1, 1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1, 1, 1, 1)), tmp5, None)
```

## 5 independent reduction axes

Sum a padded tensor with 5 independent strides. The kernel retains `R0_BLOCK` through `R4_BLOCK` and loads the input through a tensor descriptor.

```python
import torch

def fn(x):
    return x.sum()

x = torch.randn(6, 6, 6, 6, 16, device="cuda")[:3, :3, :3, :3, :8]
compiled = torch.compile(fn, options={
    "triton.max_tiles": 5,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(x)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_per_fused_sum_0(in_ptr0, out_ptr0, xnumel, r0_numel, r1_numel, r2_numel, r3_numel, r4_numel, XBLOCK : tl.constexpr):
    xnumel = 1
    r0_numel = 3
    R0_BLOCK: tl.constexpr = 4
    r1_numel = 3
    R1_BLOCK: tl.constexpr = 4
    r2_numel = 3
    R2_BLOCK: tl.constexpr = 4
    r3_numel = 3
    R3_BLOCK: tl.constexpr = 4
    r4_numel = 8
    R4_BLOCK: tl.constexpr = 8
    rnumel = r0_numel * r1_numel * r2_numel * r3_numel * r4_numel
    RBLOCK: tl.constexpr = R0_BLOCK*R1_BLOCK*R2_BLOCK*R3_BLOCK*R4_BLOCK
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None, None, None, None, None]
    xmask = tl.full([XBLOCK], True, tl.int1)[:, None, None, None, None, None]
    r0_index = tl.arange(0, R0_BLOCK)[None, :, None, None, None, None]
    r0_offset = 0
    r0_mask = r0_index < r0_numel
    r1_index = tl.arange(0, R1_BLOCK)[None, None, :, None, None, None]
    r1_offset = 0
    r1_mask = r1_index < r1_numel
    r2_index = tl.arange(0, R2_BLOCK)[None, None, None, :, None, None]
    r2_offset = 0
    r2_mask = r2_index < r2_numel
    r3_index = tl.arange(0, R3_BLOCK)[None, None, None, None, :, None]
    r3_offset = 0
    r3_mask = r3_index < r3_numel
    r4_index = tl.arange(0, R4_BLOCK)[None, None, None, None, None, :]
    r4_offset = 0
    r4_mask = r4_index < r4_numel
    roffset = r4_offset + r3_offset*r4_numel + r2_offset*r3_numel*r4_numel + r1_offset*r2_numel*r3_numel*r4_numel + r0_offset*r1_numel*r2_numel*r3_numel*r4_numel
    rindex = r4_index + r3_index*r4_numel + r2_index*r3_numel*r4_numel + r1_index*r2_numel*r3_numel*r4_numel + r0_index*r1_numel*r2_numel*r3_numel*r4_numel
    r0_0 = r0_index
    r1_1 = r1_index
    r2_2 = r2_index
    r3_3 = r3_index
    r4_4 = r4_index
    tmp0 = tl.trans(tl.make_tensor_descriptor(in_ptr0, shape=[3, 3, 3, 3, 8], strides=[3456, 576, 96, 16, 1], block_shape=[R0_BLOCK, R1_BLOCK, R2_BLOCK, R3_BLOCK, R4_BLOCK]).load([r0_offset, r1_offset, r2_offset, r3_offset, r4_offset])[:, :, :, :, :, None], [5, 0, 1, 2, 3, 4])
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, R0_BLOCK, R1_BLOCK, R2_BLOCK, R3_BLOCK, R4_BLOCK])
    tmp3 = tl.where(r0_mask & r1_mask & r2_mask & r3_mask & r4_mask, tmp1, 0)
    tmp4 = tl.reshape(tmp3, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp4, 1)[:, None, None, None, None, None].to(tl.float32)
    tl.store(out_ptr0 + (tl.full([1, 1, 1, 1, 1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1, 1, 1, 1, 1)), tmp5, None)
```

## Joint 4D pointwise + 4D reduction with permute and broadcast

Permute and multiply broadcast inputs, then sum four independent reduction axes in one kernel. `Z1_BLOCK` and `R3_BLOCK` verify both families. The four-dimensional output uses a descriptor. The eight-dimensional input tiles exceed the descriptor rank limit of five, so their loads correctly fall back to `tl.load`.

```python
import torch

def fn(a, b):
    return (a.permute(1, 3, 0, 2, 7, 6, 5, 4) * b).sum(dim=(4, 5, 6, 7))

a = torch.randn(4, 4, 4, 8, 3, 7, 5, 8, device="cuda")
b = torch.randn(4, 8, 1, 1, 8, 5, 7, 3, device="cuda")
compiled = torch.compile(fn, options={
    "triton.max_tiles": 4,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(a, b)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_per_fused_mul_permute_sum_0(in_ptr0, in_ptr1, out_ptr0, z1_numel, z0_numel, ynumel, xnumel, r0_numel, r1_numel, r2_numel, r3_numel, Z1_BLOCK : tl.constexpr, Z0_BLOCK : tl.constexpr, YBLOCK : tl.constexpr, XBLOCK : tl.constexpr):
    z1_numel = 4
    z0_numel = 8
    ynumel = 4
    xnumel = 4
    r0_numel = 8
    R0_BLOCK: tl.constexpr = 8
    r1_numel = 5
    R1_BLOCK: tl.constexpr = 8
    r2_numel = 7
    R2_BLOCK: tl.constexpr = 8
    r3_numel = 3
    R3_BLOCK: tl.constexpr = 4
    rnumel = r0_numel * r1_numel * r2_numel * r3_numel
    RBLOCK: tl.constexpr = R0_BLOCK*R1_BLOCK*R2_BLOCK*R3_BLOCK
    z1_offset = (tl.program_id(2) // tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z1_BLOCK
    z1_index = z1_offset + tl.arange(0, Z1_BLOCK)[:, None, None, None, None, None, None, None]
    z1_mask = z1_index < z1_numel
    z0_offset = (tl.program_id(2) % tl.cdiv(z0_numel, Z0_BLOCK)).to(tl.int32) * Z0_BLOCK
    z0_index = z0_offset + tl.arange(0, Z0_BLOCK)[None, :, None, None, None, None, None, None]
    z0_mask = z0_index < z0_numel
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, None, :, None, None, None, None, None]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[None, None, None, :, None, None, None, None]
    xmask = xindex < xnumel
    r0_index = tl.arange(0, R0_BLOCK)[None, None, None, None, :, None, None, None]
    r0_offset = 0
    r0_mask = r0_index < r0_numel
    r1_index = tl.arange(0, R1_BLOCK)[None, None, None, None, None, :, None, None]
    r1_offset = 0
    r1_mask = r1_index < r1_numel
    r2_index = tl.arange(0, R2_BLOCK)[None, None, None, None, None, None, :, None]
    r2_offset = 0
    r2_mask = r2_index < r2_numel
    r3_index = tl.arange(0, R3_BLOCK)[None, None, None, None, None, None, None, :]
    r3_offset = 0
    r3_mask = r3_index < r3_numel
    roffset = r3_offset + r2_offset*r3_numel + r1_offset*r2_numel*r3_numel + r0_offset*r1_numel*r2_numel*r3_numel
    rindex = r3_index + r2_index*r3_numel + r1_index*r2_numel*r3_numel + r0_index*r1_numel*r2_numel*r3_numel
    r0_4 = r0_index
    r1_5 = r1_index
    r2_6 = r2_index
    r3_7 = r3_index
    x3 = xindex
    y2 = yindex
    z0_1 = z0_index
    z1_0 = z1_index
    tmp0 = tl.load(in_ptr0 + (r0_4 + 8*r1_5 + 40*r2_6 + 280*r3_7 + 840*z0_1 + 6720*x3 + 26880*z1_0 + 107520*y2), r0_mask & r1_mask & r2_mask & r3_mask & xmask & ymask & z0_mask & z1_mask, eviction_policy='evict_first', other=0.0)
    tmp1 = tl.load(in_ptr1 + (r3_7 + 3*r2_6 + 21*r1_5 + 105*r0_4 + 840*z0_1 + 6720*z1_0), r0_mask & r1_mask & r2_mask & r3_mask & z0_mask & z1_mask, eviction_policy='evict_last', other=0.0)
    tmp2 = tmp0 * tmp1
    tmp3 = tl.broadcast_to(tmp2, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK, R0_BLOCK, R1_BLOCK, R2_BLOCK, R3_BLOCK])
    tmp5 = tl.where(r0_mask & r1_mask & r2_mask & r3_mask & xmask & ymask & z0_mask & z1_mask, tmp3, 0)
    tmp6 = tl.reshape(tmp5, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK, RBLOCK])
    tmp7 = tl.sum(tmp6, 4)[:, :, :, :, None, None, None, None].to(tl.float32)
    tl.make_tensor_descriptor(out_ptr0, shape=[4, 8, 4, 4], strides=[128, 16, 4, 1], block_shape=[Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).store([z1_offset, z0_offset, yoffset, xoffset], tl.reshape(tl.broadcast_to(tmp7, [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK, 1, 1, 1, 1]), [Z1_BLOCK, Z0_BLOCK, YBLOCK, XBLOCK]).to(tl.float32))
```

## Counterexample: padded reduction with the default two tiles

Use the same padded 4D input with `triton.max_tiles=2`. Only `R0_BLOCK` and `R1_BLOCK` are emitted, and the collapsed strides require ordinary loads rather than a descriptor.

```python
import torch

def fn(x):
    return x.sum()

x = torch.randn(6, 6, 6, 16, device="cuda")[:3, :3, :3, :8]
compiled = torch.compile(fn, options={
    "triton.max_tiles": 2,
    "triton.prefer_nd_tiling": True,
    "triton.use_tensor_descriptor": True,
    "assume_aligned_inputs": True,
    "triton.tile_reductions": True,
})
result = compiled(x)
```

```python
import triton
import triton.language as tl

triton.jit
def triton_per_fused_sum_0(in_ptr0, out_ptr0, xnumel, r0_numel, r1_numel, XBLOCK : tl.constexpr):
    xnumel = 1
    r0_numel = 27
    R0_BLOCK: tl.constexpr = 32
    r1_numel = 8
    R1_BLOCK: tl.constexpr = 8
    rnumel = r0_numel * r1_numel
    RBLOCK: tl.constexpr = R0_BLOCK*R1_BLOCK
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None, None]
    xmask = tl.full([XBLOCK], True, tl.int1)[:, None, None]
    r0_index = tl.arange(0, R0_BLOCK)[None, :, None]
    r0_offset = 0
    r0_mask = r0_index < r0_numel
    r1_index = tl.arange(0, R1_BLOCK)[None, None, :]
    r1_offset = 0
    r1_mask = r1_index < r1_numel
    roffset = r1_offset + r0_offset*r1_numel
    rindex = r1_index + r0_index*r1_numel
    r0_0 = (r0_index % 3)
    r0_1 = ((r0_index // 3) % 3)
    r0_2 = r0_index // 9
    r1_3 = r1_index
    tmp0 = tl.load(in_ptr0 + (r1_3 + 16*r0_0 + 96*r0_1 + 576*r0_2), r0_mask & r1_mask, eviction_policy='evict_first', other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, R0_BLOCK, R1_BLOCK])
    tmp3 = tl.where(r0_mask & r1_mask, tmp1, 0)
    tmp4 = tl.reshape(tmp3, [XBLOCK, RBLOCK])
    tmp5 = tl.sum(tmp4, 1)[:, None, None].to(tl.float32)
    tl.store(out_ptr0 + (tl.full([1, 1, 1], 0, tl.int32).broadcast_to(XBLOCK, 1, 1)), tmp5, None)
```

# Testing

- Add launch-grid unit tests for 3D/4D/5D kernels, checking that Z packing reconstructs each logical tile once and preserves X/Y program IDs with partial boundary tiles.
- Add block size heuristic unit tests for ND expansion, unit dimensions, block limits, malformed inputs, and unchanged existing 1D/2D/3D configurations and warp counts.
- Add regression coverage for existing reduction configuration expansion and pointwise/reduction block sizes.
- Extend prefix-helper coverage for pointwise block products and axis exclusions.
- Add numerical tests for 4D broadcasts, 5D permute/reshape refinement, and dynamic shapes including empty dimensions. Descriptor variants assert the numbered Z axes and descriptors in generated kernels.
- Execute the four pointwise examples on H100, check numerical correctness, and verify the descriptor and ordinary-load paths.
- Extend configuration expansion coverage to four and five reduction axes, checking the emitted block sizes without changing existing tuning constants.
- Add `ndims`-parametrized padded reduction tests with four and five independent reduction axes, persistent and looped reductions, and block-pointer and tensor-descriptor variants.
- Include broadcast operands that vary only along the last reduction axis, covering reduction-owned descriptor loads inside the loop for `R3_BLOCK`/`R4_BLOCK`.
- Add joint four-pointwise/two-reduction and four-pointwise/four-reduction permute/broadcast/multiply/sum tests, checking numerical correctness, one fused kernel, and both axis families.
- Execute the independent 4D/5D and joint 4D+4D examples on H100, checking descriptors and the rank-limit fallback for eight-dimensional input tiles.

# Stack

- **1/3** [D130812310: Rename `ZBLOCK` to `Z0_BLOCK`](https://www.internalfb.com/diff/D130812310) — [PyTorch PR #204273](#204273)
- **2/3** [D130812312: Generalize tiled axis handling](https://www.internalfb.com/diff/D130812312) — [PyTorch PR #204281](#204281)
- **3/3** [D130812313: Support ND pointwise and reduction tiling](https://www.internalfb.com/diff/D130812313) — [PyTorch PR #204324](#204324)

# Agent details

Revised with Codex. Example kernel bodies were captured from executed Inductor compilations.

Test Plan:
```lang=bash
buck2 test mode/opt --agent-context direct_call=true --agent-context 'reason=Validate unified GridND for 3D and ND launches, unchanged autotuning budgets, and Hopper pointwise tensor descriptors' fbcode//caffe2/test/inductor:simd_range_trees fbcode//caffe2/test/inductor:triton_heuristics fbcode//caffe2/test/inductor:torchinductor_strided_blocks -- --regex 'TestSIMDRangeTrees|pointwise_config|pointwise_budget|existing_pointwise|reduction_config_budget|test_nd_pointwise|test_pointwise_|test_broadcast_'
```
https://www.internalfb.com/intern/testinfra/testrun/20547673367950394

123 passed, 0 failed; 32 skipped and one already-disabled test omitted. Covers 3D/4D/5D launch coordinates, existing and expanded block budgets, pointwise broadcasts, dynamic shapes, and 5D Hopper tensor descriptors.

Pointwise-only revision validation: https://www.internalfb.com/intern/testinfra/testrun/20829148344842151
Seven ND pointwise cases, twelve existing configuration cases, four descriptor examples and the three-axis reduction guard passed.

ND stack split validation (2026-11-10):
H100 validation: https://www.internalfb.com/intern/testinfra/testrun/35465847120931160

- 24 independent padded reduction variants passed: four/five axes, persistent/looped, with/without last-axis broadcast, CUDA block pointers and host/device tensor descriptors.
- 11 ND pointwise/joint cases passed, including four pointwise plus four reduction axes.
- 14 existing and extended configuration cases passed.
- Eight numerical examples and generated kernels captured with `TORCH_COMPILE_DEBUG=1` passed, including descriptor generation and the `max_tiles=2` counterexamples.
- Eight Triton CPU descriptor cases skipped because that backend is unavailable on this runner.

`arc lint -a`: passed.
Conda `lintrunner -a --paths-from /tmp/nd_lintrunner_oss_paths.txt --skip CODEOWNERS_TAXONOMY,PYREFLY --output=json`: passed. CODEOWNERS_TAXONOMY needs a Git checkout and PYREFLY needs the matching built OSS environment; these previously agreed environment exclusions remain in place.

The ND reduction changes and tests are now folded into this revision. Source contents match the previously validated combined pointwise/reduction revision exactly; the successful H100 results above cover independent 4D/5D reductions and joint 4D pointwise + 4D reduction kernels.

Differential Revision: D130812313