[Bug] High Local Disk IOPS and Frequent File Open Operations with FLA 0.5.1 During Multimodal Model Training

[leooop-al]

Checklist

• I have checked FAQs and existing issues for similar problems
• Please report this bug in English to ensure wider understanding and support

Describe the Bug

During multimodal model training with FLA 0.5.1, we observed extremely frequent local disk file operations, especially open/read/stat calls. This behavior appears to be strongly correlated with degraded GPU utilization and increasing host-side gaps before CUDA kernel launches.

The issue does not seem to come from normal multimodal data loading, such as image reads. When reverting to FLA 0.4.2, the disk I/O pattern becomes normal again.

Observed Behavior

With FLA 0.5.1, the training job shows approximately:

• Read throughput: around 48 MB/s
• Read IOPS: around 5000/s
• Average read size: around 10 KB

This suggests a large number of small file reads, likely related to cache or metadata access rather than normal image/data loading.

After switching to FLA 0.4.2 under the same multimodal training setup, the disk I/O pattern becomes:

• Read throughput: around 22 MB/s
• Read IOPS: around 100/s
• Open operations: around 50/s
• Average read size: around 220 KB

This is consistent with normal multimodal data loading behavior, especially image reads.

Therefore, the excessive IOPS in FLA 0.5.1 appears to come from FLA/Triton cache-related small-file access rather than the multimodal dataset itself.

Possible Root Cause

The most suspicious change is the persistent autotune cache introduced in:

• [Cache] FLA autotune cache #798

In our setup, we run multiple ranks on the same node. When all ranks share the same node-local cache directory, FLA/Triton may repeatedly perform cache lookup and validation, causing a large number of small file operations:

Python launch thread
    -> Triton / FLA autotune cache lookup
    -> frequent open / read / stat calls
    -> host-side gap before CUDA kernel launch
    -> lower SMA / GPU utilization

This matches the profiler observation where chunk_gated_delta_rule_fwd_h has an increasing host-side gap before actual CUDA kernel execution.

Impact

The issue may cause:

• Very high Local Disk IOPS
• Frequent file open operations
• Increased host-side latency before CUDA kernel launches
• Degraded GPU utilization / SMA
• Step time instability or regression
• Possible training slowdown over time

In our trace:

• The workload is multimodal model training with image inputs.
• The number of GDN calls per step remains stable, around 720 calls/step.
• Tokens/chunks are mostly stable.
• However, chunk_gated_delta_rule_fwd_h average time increases from around 7 ms to 14–17 ms.
• The extra time mainly appears before GPU kernels are launched, suggesting a CPU-side or cache lookup bottleneck.
• The disk I/O pattern under FLA 0.5.1 is inconsistent with normal image loading, because the average read size is only around 10 KB.

Related FLA Changes

The following FLA changes may be related:

• [Cache] FLA autotune cache #798
  Introduced FLA autotune cache.
• [Enhance] Extend tensor_cache to a bounded queue #895
  Improved tensor_cache to support bounded multiple entries.
• Make varlen index helpers torch.compile-friendly #909
  Made varlen index helpers more compile-friendly.

Among them, PR #798 is the most relevant to the abnormal disk I/O pattern.

Environment

Please let me know if more details are needed. The key setup is:

FLA version: 0.5.1
Training type: multimodal model training
Data type: text + image inputs
Training style: Qwen3.5-style packed / varlen training
Distributed setup: 8 ranks per node
Cache directory: shared node-local cache directory
Kernel of interest: chunk_gated_delta_rule_fwd_h

Reproduction / Validation

A possible way to validate this issue is:

1. Run the same multimodal training job with FLA 0.5.1.
2. Monitor local disk metrics:
   • Read throughput
   • Read IOPS
   • Open operations
   • open/read/stat syscall frequency
3. Compare against FLA 0.4.2 with the same training parameters and data order.
4. Check whether chunk_gated_delta_rule_fwd_h shows increasing host-side gaps before CUDA kernel launch.

In our case, the difference is very clear:

FLA 0.5.1:
  Read:      ~48 MB/s
  Read IOPS: ~5000/s
  Avg read:  ~10 KB

FLA 0.4.2:
  Read:      ~22 MB/s
  Read IOPS: ~100/s
  Open:      ~50/s
  Avg read:  ~220 KB

