[Bug] CUDA error (Steam-layers)

[SolicTous]

Git commit

2d40a8b

Operating System & Version

Win 10

GGML backends

CUDA (max installed v12.6)

Command-line arguments used

"../sd-server.exe" -v --listen-ip 0.0.0.0 --listen-port 5777 --lora-model-dir ../loras/ --diffusion-model ../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf --vae ../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf --llm ../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf --offload-to-cpu --max-vram -1 --stream-layers --diffusion-fa --cfg-scale 1.0 --steps 10 --vae-tiling --vae-tile-overlap 0.125 -H 1024 -W 1024

Steps to reproduce

run server with stream-layers

request with any propmpt and 1024x1024 size

What you expected to happen

generation as always

without stream-layers it works

What actually happened

crash
D:\a\stable-diffusion.cpp\stable-diffusion.cpp\ggml\src\ggml-cuda\ggml-cuda.cu:102: CUDA error

Logs / error messages / stack trace

D:\SDCPP\runserv>"../sd-server.exe" -v --listen-ip 0.0.0.0 --listen-port 5777 --lora-model-dir ../loras/ --diffusion-model ../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf --vae ../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf --llm ../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf --offload-to-cpu --max-vram -1 --stream-layers --diffusion-fa --cfg-scale 1.0 --steps 10 --vae-tiling --vae-tile-overlap 0.125 -H 1024 -W 1024
[DEBUG] main.cpp:82 - version: stable-diffusion.cpp version unknown, commit 2d40a8b
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 12281 MiB):
Device 0: NVIDIA GeForce RTX 4070 Ti, compute capability 8.9, VMM: yes, VRAM: 12281 MiB
[DEBUG] main.cpp:83 - System Info:
SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | OPENMP = 1 | REPACK = 1 |
[DEBUG] main.cpp:84 - SDSvrParams {
listen_ip: 0.0.0.0,
listen_port: "5777",
serve_html_path: "",
}
[DEBUG] main.cpp:85 - SDContextParams {
n_threads: 8,
model_path: "",
clip_l_path: "",
clip_g_path: "",
clip_vision_path: "",
t5xxl_path: "",
llm_path: "../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf",
llm_vision_path: "",
diffusion_model_path: "../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf",
high_noise_diffusion_model_path: "",
embeddings_connectors_path: "",
vae_path: "../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf",
vae_format: "auto",
audio_vae_path: "",
taesd_path: "",
esrgan_path: "",
control_net_path: "",
embedding_dir: "",
embeddings: {
}
wtype: NONE,
tensor_type_rules: "",
lora_model_dir: "../loras/",
hires_upscalers_dir: "",
photo_maker_path: "",
rng_type: cuda,
sampler_rng_type: NONE,
offload_params_to_cpu: true,
max_vram: -1,
stream_layers: true,
backend: "",
params_backend: "",
enable_mmap: false,
control_net_cpu: false,
clip_on_cpu: false,
vae_on_cpu: false,
flash_attn: false,
diffusion_flash_attn: true,
diffusion_conv_direct: false,
vae_conv_direct: false,
circular: false,
circular_x: false,
circular_y: false,
chroma_use_dit_mask: true,
qwen_image_zero_cond_t: false,
chroma_use_t5_mask: false,
chroma_t5_mask_pad: 1,
prediction: NONE,
lora_apply_mode: auto,
force_sdxl_vae_conv_scale: false
}
[DEBUG] main.cpp:86 - SDGenerationParams {
loras: "{
}",
high_noise_loras: "{
}",
prompt: "",
negative_prompt: "",
clip_skip: -1,
width: 1024,
height: 1024,
batch_count: 1,
init_image_path: "",
end_image_path: "",
mask_image_path: "",
control_image_path: "",
ref_image_paths: [],
control_video_path: "",
auto_resize_ref_image: true,
increase_ref_index: false,
pm_id_images_dir: "",
pm_id_embed_path: "",
pm_style_strength: 20,
skip_layers: [7, 8, 9],
sample_params: (txt_cfg: 1.00, img_cfg: 1.00, distilled_guidance: 3.50, slg.layer_count: 0, slg.layer_start: 0.01, slg.layer_end: 0.20, slg.scale: 0.00, scheduler: NONE, sample_method: NONE, sample_steps: 10, eta: inf, shifted_timestep: 0, flow_shift: inf, extra_sample_args: ),
high_noise_skip_layers: [7, 8, 9],
high_noise_sample_params: (txt_cfg: 7.00, img_cfg: 7.00, distilled_guidance: 3.50, slg.layer_count: 0, slg.layer_start: 0.01, slg.layer_end: 0.20, slg.scale: 0.00, scheduler: NONE, sample_method: NONE, sample_steps: 20, eta: inf, shifted_timestep: 0, flow_shift: inf, extra_sample_args: ),
custom_sigmas: [],
cache_mode: "",
cache_option: "",
cache: disabled (threshold=inf, start=0.15, end=0.95),
moe_boundary: 0.875,
video_frames: 1,
fps: 16,
vace_strength: 1,
strength: 0.75,
control_strength: 0.9,
seed: 42,
upscale_repeats: 1,
upscale_tile_size: 128,
hires: { enabled: false, upscaler: "Latent", model_path: "", scale: 2, target_width: 0, target_height: 0, steps: 0, denoising_strength: 0.7, custom_sigmas: [], upscale_tile_size: 128 },
vae_tiling_params: { 1, 0, 0, 0, 0.125, 0, 0, "" },
}
[DEBUG] ggml_extend_backend.cpp:367 - Found 2 backend devices:
[DEBUG] ggml_extend_backend.cpp:370 - #0: CUDA0
[DEBUG] ggml_extend_backend.cpp:370 - #1: CPU
[DEBUG] ggml_extend_backend.cpp:292 - Initializing backend: CUDA0
[DEBUG] ggml_extend_backend.cpp:292 - Initializing backend: CPU
[INFO ] ggml_graph_cut.cpp:127 - --max-vram < 0 auto-detected 10.78 GiB free VRAM (11.99 GiB total), reserving 1.00 GiB; using 9.78 GiB
[INFO ] stable-diffusion.cpp:305 - loading diffusion model from '../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf'
[INFO ] model.cpp:216 - load ../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf using gguf format
[DEBUG] model.cpp:265 - init from '../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf'
[INFO ] stable-diffusion.cpp:352 - loading llm from '../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf'
[INFO ] model.cpp:216 - load ../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf using gguf format
[DEBUG] model.cpp:265 - init from '../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf'
[INFO ] stable-diffusion.cpp:366 - loading vae from '../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf'
[INFO ] model.cpp:216 - load ../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf using gguf format
[DEBUG] model.cpp:265 - init from '../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf'
[INFO ] stable-diffusion.cpp:408 - Version: Z-Image
[INFO ] stable-diffusion.cpp:436 - Weight type stat: f32: 570 | f16: 70 | q4_0: 192 | q8_0: 263
[INFO ] stable-diffusion.cpp:437 - Conditioner weight type stat: f32: 145 | q8_0: 253
[INFO ] stable-diffusion.cpp:438 - Diffusion model weight type stat: f32: 251 | q4_0: 192 | q8_0: 10
[INFO ] stable-diffusion.cpp:439 - VAE weight type stat: f32: 174 | f16: 70
[DEBUG] stable-diffusion.cpp:441 - ggml tensor size = 400 bytes
[DEBUG] qwen2_tokenizer.cpp:14 - merges size 151387
[DEBUG] qwen2_tokenizer.cpp:39 - vocab size: 151674
[DEBUG] llm.hpp:1568 - llm: num_layers = 36, vocab_size = 151936, hidden_size = 2560, intermediate_size = 9728
[INFO ] stable-diffusion.cpp:855 - using VAE for encoding / decoding
[INFO ] auto_encoder_kl.hpp:525 - vae decoder: ch = 128
[INFO ] stable-diffusion.cpp:957 - Using flash attention in the diffusion model
[DEBUG] stable-diffusion.cpp:986 - loading weights
[DEBUG] ggml_extend.hpp:3014 - qwen3 params backend buffer size = 4076.43 MB(RAM) (398 tensors)
[DEBUG] ggml_extend.hpp:3014 - z_image params backend buffer size = 3512.71 MB(RAM) (453 tensors)
[DEBUG] ggml_extend.hpp:3014 - vae params backend buffer size = 160.00 MB(RAM) (244 tensors)
[INFO ] model.cpp:814 - NOT using mmap for '../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf' (mmap disabled by caller)
[INFO ] model.cpp:814 - NOT using mmap for '../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf' (mmap disabled by caller)
[INFO ] model.cpp:814 - NOT using mmap for '../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf' (mmap disabled by caller)
[INFO ] model.cpp:825 - model files processing completed in 0.00s
[DEBUG] model.cpp:929 - using 8 threads for model loading
[DEBUG] model.cpp:945 - loading tensors from ../../Models/Diffusers/Transformer/Z-Image/z_image_turbo-Q4_0.gguf
|====================> | 453/1095 - 4.19GB/s
[DEBUG] model.cpp:945 - loading tensors from ../../Models/Diffusers/LLM/Qwen/Qwen3-4B-Instruct-2507-Q8_0.gguf
|======================================> | 851/1095 - 4.50GB/s
[DEBUG] model.cpp:945 - loading tensors from ../../Models/Diffusers/Vae/pig_flux_vae_fp32-f16.gguf
|==================================================| 1095/1095 - 4.08GB/s
[INFO ] model.cpp:1169 - loading tensors completed, taking 1.86s (read: 1.49s, memcpy: 0.00s, convert: 0.00s, copy_to_backend: 0.00s)
[DEBUG] stable-diffusion.cpp:1088 - finished loaded file
[INFO ] stable-diffusion.cpp:1170 - total params memory size = 7749.14MB (VRAM 0.00MB, RAM 7749.14MB): text_encoders 4076.43MB(RAM), diffusion_model 3512.71MB(RAM), vae 160.00MB(RAM), controlnet 0.00MB(N/A), pmid 0.00MB(N/A)
[INFO ] stable-diffusion.cpp:1257 - running in FLOW mode
[INFO ] main.cpp:148 - listening on: http://0.0.0.0:5777
[INFO ] stable-diffusion.cpp:4449 - generate_image 1024x1024
[INFO ] denoiser.hpp:579 - get_sigmas with discrete scheduler
[INFO ] stable-diffusion.cpp:3510 - sampling using Euler method
[DEBUG] conditioner.hpp:1854 - parse '<|im_start|>user
Paper Dragon<|im_end|>
<|im_start|>assistant
' to [['<|im_start|>user
', 1], ['Paper Dragon', 1], ['<|im_end|>
<|im_start|>assistant
', 1], ]
[DEBUG] bpe_tokenizer.cpp:207 - split prompt "<|im_start|>user
" to tokens ["<|im_start|>", "user", "Ċ", ]
[DEBUG] bpe_tokenizer.cpp:207 - split prompt "Paper Dragon" to tokens ["Paper", "ĠDragon", ]
[DEBUG] bpe_tokenizer.cpp:207 - split prompt "<|im_end|>
<|im_start|>assistant
" to tokens ["<|im_end|>", "Ċ", "<|im_start|>", "assistant", "Ċ", ]
[INFO ] ggml_graph_cut.cpp:731 - qwen3 build cached graph cut plan done (taking 1 ms)
[INFO ] ggml_graph_cut.cpp:698 - qwen3 graph cut max_vram=10012.00 MB merged 36 segments -> 1 segments
[INFO ] ggml_graph_cut.cpp:704 - qwen3 graph cut max_vram budget merge took 14 ms
[DEBUG] ggml_extend.hpp:1914 - qwen3 compute buffer size: 1.09 MB(VRAM)
[INFO ] ggml_extend.hpp:2154 - qwen3 offload params (4076.43 MB, 398 tensors) to runtime backend (CUDA0), taking 0.91s
[DEBUG] conditioner.hpp:2256 - computing condition graph completed, taking 1089 ms
[INFO ] stable-diffusion.cpp:4205 - get_learned_condition completed, taking 1.09s
[INFO ] stable-diffusion.cpp:4483 - generating image: 1/1 - seed 1
[INFO ] ggml_graph_cut.cpp:731 - z_image build cached graph cut plan done (taking 2 ms)
[INFO ] ggml_graph_cut.cpp:698 - z_image graph cut max_vram=10012.00 MB merged 36 segments -> 1 segments
[INFO ] ggml_graph_cut.cpp:704 - z_image graph cut max_vram budget merge took 29 ms
[INFO ] ggml_extend.hpp:2466 - z_image streaming budget = 10012.00 MB
[DEBUG] ggml_extend.hpp:1914 - z_image compute buffer size: 2665.29 MB(VRAM)
[INFO ] ggml_extend.hpp:2154 - z_image offload params (3512.75 MB, 453 tensors) to runtime backend (CUDA0), taking 0.58s
|=====> | 1/10 - 1.70s/it[INFO ] ggml_extend.hpp:2447 - z_image clamping streaming budget: actual free VRAM 4268.00 MB < user cap 10012.00 MB
[INFO ] ggml_graph_cut.cpp:698 - z_image graph cut max_vram=4268.00 MB merged 36 segments -> 3 segments
[INFO ] ggml_graph_cut.cpp:704 - z_image graph cut max_vram budget merge took 15 ms
[INFO ] ggml_extend.hpp:2466 - z_image streaming budget = 4268.00 MB
[INFO ] ggml_extend.hpp:2372 - z_image offload resident params (1759.32 MB, 214 tensors) to runtime backend (CUDA0)
[DEBUG] ggml_extend.hpp:2876 - z_image streaming-cut executing segment 1/3: z_image.prelude..z_image.layers.10 (residency=STREAMED)
[DEBUG] ggml_extend.hpp:2881 - z_image streaming: no registry entry for group 'z_image.prelude..z_image.layers.10' (using upstream offload path)
[DEBUG] ggml_extend.hpp:1914 - z_image compute buffer size: 1516.38 MB(VRAM)
[ERROR] ggml_extend.hpp:70 - CUDA error: an illegal memory access was encountered
[ERROR] ggml_extend.hpp:70 - current device: 0, in function ggml_backend_cuda_synchronize at D:\a\stable-diffusion.cpp\stable-diffusion.cpp\ggml\src\ggml-cuda\ggml-cuda.cu:3235
[ERROR] ggml_extend.hpp:70 - cudaStreamSynchronize(cuda_ctx->stream())
D:\a\stable-diffusion.cpp\stable-diffusion.cpp\ggml\src\ggml-cuda\ggml-cuda.cu:102: CUDA error

