Compute dispatches in one submit produce all-zero output on RADV STRIX1 (Phoenix iGPU)

[kvark]

Summary

Chaining a handful of large compute dispatches into a single submit on
RADV STRIX1 (AMD Radeon 890M, Phoenix GFX11.5 iGPU) produces
deterministically all-zero output at the end of the chain, regardless
of whether the dispatches share one compute pass with pass.barrier()
between groups or are split into one pass per dispatch. Putting the
identical dispatches into separate submits with vkWaitForFences between
each one produces correct output.

Threshold: 7 chained conv-T + slice dispatches is enough to trigger
the bug. Each conv-T writes ~30 MiB into a temp buffer; slice copies
~14 MiB back. The shader is meganeura's conv2d_grad_input_hw.wgsl
(verbatim — it's just a tiled conv-transpose), workgroup_size(16,16,1),
dispatch [61, 4, 128].

single mode (one pass, barrier-between-groups) at N=6 reliably
triggers a GPU hang / TDR:

radv/amdgpu: The CS has been cancelled because the context is lost.
This context is guilty of a hard recovery.

multipass mode (one submit, one encoder.compute() per dispatch ⇒
blade emits its standard vkCmdPipelineBarrier at each pass boundary)
silently produces zeros at N ≥ 7. split mode (one submit per
dispatch + vkWaitForFences) always gives correct output, up to at
least N=20.

Originally found while porting Google's magenta-realtime SpectroStream
decoder to meganeura — the decoder produced all-zero output past block
5. I've reduced it to a pure-blade repro that uses only meganeura's
verbatim conv-T shader and a slice2d shader, with no meganeura runtime
involvement.

Environment

Device:    AMD Radeon 890M Graphics (RADV STRIX1)   (Phoenix, RDNA 3.5)
Driver:    Mesa 26.0.2
Vulkan:    1.4.335
Blade rev: 72c30c1

Have not tested other RADV GPUs / driver versions.

Reproduction

Three execution modes record the exact same N (conv-T → slice) chain:

mode	submits	barrier between dispatches
split	2N	vkWaitForFences (queue drain)
multipass	1	one compute() per dispatch ⇒ blade emits MemoryBarrier at begin_pass
single	1	one compute(), pass.barrier() between every dispatch

All three are equivalent under the Vulkan spec.

$ N=7 cargo run --release split
Device: AMD Radeon 890M Graphics (RADV STRIX1)
small=3686400 big=7634432 (each conv-T writes 7634432 elems = 29 MiB)
Chaining 7 conv-T + slice ops
conv-T wgs=[61, 4, 128], slice wgs=[14400, 1, 1]
  mode=split     nz=3686400/3686400 (100.0%) sum=5.6694e-2 max_abs=5.0715e-4

$ N=7 cargo run --release multipass
[same setup]
  mode=multipass nz=0/3686400 (0.0%) sum=0.0000e0 max_abs=0.0000e0

$ N=6 cargo run --release single
[same setup]
  mode=split     nz=3686400/3686400 (100.0%) sum=-3.7416e-2 max_abs=1.9642e-3
  mode=multipass nz=3686400/3686400 (100.0%) sum=-3.7416e-2 max_abs=1.9642e-3
radv/amdgpu: The CS has been cancelled because the context is lost.
This context is guilty of a hard recovery.
thread 'main' panicked: GPU has crashed, and no debug information is available.

Deterministic across runs.

Sync validation layer flags nothing in either failing mode.

What the barrier blade emits looks like

blade-graphics/src/vulkan/command.rs (rev 72c30c1, ~line 369):

fn barrier(&mut self) {
    let wa = &self.device.workarounds;
    let barrier = vk::MemoryBarrier {
        src_access_mask: vk::AccessFlags::MEMORY_WRITE | wa.extra_sync_src_access,
        dst_access_mask: vk::AccessFlags::MEMORY_READ
            | vk::AccessFlags::MEMORY_WRITE
            | wa.extra_sync_dst_access,
        ..Default::default()
    };
    unsafe {
        self.device.core.cmd_pipeline_barrier(
            self.buffers[0].raw,
            vk::PipelineStageFlags::ALL_COMMANDS,
            vk::PipelineStageFlags::ALL_COMMANDS,
            vk::DependencyFlags::empty(),
            &[barrier], &[], &[],
        );
    }
}

By the spec this is a maximal global memory barrier. Yet the
write-then-read chain breaks somewhere around the 6th–7th iteration on
this device.

What I've already ruled out

• Shader race in my code. The shader is meganeura's
  conv2d_grad_input_hw.wgsl, used in production. I verified one
  conv-T in isolation always produces 100% non-zero output. The bug is
  contextual.
• Coop matrix shaders. This shader does not go through coop
  selection. (My SmolVLA/SpectroStream path with
  MEGANEURA_DISABLE_COOP=1 showed the same failure pattern.)
• Dispatch dimensions vs shader assumption. Carefully audited the
  shader; the iw < params.in_w && ih < params.in_h && n < params.batch
  bounds check covers all the workgroup ceiling overspill. No collisions
  across threads.
• Buffer size or count alone. The same buffers, same shapes, same
  init data — just regrouped into separate submits — works fine.
• f16 / cooperative matrix workgroup arrays. Not used by this
  shader. (Naga does emit invalid SPIR-V workgroup ArrayStride
  decorations for f16 workgroup arrays elsewhere — flagged by
  spirv-val — but unrelated to this case.)

Hypothesis

Either:

1. Phoenix GFX11.5 L2/scalar cache requires a finer-grained flush that
   MEMORY_WRITE → MEMORY_READ doesn't actually emit on RADV — a Mesa
   bug to escalate upstream.
2. There's a queue-internal staging state on this IP that only drains
   on submit boundary, and the maximal in-stream barrier is
   mis-translated.

The GPU hang in single mode at N=6 suggests the in-stream barrier may
also fail to order the dispatches enough to keep the command processor
in a valid state on this iGPU, not just to flush caches.

Possible workarounds in blade

(Haven't tried — flagging as next steps.)

• Use Synchronization2 with explicit
  SHADER_STORAGE_WRITE → SHADER_STORAGE_READ access masks (in case
  RADV's translation of MEMORY_* is partial).
• Add an optional auto-split policy: end the current command buffer
  and submit + fence-wait when the cumulative write footprint of the
  pass crosses N MiB. Coarse, but matches the only mode that works
  on this device.

Happy to send a PR for either if you have a preference.

Source

Drop these three files into a fresh directory and cargo run --release.

blade_radv_repro/
├── Cargo.toml
└── src/
    ├── main.rs
    └── shaders.wgsl

Usage

$ cargo run --release split        # one-submit-per-dispatch, fence wait — correct
$ N=7 cargo run --release multipass  # one submit, blade barrier per dispatch — ZERO
$ N=6 cargo run --release single     # one submit, one pass, pass.barrier() — GPU HANG / TDR

Exit code 2 = bug reproduced (split non-zero while another mode is zero).

The chain repeats conv-T (29 MiB output) → slice (14 MiB output) N times.

---

Cargo.toml

[package]
name = "blade_radv_repro"
version = "0.1.0"
edition = "2021"

[dependencies]
blade-graphics = { git = "https://github.com/kvark/blade", rev = "72c30c1" }
blade-macros = "0.3"
bytemuck = { version = "1", features = ["derive"] }
env_logger = "0.11"

[[bin]]
name = "blade_radv_repro"
path = "src/main.rs"

---

src/shaders.wgsl

// Conv2d backward w.r.t. input with SEPARATE stride_h, stride_w.
// Doubles as ConvTranspose2D with asymmetric strides (used by SpectroStream).
// Copied verbatim from meganeura/src/shaders/conv2d_grad_input_hw.wgsl.

struct Params {
    batch: u32,
    in_channels: u32,
    in_h: u32,
    in_w: u32,
    out_channels: u32,
    kernel_h: u32,
    kernel_w: u32,
    stride_h: u32,
    padding_h: u32,
    out_h: u32,
    out_w: u32,
    padding_w: u32,
    stride_w: u32,
    _pad0: u32,
    _pad1: u32,
    _pad2: u32,
}

var<storage> grad_out: array<f32>;
var<storage> weight: array<f32>;
var<storage, read_write> dst: array<f32>;
var<uniform> params: Params;

var<workgroup> wg_weight: array<f32, 49>;

@compute @workgroup_size(16, 16, 1)
fn conv_t(
    @builtin(global_invocation_id) gid: vec3<u32>,
    @builtin(local_invocation_id) lid: vec3<u32>,
) {
    let iw = gid.x;
    let ih = gid.y;
    let nci = gid.z;

    let n = nci / params.in_channels;
    let ci = nci % params.in_channels;
    let in_bounds = iw < params.in_w && ih < params.in_h && n < params.batch;

    let tid = lid.y * 16u + lid.x;
    let kernel_size = params.kernel_h * params.kernel_w;
    let i_padding_h = i32(params.padding_h);
    let i_padding_w = i32(params.padding_w);
    let i_stride_h = i32(params.stride_h);
    let i_stride_w = i32(params.stride_w);

    var sum = 0.0;

    for (var co = 0u; co < params.out_channels; co++) {
        if tid < kernel_size {
            wg_weight[tid] = weight[(co * params.in_channels + ci) * kernel_size + tid];
        }
        workgroupBarrier();

        if in_bounds {
            let go_base = (n * params.out_channels + co) * params.out_h * params.out_w;

            for (var kh = 0u; kh < params.kernel_h; kh++) {
                let h_off = i32(ih) + i_padding_h - i32(kh);
                if h_off >= 0 && (h_off % i_stride_h) == 0 {
                    let oh = u32(h_off) / params.stride_h;
                    if oh < params.out_h {
                        for (var kw = 0u; kw < params.kernel_w; kw++) {
                            let w_off = i32(iw) + i_padding_w - i32(kw);
                            if w_off >= 0 && (w_off % i_stride_w) == 0 {
                                let ow = u32(w_off) / params.stride_w;
                                if ow < params.out_w {
                                    sum += grad_out[go_base + oh * params.out_w + ow]
                                         * wg_weight[kh * params.kernel_w + kw];
                                }
                            }
                        }
                    }
                }
            }
        }

        workgroupBarrier();
    }

    if in_bounds {
        let in_idx = ((n * params.in_channels + ci) * params.in_h + ih) * params.in_w + iw;
        dst[in_idx] = sum;
    }
}

// Slice2D — input[N,C,in_h,in_w] → output[N,C, in_h-start_h-end_h, in_w-start_w-end_w]
struct SliceParams {
    batch: u32,
    channels: u32,
    in_h: u32,
    in_w: u32,
    start_h: u32,
    end_h: u32,
    start_w: u32,
    end_w: u32,
}

var<storage> slice_src: array<f32>;
var<storage, read_write> slice_dst: array<f32>;
var<uniform> slice_params: SliceParams;

@compute @workgroup_size(256)
fn slice(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    let out_h = slice_params.in_h - slice_params.start_h - slice_params.end_h;
    let out_w = slice_params.in_w - slice_params.start_w - slice_params.end_w;
    let total = slice_params.batch * slice_params.channels * out_h * out_w;
    if i >= total { return; }

    let ow = i % out_w;
    let oh = (i / out_w) % out_h;
    let c = (i / (out_w * out_h)) % slice_params.channels;
    let n = i / (slice_params.channels * out_h * out_w);

    let ih = oh + slice_params.start_h;
    let iw = ow + slice_params.start_w;
    let src_idx = ((n * slice_params.channels + c) * slice_params.in_h + ih) * slice_params.in_w + iw;
    slice_dst[i] = slice_src[src_idx];
}

---

src/main.rs

//! Pure-blade repro for the SpectroStream Block(5) zero-output bug.
//!
//! Chain N copies of:
//!   conv-T (Conv2dGradInputHW shader, dispatches [61, 4, 128] workgroups,
//!           writes 7.6M f32 ≈ 30 MiB output)
//!   slice2d (crops back to 3.6M f32 ≈ 14 MiB so next iter can chain)
//!
//! On RADV STRIX1 (AMD Radeon 890M / Phoenix), the chain produces
//! all-zero output once N ≥ 6 when recorded into a single submit.
//! Splitting into one-submit-per-group with vkWaitForFences works.

use blade_graphics as gpu;
use gpu::ShaderData as _;

#[derive(Clone, Copy, bytemuck::Zeroable, bytemuck::Pod)]
#[repr(C)]
struct ConvParams {
    batch: u32,
    in_channels: u32,
    in_h: u32,
    in_w: u32,
    out_channels: u32,
    kernel_h: u32,
    kernel_w: u32,
    stride_h: u32,
    padding_h: u32,
    out_h: u32,
    out_w: u32,
    padding_w: u32,
    stride_w: u32,
    _pad0: u32,
    _pad1: u32,
    _pad2: u32,
}

#[derive(Clone, Copy, bytemuck::Zeroable, bytemuck::Pod)]
#[repr(C)]
struct SliceParams {
    batch: u32,
    channels: u32,
    in_h: u32,
    in_w: u32,
    start_h: u32,
    end_h: u32,
    start_w: u32,
    end_w: u32,
}

#[derive(blade_macros::ShaderData)]
struct ConvData {
    grad_out: gpu::BufferPiece,
    weight: gpu::BufferPiece,
    dst: gpu::BufferPiece,
    params: gpu::BufferPiece,
}

#[derive(blade_macros::ShaderData)]
struct SliceData {
    slice_src: gpu::BufferPiece,
    slice_dst: gpu::BufferPiece,
    slice_params: gpu::BufferPiece,
}

fn main() {
    env_logger::init();
    let ctx = unsafe {
        gpu::Context::init(gpu::ContextDesc {
            validation: false,
            ..Default::default()
        })
        .expect("init")
    };
    println!("Device: {}", ctx.device_information().device_name);

    let source = include_str!("shaders.wgsl");
    let shader = ctx.create_shader(gpu::ShaderDesc { source, naga_module: None });
    let mut conv_pipe = ctx.create_compute_pipeline(gpu::ComputePipelineDesc {
        name: "conv_t",
        data_layouts: &[&ConvData::layout()],
        compute: shader.at("conv_t"),
    });
    let mut slice_pipe = ctx.create_compute_pipeline(gpu::ComputePipelineDesc {
        name: "slice",
        data_layouts: &[&SliceData::layout()],
        compute: shader.at("slice"),
    });

    // Shapes (matching SpectroStream d67).
    let batch = 1u32;
    let in_c = 128u32;
    let in_h = 60u32;
    let in_w = 480u32;
    let out_c = 128u32;
    let kh = 3u32;
    let kw = 4u32;
    let stride_h = 1u32;
    let stride_w = 2u32;
    let pad_h = 0u32;
    let pad_w = 0u32;
    let out_h = (in_h - 1) * stride_h + kh - 2 * pad_h; // 62
    let out_w = (in_w - 1) * stride_w + kw - 2 * pad_w; // 962

    let conv_in_h = out_h;   let conv_in_w = out_w;
    let conv_out_h = in_h;   let conv_out_w = in_w;
    let conv_in_c = out_c;   let conv_out_c = in_c;

    let small_size = (batch * in_c * in_h * in_w) as usize;   // 3.6M
    let big_size = (batch * out_c * out_h * out_w) as usize;  // 7.6M
    println!("small={small_size} big={big_size} (each conv-T writes {big_size} elems = {} MiB)",
        big_size * 4 / 1024 / 1024);

    let n_chain = std::env::var("N").ok().and_then(|s| s.parse().ok()).unwrap_or(6usize);
    println!("Chaining {n_chain} conv-T + slice ops");

    let mode = std::env::args().nth(1).unwrap_or_else(|| "all".into());
    let modes: &[&str] = match mode.as_str() {
        "single" => &["single"],
        "multipass" => &["multipass"],
        "split" => &["split"],
        _ => &["split", "multipass", "single"],
    };

    let conv_params = ConvParams {
        batch, in_channels: conv_in_c, in_h: conv_in_h, in_w: conv_in_w,
        out_channels: conv_out_c, kernel_h: kh, kernel_w: kw,
        stride_h, padding_h: pad_h, out_h: conv_out_h, out_w: conv_out_w,
        padding_w: pad_w, stride_w, _pad0: 0, _pad1: 0, _pad2: 0,
    };
    let conv_params_buf = ctx.create_buffer(gpu::BufferDesc {
        name: "conv_params",
        size: std::mem::size_of::<ConvParams>() as u64,
        memory: gpu::Memory::Shared,
    });
    unsafe {
        std::ptr::copy_nonoverlapping(
            &conv_params as *const ConvParams as *const u8,
            conv_params_buf.data(),
            std::mem::size_of::<ConvParams>(),
        );
    }

    let slice_params = SliceParams {
        batch, channels: in_c, in_h: out_h, in_w: out_w,
        start_h: 1, end_h: 1, start_w: 1, end_w: 481,
    };
    let slice_params_buf = ctx.create_buffer(gpu::BufferDesc {
        name: "slice_params",
        size: std::mem::size_of::<SliceParams>() as u64,
        memory: gpu::Memory::Shared,
    });
    unsafe {
        std::ptr::copy_nonoverlapping(
            &slice_params as *const SliceParams as *const u8,
            slice_params_buf.data(),
            std::mem::size_of::<SliceParams>(),
        );
    }

    let weight_buf = ctx.create_buffer(gpu::BufferDesc {
        name: "weight",
        size: (conv_out_c * conv_in_c * kh * kw * 4) as u64,
        memory: gpu::Memory::Shared,
    });
    let weight_data: Vec<f32> = (0..(conv_out_c * conv_in_c * kh * kw))
        .map(|i| (i as f32 * 0.01).cos() * 0.05)
        .collect();
    unsafe {
        std::ptr::copy_nonoverlapping(
            weight_data.as_ptr() as *const u8,
            weight_buf.data(),
            weight_data.len() * 4,
        );
    }

    let mut small_bufs: Vec<gpu::Buffer> = Vec::with_capacity(n_chain + 1);
    let mut big_bufs: Vec<gpu::Buffer> = Vec::with_capacity(n_chain);
    for i in 0..(n_chain + 1) {
        let memory = if i == 0 || i == n_chain { gpu::Memory::Shared } else { gpu::Memory::Device };
        small_bufs.push(ctx.create_buffer(gpu::BufferDesc {
            name: "small",
            size: (small_size * 4) as u64,
            memory,
        }));
    }
    for _ in 0..n_chain {
        big_bufs.push(ctx.create_buffer(gpu::BufferDesc {
            name: "big",
            size: (big_size * 4) as u64,
            memory: gpu::Memory::Device,
        }));
    }

    let init_data: Vec<f32> = (0..small_size).map(|i| (i as f32 * 0.001).sin() + 1.0).collect();
    unsafe {
        std::ptr::copy_nonoverlapping(
            init_data.as_ptr() as *const u8, small_bufs[0].data(), init_data.len() * 4,
        );
    }

    let conv_dispatch_wgs = [conv_in_w.div_ceil(16), conv_in_h.div_ceil(16), batch * conv_in_c];
    let slice_total = small_size as u32;
    let slice_dispatch_wgs = [slice_total.div_ceil(256), 1, 1];
    println!("conv-T wgs={:?}, slice wgs={:?}", conv_dispatch_wgs, slice_dispatch_wgs);

    let mut encoder = ctx.create_command_encoder(gpu::CommandEncoderDesc {
        name: "repro", buffer_count: 1,
    });

    let run = |encoder: &mut gpu::CommandEncoder, ctx: &gpu::Context, mode: &str| {
        match mode {
            "single" => {
                encoder.start();
                {
                    let mut pass = encoder.compute("step");
                    for i in 0..n_chain {
                        if i > 0 { pass.barrier(); }
                        let mut pc = pass.with(&conv_pipe);
                        pc.bind(0, &ConvData {
                            grad_out: small_bufs[i].into(),
                            weight: weight_buf.into(),
                            dst: big_bufs[i].into(),
                            params: conv_params_buf.into(),
                        });
                        pc.dispatch(conv_dispatch_wgs);
                        pass.barrier();
                        let mut pc = pass.with(&slice_pipe);
                        pc.bind(0, &SliceData {
                            slice_src: big_bufs[i].into(),
                            slice_dst: small_bufs[i + 1].into(),
                            slice_params: slice_params_buf.into(),
                        });
                        pc.dispatch(slice_dispatch_wgs);
                    }
                }
                let sp = ctx.submit(encoder);
                let _ = ctx.wait_for(&sp, !0);
            }
            "multipass" => {
                encoder.start();
                for i in 0..n_chain {
                    { let mut pass = encoder.compute("conv_t");
                      let mut pc = pass.with(&conv_pipe);
                      pc.bind(0, &ConvData {
                          grad_out: small_bufs[i].into(),
                          weight: weight_buf.into(),
                          dst: big_bufs[i].into(),
                          params: conv_params_buf.into(),
                      });
                      pc.dispatch(conv_dispatch_wgs);
                    }
                    { let mut pass = encoder.compute("slice");
                      let mut pc = pass.with(&slice_pipe);
                      pc.bind(0, &SliceData {
                          slice_src: big_bufs[i].into(),
                          slice_dst: small_bufs[i + 1].into(),
                          slice_params: slice_params_buf.into(),
                      });
                      pc.dispatch(slice_dispatch_wgs);
                    }
                }
                let sp = ctx.submit(encoder);
                let _ = ctx.wait_for(&sp, !0);
            }
            "split" => {
                for i in 0..n_chain {
                    encoder.start();
                    { let mut pass = encoder.compute("conv_t");
                      let mut pc = pass.with(&conv_pipe);
                      pc.bind(0, &ConvData {
                          grad_out: small_bufs[i].into(),
                          weight: weight_buf.into(),
                          dst: big_bufs[i].into(),
                          params: conv_params_buf.into(),
                      });
                      pc.dispatch(conv_dispatch_wgs);
                    }
                    let sp = ctx.submit(encoder);
                    let _ = ctx.wait_for(&sp, !0);

                    encoder.start();
                    { let mut pass = encoder.compute("slice");
                      let mut pc = pass.with(&slice_pipe);
                      pc.bind(0, &SliceData {
                          slice_src: big_bufs[i].into(),
                          slice_dst: small_bufs[i + 1].into(),
                          slice_params: slice_params_buf.into(),
                      });
                      pc.dispatch(slice_dispatch_wgs);
                    }
                    let sp = ctx.submit(encoder);
                    let _ = ctx.wait_for(&sp, !0);
                }
            }
            _ => panic!("unknown mode {mode}"),
        }
    };

    println!("--- results ---");
    let mut results = Vec::new();
    for &m in modes {
        run(&mut encoder, &ctx, m);
        let final_buf = &small_bufs[n_chain];
        let slice = unsafe {
            std::slice::from_raw_parts(final_buf.data() as *const f32, small_size)
        };
        let nz = slice.iter().filter(|&&v| v != 0.0 && v.is_finite()).count();
        let sum: f64 = slice.iter().map(|&v| v as f64).sum();
        let max_abs = slice.iter().fold(0.0_f32, |a, &v| a.max(v.abs()));
        println!(
            "  mode={m:9} nz={nz}/{small_size} ({:.1}%) sum={sum:.4e} max_abs={max_abs:.4e}",
            100.0 * nz as f32 / small_size as f32
        );
        results.push((m, nz, sum));
    }

    let any_zero = results.iter().any(|&(_, nz, _)| nz == 0);
    let any_nonzero = results.iter().any(|&(_, nz, _)| nz > 0);
    if any_zero && any_nonzero {
        println!("--- BUG REPRODUCED: at least one mode gave all-zero output");
        println!("--- while another mode produced non-zero data");
        std::process::exit(2);
    } else if any_zero {
        println!("--- All modes gave zero");
    } else {
        println!("--- All modes produced non-zero output");
    }

    ctx.destroy_compute_pipeline(&mut conv_pipe);
    ctx.destroy_compute_pipeline(&mut slice_pipe);
    ctx.destroy_command_encoder(&mut encoder);
}