The FLA 0.4.2 pattern is consistent with normal image loading in multimodal training, while the FLA 0.5.1 pattern indicates many tiny cache or metadata reads.

Temporary Workaround to Test

We plan to test FLA 0.5.1 with disk autotune result caching disabled:

export FLA_CACHE_RESULTS=0
export FLA_CACHE_MODE=disabled

These variables need to be set before launching training.

We do not plan to set:

export FLA_DISABLE_TENSOR_CACHE=1

because the tensor cache is an in-memory cache, and disabling it may increase CPU overhead further.

Expected Behavior

With FLA 0.5.1, cache lookup should not cause thousands of small disk reads or extremely frequent file open operations during steady-state multimodal training.

Ideally:

• Local Disk IOPS should remain close to the normal data-loading level.
• Image loading should remain the dominant and expected source of disk reads.
• Cache metadata access should not block or delay CUDA kernel launches.
• chunk_gated_delta_rule_fwd_h host-side gap should remain stable.
• GPU utilization / SMA should not degrade over time due to cache file operations.

Questions

1. Is the persistent autotune cache in FLA 0.5.1 expected to perform frequent disk file checks during training?
2. Are FLA_CACHE_RESULTS=0 and FLA_CACHE_MODE=disabled the recommended way to disable disk-backed autotune cache?
3. Is there a recommended per-rank cache directory setup to avoid contention when multiple ranks share the same node-local cache?
4. Could FLA avoid repeated disk checks after autotune results are loaded into memory?
5. Are there additional debug flags to trace FLA autotune cache hit/miss behavior?
6. For multimodal training, are there any recommended FLA/Triton cache settings to avoid interfering with normal image data loading?

Additional Context

This issue is important because the regression is not simply higher disk bandwidth. The main difference is the number of small file operations:

0.5.1: many tiny reads, very high IOPS
0.4.2: fewer larger reads, consistent with normal image loading

This strongly suggests cache or metadata lookup overhead in FLA 0.5.1. If disabling disk autotune cache brings IOPS back to around 100/s and stabilizes SMA, then the best workaround may be:

Use FLA 0.5.1
+ keep newer varlen / tensor-cache fixes
+ disable disk-backed autotune cache

rather than reverting permanently to FLA 0.4.2.

Steps to Reproduce the Bug

#!/usr/bin/env bash

set -euo pipefail

export FLA_TILELANG=0
export SWIFT_USE_MCORE_GDN=1
export PYTORCH_CUDA_ALLOC_CONF='expandable_segments:True'

export IMAGE_MIN_TOKEN_NUM=64
export IMAGE_MAX_TOKEN_NUM=64
export VIDEO_MAX_TOKEN_NUM=128
export FPS_MAX_FRAMES=12

export CACHE_PATH=/tmp/qwen35_swift_cache
export HF_DATASETS_CACHE=${CACHE_PATH}/hf_datasets
export MODELSCOPE_CACHE=${CACHE_PATH}/modelscope
export XDG_CACHE_HOME=${CACHE_PATH}/xdg
mkdir -p "${CACHE_PATH}"
sudo chmod 777 "${CACHE_PATH}"

MODEL="Qwen3.5-35B-A3B"

DATASET="train_local.jsonl"

TP_SIZE=${TP_SIZE:-2}
PP_SIZE=${PP_SIZE:-2}
VPP_SIZE=${VPP_SIZE:-null}
CP_SIZE=${CP_SIZE:-1}
EP_SIZE=${EP_SIZE:-8}
ETP_SIZE=${ETP_SIZE:-1}

SHARD_CONFIG="tp${TP_SIZE}_pp${PP_SIZE}_ep${EP_SIZE}_cp${CP_SIZE}_2nnodes_anayls"

OUTPUT_DIR="${OUTPUT_DIR}"

sudo mkdir -p "${OUTPUT_DIR}"
sudo chmod 777 "${OUTPUT_DIR}"

PROFILE_DIR=${PROFILE_DIR:-"${OUTPUT_DIR}/profiler"}
ENABLE_PROFILE=${ENABLE_PROFILE:-True}
mkdir -p "${PROFILE_DIR}"