Additional context / environment details

version: stable-diffusion.cpp version unknown, commit 2d40a8b
Device 0: NVIDIA GeForce RTX 4070 Ti
System Info:
SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | OPENMP = 1 | REPACK = 1

[wbruna]

When trying to reproduce this, I hit memory errors both on ROCm and Vulkan:

[INFO ] ggml_graph_cut.cpp:731  - qwen3 build cached graph cut plan done (taking 1 ms)
[INFO ] ggml_graph_cut.cpp:694  - qwen3 graph cut max_vram=14336.00 MB merged 29 segments -> 1 segments
[INFO ] ggml_graph_cut.cpp:702  - qwen3 graph cut max_vram budget merge took 10 ms
[DEBUG] ggml_extend.hpp:1911 - qwen3 compute buffer size: 289.00 MB(VRAM)
[INFO ] ggml_extend.hpp:2149 - qwen3 offload params (6445.43 MB, 398 tensors) to runtime backend (ROCm0), taking 1.28s
[DEBUG] conditioner.hpp:2256 - computing condition graph completed, taking 1779 ms
[INFO ] stable-diffusion.cpp:4205 - get_learned_condition completed, taking 1.78s
[INFO ] stable-diffusion.cpp:4483 - generating image: 1/1 - seed 6102819
[INFO ] ggml_graph_cut.cpp:731  - flux build cached graph cut plan done (taking 1 ms)
[INFO ] ggml_graph_cut.cpp:694  - flux graph cut max_vram=14336.00 MB merged 34 segments -> 1 segments
[INFO ] ggml_graph_cut.cpp:702  - flux graph cut max_vram budget merge took 21 ms
[INFO ] ggml_extend.hpp:2464 - flux streaming budget = 14336.00 MB
[DEBUG] ggml_extend.hpp:1911 - flux compute buffer size: 3440.04 MB(VRAM)
[INFO ] ggml_extend.hpp:2149 - flux offload params (5356.04 MB, 201 tensors) to runtime backend (ROCm0), taking 1.12s
  |================>                                 | 1/3 - 7.27s/it[INFO ] ggml_extend.hpp:2444 - flux clamping streaming budget: actual free VRAM 6648.00 MB < user cap 14336.00 MB
