pytorch - 💡(How to fix) Fix [inductor] realize_reads_threshold could be tuned differently with auto-chunking [1 comments, 2 participants]

pytorch2026-03-10 18:58:52

ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

GitHub issue URL

Helpful · Quick feedback

GitHub stats

pytorch/pytorch#177053•Fetched 2026-04-08 00:22:29

View on GitHub

Comments

Participants

Timeline

Reactions

Author

Participants

Assignees

Timeline (top)

mentioned ×3subscribed ×3labeled ×2assigned ×1

Error Message

Error logs

Code Example

From the actual Inductor output (chunking_subgraph_0), the materialized fp32 buffer:

  # Inside chunking_subgraph_0:
  chunking_subgraph_0_buf0 = empty_strided_cuda((4096, 200003), (200064, 1), torch.bfloat16)  # logits
  extern_kernels.mm(x_chunk, weight, out=chunking_subgraph_0_buf0)

  chunking_subgraph_0_buf3 = empty_strided_cuda((4096, 200003), (200032, 1), torch.float32)   # <--- THE PROBLEM: ~3.1 GiB fp32 buffer!
  chunking_subgraph_0_buf5 = empty_strided_cuda((4096, 200003), (200064, 1), torch.bfloat16)  # gradient output

  triton_fused_kernel_27.run(..., chunking_subgraph_0_buf3, chunking_subgraph_0_buf5, ...)
  del chunking_subgraph_0_buf3  # freed immediately after, but damage is done — peak memory already spiked

  Inside the fused triton kernel, the first loop writes to this fp32 buffer:

  # First loop: online softmax + CE backward sparse gradient
  for r0_offset in tl.range(0, r0_numel, R0_BLOCK):    # r0_numel = 200003
      # ... online softmax accumulation ...
      # ... compute sparse CE gradient (one-hot * loss_weight) ...
      tmp22 = tmp12 * tmp21                               # sparse gradient * loss weight
      tl.store(out_ptr2 + (r0_1 + 200032*x0), tmp22, r0_mask)  # <-- WRITES to 3.1 GiB fp32 buf3

  And the second loop reads it back:

  # Second loop: compute final dlogits using softmax + the stored intermediate
  for r0_offset in tl.range(0, r0_numel, R0_BLOCK):
      tmp26 = tl.load(out_ptr2 + (r0_1 + 200032*x0), ...)  # <-- READS from 3.1 GiB fp32 buf3
      # ... recompute bf16→fp32 logits, softmax ...
      tmp36 = tmp26 - tmp35                                  # sparse_grad - softmax * sum_grad
      tl.store(out_ptr4 + (r0_1 + 200064*x0), tmp36.to(tl.float32), r0_mask)  # bf16 gradient

RAW_BUFFERClick to expand / collapse

🐛 Describe the bug

I don't think this is a bug, but filing an issue so that we don't lose track of it.

As discovered in my post (Meta-access only), when using Inductor's auto-chunker, it creates a single subgraph with both forward and backward ops for the cross-entropy loss. This gives intermediate tensors more users than they'd have in separate forward/backward only graphs. This causes Inductor's realization heuristic (realize_reads_threshold > 4) to materialize a large fp32 buffer (for online softmax) that increases memory footprint significantly.

Manually bumping up the threshold to 6 avoids the materialization , minimizing footprint without degrading perf. There is the opportunity to tune this threshold differently when auto-chunking is enabled.

This heuristic's tuning could also become salient if we start capturing fwd + bwd as a single graph.

To reproduce -- consult the workplace post. TLParse with inductor output code

Pointer to fp32 buffer:

From the actual Inductor output (chunking_subgraph_0), the materialized fp32 buffer:

  # Inside chunking_subgraph_0:
  chunking_subgraph_0_buf0 = empty_strided_cuda((4096, 200003), (200064, 1), torch.bfloat16)  # logits
  extern_kernels.mm(x_chunk, weight, out=chunking_subgraph_0_buf0)

  chunking_subgraph_0_buf3 = empty_strided_cuda((4096, 200003), (200032, 1), torch.float32)   # <--- THE PROBLEM: ~3.1 GiB fp32 buffer!
  chunking_subgraph_0_buf5 = empty_strided_cuda((4096, 200003), (200064, 1), torch.bfloat16)  # gradient output

  triton_fused_kernel_27.run(..., chunking_subgraph_0_buf3, chunking_subgraph_0_buf5, ...)
  del chunking_subgraph_0_buf3  # freed immediately after, but damage is done — peak memory already spiked

  Inside the fused triton kernel, the first loop writes to this fp32 buffer:

  # First loop: online softmax + CE backward sparse gradient
  for r0_offset in tl.range(0, r0_numel, R0_BLOCK):    # r0_numel = 200003
      # ... online softmax accumulation ...
      # ... compute sparse CE gradient (one-hot * loss_weight) ...
      tmp22 = tmp12 * tmp21                               # sparse gradient * loss weight
      tl.store(out_ptr2 + (r0_1 + 200032*x0), tmp22, r0_mask)  # <-- WRITES to 3.1 GiB fp32 buf3

  And the second loop reads it back:

  # Second loop: compute final dlogits using softmax + the stored intermediate
  for r0_offset in tl.range(0, r0_numel, R0_BLOCK):
      tmp26 = tl.load(out_ptr2 + (r0_1 + 200032*x0), ...)  # <-- READS from 3.1 GiB fp32 buf3
      # ... recompute bf16→fp32 logits, softmax ...
      tmp36 = tmp26 - tmp35                                  # sparse_grad - softmax * sum_grad
      tl.store(out_ptr4 + (r0_1 + 200064*x0), tmp36.to(tl.float32), r0_mask)  # bf16 gradient

Error logs

No response

Versions

Collecting environment information... PyTorch version: N/A Is debug build: N/A CUDA used to build PyTorch: N/A ROCM used to build PyTorch: N/A

OS: CentOS Stream 9 (x86_64) GCC version: (GCC) 11.5.0 20240719 (Red Hat 11.5.0-14) Clang version: Could not collect CMake version: version 3.31.6 Libc version: glibc-2.34

Python version: 3.12.12 | packaged by Anaconda, Inc. | (main, Oct 21 2025, 20:16:04) [GCC 11.2.0] (64-bit runtime) Python platform: Linux-6.4.3-0_fbk15_hardened_2630_gf27365f948db-x86_64-with-glibc2.34 Is CUDA available: N/A CUDA runtime version: 12.9.86 CUDA_MODULE_LOADING set to: N/A GPU models and configuration: GPU 0: NVIDIA H100 GPU 1: NVIDIA H100 GPU 2: NVIDIA H100 GPU 3: NVIDIA H100

Nvidia driver version: 550.90.07 cuDNN version: Could not collect Is XPU available: N/A HIP runtime version: N/A MIOpen runtime version: N/A Is XNNPACK available: N/A Caching allocator config: N/A

CPU: Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Address sizes: 52 bits physical, 57 bits virtual Byte Order: Little Endian CPU(s): 184 On-line CPU(s) list: 0-183 Vendor ID: AuthenticAMD Model name: AMD EPYC 9654 96-Core Processor CPU family: 25 Model: 17 Thread(s) per core: 1 Core(s) per socket: 184 Socket(s): 1 Stepping: 1 BogoMIPS: 4792.80 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm rep_good nopl cpuid extd_apicid tsc_known_freq pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy svm cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw perfctr_core invpcid_single ssbd ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves avx512_bf16 clzero xsaveerptr wbnoinvd arat npt lbrv nrip_save tsc_scale vmcb_clean flushbyasid pausefilter pfthreshold v_vmsave_vmload vgif vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq rdpid fsrm flush_l1d arch_capabilities Virtualization: AMD-V Hypervisor vendor: KVM Virtualization type: full L1d cache: 11.5 MiB (184 instances) L1i cache: 11.5 MiB (184 instances) L2 cache: 92 MiB (184 instances) L3 cache: 2.9 GiB (184 instances) NUMA node(s): 1 NUMA node0 CPU(s): 0-183 Vulnerability Gather data sampling: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Retbleed: Not affected Vulnerability Spec store bypass: Vulnerable Vulnerability Spectre v1: Vulnerable: __user pointer sanitization and usercopy barriers only; no swapgs barriers Vulnerability Spectre v2: Vulnerable, IBPB: disabled, STIBP: disabled, PBRSB-eIBRS: Not affected Vulnerability Srbds: Not affected Vulnerability Tsx async abort: Not affected