megatron sft
--model "${MODEL}"
--save_safetensors true
--dataset "${DATASET}"
--load_from_cache_file true
--split_dataset_ratio 0.05
--tuner_type full
--context_parallel_size ${CP_SIZE}
--tensor_model_parallel_size ${TP_SIZE}
--pipeline_model_parallel_size ${PP_SIZE}
--expert_model_parallel_size ${EP_SIZE}
--overlap_p2p_comm true
--overlap_grad_reduce true
--overlap_param_gather true
--sequence_parallel true
--moe_permute_fusion true
--moe_grouped_gemm true
--moe_shared_expert_overlap true
--moe_aux_loss_coeff 1e-6
--micro_batch_size 2
--global_batch_size 128
--recompute_granularity full
--recompute_method uniform
--recompute_num_layers 1
--num_train_epochs 5
--packing true
--finetune true
--freeze_llm false
--freeze_vit true
--freeze_aligner true
--cross_entropy_loss_fusion true
--lr 3e-6
--lr_warmup_fraction 0.05
--min_lr 1e-6
--max_length 14000
--dataloader_num_workers 2
--dataset_num_proc 32
--eval_steps 100
--save_steps 100
--no_save_optim true
--no_save_rng true
--moe_expert_capacity_factor 2
--optimizer_cpu_offload true
--use_precision_aware_optimizer true
--optimizer_offload_fraction 0.95
--attention_backend flash
--report_to wandb
--wandb_project "${WANDB_PROJECT}"
--wandb_exp_name "${WANDB_NAME}"
--logging_steps 10
--output_dir "${OUTPUT_DIR}"

Expected Behavior

1. Local Disk IOPS should remain close to the normal data-loading baseline.
   For multimodal training, disk reads are expected to mainly come from image/data loading. The I/O pattern should look like fewer, larger reads rather than thousands of tiny reads per second.
2. FLA/Triton cache lookup should not repeatedly hit the filesystem during steady-state training.
   After autotune results or compiled kernels are available, repeated kernel calls should avoid frequent open/read/stat operations on disk.
3. Multiple ranks on the same node should not cause cache contention or repeated metadata checks.
   When 8 ranks share the same node-local cache directory, cache lookup should remain efficient and should not generate thousands of small file operations per second.
4. CUDA kernel launches should not be delayed by cache file operations.
   The host-side gap before chunk_gated_delta_rule_fwd_h CUDA kernels should remain stable and should not increase over training steps due to cache lookup overhead.
5. GPU utilization / SMA should remain stable over time.
   There should not be a progressive drop in SMA caused by CPU-side stalls or disk-backed cache access.
6. Step time should remain stable after warmup.
   Once Triton compilation/autotuning is finished, step time should not regress due to ongoing cache metadata access.

Environment Information

Torch: CUDA 12.9 build
Triton: TBD

[leooop-al]

Update: forcing FLA_TILELANG=0 did not reduce the abnormal file-open/read rate, and the issue still reproduced. This makes the TileLang backend unlikely to be the root cause and narrows the problem to the Triton path, particularly autotune-result or JIT-cache access. The next controlled experiment is FLA_CACHE_RESULTS=0, followed by FLA_CACHE_MODE=default to bypass dynamic-key autotuning.

export FLA_TILELANG=0
export FLA_CACHE_MODE=disabled
export FLA_CACHE_RESULTS=0
either

[yurekami]

Two different caches are in play here, and I think they're being conflated — which matters because the first experiment on your list is a no-op.