[INFO ] ggml_graph_cut.cpp:694  - flux graph cut max_vram=6648.00 MB merged 34 segments -> 2 segments
[INFO ] ggml_graph_cut.cpp:702  - flux graph cut max_vram budget merge took 11 ms
[INFO ] ggml_extend.hpp:2464 - flux streaming budget = 6648.00 MB
[INFO ] ggml_extend.hpp:2368 - flux offload resident params (1919.01 MB, 78 tensors) to runtime backend (ROCm0)
[DEBUG] ggml_extend.hpp:2871 - flux streaming-cut executing segment 1/2: flux.prelude..flux.single_blocks.5 (residency=STREAMED)
[DEBUG] ggml_extend.hpp:2879 - flux streaming: no registry entry for group 'flux.prelude..flux.single_blocks.5' (using upstream offload path)
[DEBUG] ggml_extend.hpp:2232 - flux offload partial params (1266.01 MB, 49 tensors) to runtime backend (ROCm0)
[DEBUG] ggml_extend.hpp:1911 - flux compute buffer size: 2184.54 MB(VRAM)
Memory access fault by GPU node-1 (Agent handle: 0x55912cab4370) on address 0x55913cbd0000. Reason: Page not present or supervisor privilege.

This is with a Klein 9b Q4 quant. It's a bit hard to reproduce because it seems to depend on the model size and --max-vram values, but it's consistent for the same parameters. It does not crash without --stream-layers. It also seems associated with the "clamping streaming budget: actual free VRAM 6648.00 MB < user cap 14336.00 MB" message (a larger Q8 quant shows a different value for the free VRAM, but crashes in the same way).