Versions of relevant libraries: [pip3] flake8==7.3.0 [pip3] flake8-bugbear==24.12.12 [pip3] flake8-comprehensions==3.16.0 [pip3] flake8-executable==2.1.3 [pip3] flake8-logging-format==2024.24.12 [pip3] flake8-pyi==25.5.0 [pip3] flake8_simplify==0.22.0 [pip3] intel-cmplr-lib-ur==2025.3.1 [pip3] intel-openmp==2025.3.1 [pip3] mkl-include==2025.3.0 [pip3] mkl-static==2025.3.0 [pip3] mypy==1.16.0 [pip3] mypy_extensions==1.1.0 [pip3] numpy==2.4.2 [pip3] nvidia-cublas-cu12==12.9.1.4 [pip3] nvidia-cuda-cupti-cu12==12.9.79 [pip3] nvidia-cuda-nvrtc-cu12==12.9.86 [pip3] nvidia-cuda-runtime-cu12==12.9.79 [pip3] nvidia-cudnn-cu12==9.17.1.4 [pip3] nvidia-cufft-cu12==11.4.1.4 [pip3] nvidia-curand-cu12==10.3.10.19 [pip3] nvidia-cusolver-cu12==11.7.5.82 [pip3] nvidia-cusparse-cu12==12.5.10.65 [pip3] nvidia-cusparselt-cu12==0.7.1 [pip3] nvidia-nccl-cu12==2.28.9 [pip3] nvidia-nvjitlink-cu12==12.9.86 [pip3] nvidia-nvtx-cu12==12.9.79 [pip3] onemkl-license==2025.3.0 [pip3] onnx==1.19.1 [pip3] onnx-ir==0.1.12 [pip3] onnxscript==0.5.4 [pip3] optree==0.17.0 [pip3] tbb==2022.3.0 [pip3] tbb-devel==2022.3.0 [pip3] tcmlib==1.4.1 [pip3] torch==2.12.0a0+git536ab83 [pip3] torchdata==0.11.0 [pip3] torchtitan==0.2.1 [pip3] torchvision==0.26.0.dev20260225+cu129 [pip3] triton==3.6.0+git9844da95 [pip3] umf==1.0.2 [conda] intel-cmplr-lib-ur 2025.3.1 pypi_0 pypi [conda] intel-openmp 2025.3.1 pypi_0 pypi [conda] magma-cuda126 2.6.1 1 pytorch [conda] mkl-include 2025.3.0 pypi_0 pypi [conda] mkl-static 2025.3.0 pypi_0 pypi [conda] numpy 2.4.2 pypi_0 pypi [conda] nvidia-cublas-cu12 12.9.1.4 pypi_0 pypi [conda] nvidia-cuda-cupti-cu12 12.9.79 pypi_0 pypi [conda] nvidia-cuda-nvrtc-cu12 12.9.86 pypi_0 pypi [conda] nvidia-cuda-runtime-cu12 12.9.79 pypi_0 pypi [conda] nvidia-cudnn-cu12 9.17.1.4 pypi_0 pypi [conda] nvidia-cufft-cu12 11.4.1.4 pypi_0 pypi [conda] nvidia-curand-cu12 10.3.10.19 pypi_0 pypi [conda] nvidia-cusolver-cu12 11.7.5.82 pypi_0 pypi [conda] nvidia-cusparse-cu12 12.5.10.65 pypi_0 pypi [conda] nvidia-cusparselt-cu12 0.7.1 pypi_0 pypi [conda] nvidia-nccl-cu12 2.28.9 pypi_0 pypi [conda] nvidia-nvjitlink-cu12 12.9.86 pypi_0 pypi [conda] nvidia-nvtx-cu12 12.9.79 pypi_0 pypi [conda] onemkl-license 2025.3.0 pypi_0 pypi [conda] optree 0.17.0 pypi_0 pypi [conda] tbb 2022.3.0 pypi_0 pypi [conda] tbb-devel 2022.3.0 pypi_0 pypi [conda] tcmlib 1.4.1 pypi_0 pypi [conda] torch 2.12.0a0+git536ab83 pypi_0 pypi [conda] torchdata 0.11.0 pypi_0 pypi [conda] torchfix 0.4.0 pypi_0 pypi [conda] torchtitan 0.2.1 pypi_0 pypi [conda] torchvision 0.26.0.dev20260225+cu129 pypi_0 pypi [conda] triton 3.6.0+git9844da95 pypi_0 pypi [conda] umf 1.0.2 pypi_0 pypi