FLA_CACHE_MODE (the PR #798 config cache) is already off by default, in 0.5.1 and on main. fla/ops/utils/cache.py:52:

        mode_str = os.environ.get("FLA_CACHE_MODE", cls.DISABLED.value)

DISABLED short-circuits before any file access (should_check_fla_cache, cache.py:368), and the only open() in the whole package — cache.py:257 — sits behind both that gate and a @cache memo. So setting FLA_CACHE_MODE=disabled changes nothing from your current run, and PR #798's cache is unlikely to be your source of small-file I/O.

FLA_CACHE_RESULTS is the one that defaults to on. fla/utils/_config.py:11:

FLA_CACHE_RESULTS = os.getenv('FLA_CACHE_RESULTS', '1') == '1'

and fla/utils/_compat.py:28-29:

SUPPORTS_AUTOTUNE_CACHE = "cache_results" in inspect.signature(triton.autotune).parameters
autotune_cache_kwargs = {"cache_results": FLA_CACHE_RESULTS} if SUPPORTS_AUTOTUNE_CACHE else {}

That dict is then spread as **autotune_cache_kwargs into the @triton.autotune decorators at 298 sites across fla/ (grep -rn autotune_cache_kwargs --include=*.py fla/ | wc -l). So by default every autotuned FLA kernel asks Triton to persist its autotune timings to disk — which is the small-file open/read/stat pattern you're measuring, and it is Triton's cache directory rather than FLA's.

This matches your regression window exactly. Neither file exists in v0.4.2:

$ gh api "repos/fla-org/flash-linear-attention/contents/fla/utils/_config.py?ref=v0.4.2"
404 Not Found
$ gh api "repos/fla-org/flash-linear-attention/contents/fla/utils/_compat.py?ref=v0.4.2"
404 Not Found

Both are present in v0.5.1 with the default above. So the cache_results plumbing is new in exactly the interval where your I/O profile changed.

It also explains the FLA_TILELANG=0 result. The TileLang probe cannot be your per-call I/O either — find_spec_cached and has_usable_nvcc are both @functools.cache-memoized (fla/utils/_compat.py:32-33 and :37-38) and _TILELANG_AVAILABLE is evaluated once at import (fla/ops/common/backends/tilelang/__init__.py:23), so that path does its filesystem work once per process regardless of the flag.

Suggested experiment order

Your list has the right variable in it, just not first. I'd reorder to:

export FLA_CACHE_RESULTS=0     # <- this one; the others are already the default

and leave FLA_CACHE_MODE unset. If the IOPS drop, it is Triton's autotune-results cache. ENVs.md:70 documents this variable as "Cache Triton autotune timings to disk (passed through to triton.autotune(cache_results=...) when supported)".

If it does not drop, the remaining candidate is Triton's ordinary JIT compile cache rather than anything FLA-specific, and TRITON_CACHE_DIR pointed at a node-local path (or tmpfs) would be the next thing to try — that would also address the multi-rank shared-directory contention you described, which cache_results=True will make worse since each rank writes timing files keyed by its own autotune keys.

Worth noting for the maintainers: SUPPORTS_AUTOTUNE_CACHE is gated on cache_results being present in triton.autotune's signature, so whether this is active at all depends on the Triton version — which would make the regression look inconsistent across environments.

Verification

• All quotes and line numbers read from main @ 0f0f0c9 this session, and the v0.5.1 / v0.4.2 files fetched at those tags to confirm the regression window.
• The 298 count is grep -rn autotune_cache_kwargs --include=*.py fla/ | wc -l on that checkout.
• Not reproduced. I do not have a multi-GPU training environment, so I have not measured IOPS before or after changing FLA_CACHE_RESULTS. The above is from reading the 0.4.2/0.5.1 sources; the experiment is yours to run and I may be wrong about which cache dominates.

[leooop-al]

However, the abnormal file-open/read rate and the progressive SMA degradation still reproduced. Therefore, Triton autotune timing persistence does not appear to be the dominant source in our environment.
The controlled version comparison remains:
FLA 0.5.1: thousands of small read/open operations per second and progressively decreasing SMA.
FLA 0.4.2: approximately normal image-loading I/O and stable SMA beyond 350 steps.
We will next:
Capture the exact paths involved in repeated openat/stat calls.
Test TRITON_CACHE_DIR on node-local tmpfs and, if possible, isolate it per rank.
A/B test the GDN forward implementation introduced around PR #789, especially the fused KKT + triangular-solve path.
This should distinguish ordinary Triton JIT-cache access from a regression in the newer GDN execution path.

[zhiyuan1i]

No repro on our side so far — waiting for more details, e.g. the exact training shapes (tokens per step, packed sequence count/lengths, head dims) and ideally the actual file paths showing up in the repeated openat/stat calls.

[leooop-al]

This MR only affects FLA config lookups when FLA_CACHE_MODE is enabled. With FlaCacheMode.DISABLED, CachedAutotuner.should_check_fla_cache() returns False, so load_cached_config() and the newly introduced load_kernel_config_file() are never called. Therefore, it cannot explain or fix the high filesystem IOPS reproduced with FLA_CACHE_MODE=disabled.
It also introduces persistent negative caching: if a config file is missing during the first lookup and is created later, non-ALWAYS modes will never discover it until the process restarts or the function cache is cleared. This is a behavioral change from the previous implementation and should be documented or reconsidered.
It would be useful to provide syscall-level evidence showing that the excessive open/stat operations actually originate from load_cached_config(), rather than Triton’s JIT/compiler cache.
@EastZeus