The Vulkan log:

[INFO ] ggml_extend.hpp:2464 - flux streaming budget = 14336.00 MB
[DEBUG] ggml_extend.hpp:1911 - flux compute buffer size: 3799.66 MB(VRAM)
[INFO ] ggml_extend.hpp:2149 - flux offload params (5356.04 MB, 201 tensors) to runtime backend (Vulkan1), taking 3.23s
  |================>                                 | 1/3 - 11.72s/it[INFO ] ggml_extend.hpp:2444 - flux clamping streaming budget: actual free VRAM 6592.85 MB < user cap 14336.00 MB
[INFO ] ggml_graph_cut.cpp:694  - flux graph cut max_vram=6592.85 MB merged 34 segments -> 3 segments
[INFO ] ggml_graph_cut.cpp:702  - flux graph cut max_vram budget merge took 10 ms
[INFO ] ggml_extend.hpp:2464 - flux streaming budget = 6592.85 MB
[INFO ] ggml_extend.hpp:2368 - flux offload resident params (1919.01 MB, 78 tensors) to runtime backend (Vulkan1)
[DEBUG] ggml_extend.hpp:2871 - flux streaming-cut executing segment 1/3: flux.prelude..flux.single_blocks.3 (residency=STREAMED)
[DEBUG] ggml_extend.hpp:2879 - flux streaming: no registry entry for group 'flux.prelude..flux.single_blocks.3' (using upstream offload path)
[DEBUG] ggml_extend.hpp:2232 - flux offload partial params (1032.01 MB, 41 tensors) to runtime backend (Vulkan1)
[DEBUG] ggml_extend.hpp:1911 - flux compute buffer size: 2637.91 MB(VRAM)
./test.sh, linha 44: 86358 Falha de segmentação  (imagem do núcleo gravada)"${ARGS[@]}"

[wbruna]

Hm, I can't reproduce this anymore on a7f2e03 . @fszontagh , any chance this could have been fixed by that change?

[fszontagh]

Yes, a7f2e03 is the squash-merge of #1598 and it bundles two pieces: the chunk-K-engaged-with-LoRA change, plus a worst-merged-segment reservation I pushed to the same branch after observing OOM under heavy chunk-K residency on bf16 Z-Image at higher resolutions. Together they keep the residency budget stable across sampling steps, so the re-clamp + re-merge mid-generation that triggered the page fault doesn't happen anymore. Glad it's fixed for you too.

[SolicTous]

looks fixed