cc @chauhang @penguinwu

extent analysis

Fix Plan

Problem: Large fp32 buffer materialization due to Inductor's auto-chunker

Solution:

Tune the realization heuristic threshold: Manually bump up the realize_reads_threshold to 6 to avoid materializing the large fp32 buffer.

In your PyTorch model or training script

import torch

Set the realization heuristic threshold to 6

torch.backends.triton.realize_reads_threshold = 6


2. **Alternative solution: Optimize the model for auto-chunking**: If possible, optimize the model to reduce the number of intermediate tensors and minimize the memory footprint.

   ```python
# In your PyTorch model or training script
import torch

# Use the `torch.jit.optimize_for_inference` function to optimize the model
model = torch.jit.optimize_for_inference(model)

Temporary workaround: Disable auto-chunking: If the above solutions are not feasible, you can temporarily disable auto-chunking by setting the triton_auto_chunk flag to False.

In your PyTorch model or training script

import torch

Set the `triton_auto_chunk` flag to `False`

torch.backends.triton.auto_chunk = False


## Verification

1. **Monitor memory usage**: Use tools like `nvidia-smi` or `top` to monitor the memory usage of your GPU.
2. **Check the Inductor output**: Verify that the large fp32 buffer is not materialized by checking the Inductor output.
3. **Run the model with the fix**: Run the model with the fix and verify that it produces the expected results.

## Extra Tips

* **Profile the model**: Use tools like `pytorch-profiler` to profile the model and identify performance bottlenecks.
* **Optimize the model**: Optimize the model to reduce the number of

Vote matrix · Quick signals

Works

Did the solution work? Tap to confirm.

Easy Fix

Was it a quick fix?

Time Saver

Did it save you time?

Blocking

Was it severely blocking?

Common Issue

Are others likely hitting this too?

Flaky / Intermittent

Is it intermittent?

Verified / Reproducible

Can you reproduce it reliably?

#api #ssr #installation #tensor shape #autograd error #environment setup #docker error #permission error #memory optimization

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Data

Security

Network

Code

UI/UX

Text

System

Multimedia

Protocol

API

Engineering

pytorch - 💡(How to fix) Fix [inductor] realize_reads_threshold could be tuned differently with auto-chunking [1 comments, 2 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Error logs

Code Example

🐛 Describe the bug

Error logs

Versions

extent analysis

Fix Plan

Problem: Large fp32 buffer materialization due to Inductor's auto-chunker

Solution:

In your PyTorch model or training script

Set the realization heuristic threshold to 6

In your PyTorch model or training script

Set the `triton_auto_chunk` flag to `False`

Still need to ship something?

TRENDING

pytorch - 💡(How to fix) Fix [inductor] realize_reads_threshold could be tuned differently with auto-chunking [1 comments, 2 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Error logs

Code Example

🐛 Describe the bug

Error logs

Versions

extent analysis

Fix Plan

Problem: Large fp32 buffer materialization due to Inductor's auto-chunker

Solution:

In your PyTorch model or training script

Set the realization heuristic threshold to 6

In your PyTorch model or training script

Set the triton_auto_chunk flag to False

Still need to ship something?

RELATED_DISCOVERY

TRENDING

Set the `triton_auto_chunk` flag to `False`