#Problems adding a storage buffer to my compute shader

77 messages · Page 1 of 1 (latest)

night kite
#

Hello, I tried adding a storage buffer to my shader but now it crashes after a few frames with:

error: process didn't exit successfully: `target\debug\bevy_dig.exe` (exit code: 0xc0000005, STATUS_ACCESS_VIOLATION)

I'm not sure where to go from here, this is all a bit confusing to me.
Here's a link to the repo https://github.com/Sigma-dev/bevy_dig/tree/replace-array (make sure to look at the replace-array branch)

GitHub
GitHub - Sigma-dev/bevy_dig at replace-array

Contribute to Sigma-dev/bevy_dig development by creating an account on GitHub.

rancid tulip
#

Do you have renderdoc? It'd help figuring stuff like this out. Besides that, access violation more often than not means your shader tried to get something out of bounds in my experience

#

Also is there no logging besides the access violation? I'd assume in most cases that wgpu/validation would catch what you're doing wrong

night kite
#

Maybe I'll give render doc a try

night kite
#

Here is an even smaller minimal reproduction:

const SHADER_ASSET_PATH: &str = "shaders/marching_cubes.wgsl";
pub const CHUNK_WIDTH: usize = 31;
pub const INPUT_CHUNK_WIDTH: usize = CHUNK_WIDTH + 2;
pub const BUFFER_LEN_UNCOMPRESSED: usize =
    INPUT_CHUNK_WIDTH * INPUT_CHUNK_WIDTH * INPUT_CHUNK_WIDTH;
pub const BUFFER_LEN: usize = BUFFER_LEN_UNCOMPRESSED / 32;
const MAX_VERTICES_PER_CUBE: usize = 12;
const TRI_BUFFER_LEN: usize =
    (CHUNK_WIDTH + 2) * (CHUNK_WIDTH + 2) * (CHUNK_WIDTH + 2) * MAX_VERTICES_PER_CUBE;

fn main() {
    App::new()
        .add_plugins((DefaultPlugins, GpuReadbackPlugin))
        .run();
}

pub(crate) struct GpuReadbackPlugin;
impl Plugin for GpuReadbackPlugin {
    fn build(&self, _app: &mut App) {}

    fn finish(&self, app: &mut App) {
        let render_app = app.sub_app_mut(RenderApp);
        let world = render_app.world();
        let render_device = world.resource::<RenderDevice>();
        let layout = render_device.create_bind_group_layout(
            None,
            &BindGroupLayoutEntries::sequential(
                ShaderStages::COMPUTE,
                (
                    storage_buffer::<[u32; BUFFER_LEN]>(false),
                    storage_buffer::<[Vec4; TRI_BUFFER_LEN]>(false),
                    storage_buffer_read_only::<[[i32; 16]; 256]>(false),
                ),
            ),
        );
        let shader: Handle<Shader> = world.load_asset(SHADER_ASSET_PATH);
        let pipeline_cache = world.resource::<PipelineCache>();
        let _pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor {
            label: Some("GPU readback compute shader".into()),
            layout: vec![layout.clone()],
            push_constant_ranges: Vec::new(),
            shader: shader.clone(),
            shader_defs: Vec::new(),
            entry_point: "main".into(),
            zero_initialize_workgroup_memory: false,
        });
    }
}
royal stratus
#

You probably did alreay, but did you try removing the code from the compute shader, so that it just takes int the fields and does nothing with them? Then one by one do something with them?
That way you will know, if it's your setup or your shader

night kite
# royal stratus You probably did alreay, but did you try removing the code from the compute shad...

Good advice, I def know it's the shader, but I've been able to narrow it down:

@group(0) @binding(0) var<storage, read_write> input_data: array<u32, INPUT_LENGTH>;
@group(0) @binding(1) var<storage, read_write> output_data: array<vec4<f32>, OUTPUT_LENGTH>;
@group(0) @binding(2) var<storage, read> triangles_table: array<array<i32, 16>, 256>;

@compute @workgroup_size(4, 4, 4)
fn main(
    @builtin(global_invocation_id) index: vec3<u32>,
) {
    let edges = triangles_table[0];
     for (var i: u32 = 0; i <= 12 && edges[i] != -1; i += 3u) {
        let edge = edges[i];
        output_data[0] = vec4<f32>(0.);
    }
}

This still does the same error, but removing the edges[1] != -1 makes it not crash:

@group(0) @binding(0) var<storage, read_write> input_data: array<u32, INPUT_LENGTH>;
@group(0) @binding(1) var<storage, read_write> output_data: array<vec4<f32>, OUTPUT_LENGTH>;
@group(0) @binding(2) var<storage, read> triangles_table: array<array<i32, 16>, 256>;

@compute @workgroup_size(4, 4, 4)
fn main(
    @builtin(global_invocation_id) index: vec3<u32>,
) {
    let edges = triangles_table[0];
     for (var i: u32 = 0; i <= 12; i += 3u) {
        let edge = edges[i];
        output_data[0] = vec4<f32>(0.);
    }
}

Doesn't make it any less weird though :(

royal stratus
#

I think what you want is:
(i <= 12) && (edges[i] != 1)

#

oh no wait

#

Now I'm confused what even happens in this kind of condition if you try to access the data up there and down in the loop too? 🤔

#

anyway, that is probably the problem

night kite
#

So i'm pretty sure condition is properly computed, that's how most languages work afaik

royal stratus
#

but you still got the i <= 12 && edges[i] != -1; in there

night kite
#

forgot to remove it but it also crashes without it

#

mb

royal stratus
#

also you used == -1 now

#

instead of !=

night kite
#

another mistake because I did these tests before, but yeah it still crashes

#

(i redid them ofc)

royal stratus
#

Is there anything after that loop?

#

because if there isn't anything after that loop and the changed code does not compute, I'm really confused 😂

#

oh you also didnt use edge here. You used 0. instead

#

let edge = edges[i];
output_data[0] = vec4<f32>(0.);

#

So maybe you tried too many things together and you actually just needed to change the && expression

night kite
night kite
night kite
royal stratus
#

maybe you will want to change the error in "rendering" too, since the && expression will not compute the way you want it to

night kite
royal stratus
night kite
royal stratus
#

but the way you wrote it, it would compute from left to right

#

let's assume i= 0 and that edges[0] = 13
then
i <= 12 && edges[i] != -1
would become
true && edges[0] != -1
or
true && 13 != -1
that would become (if any number than 0 is considered true)
true != -1

#

I'm pretty sure you did not intend it that way

#

I think what you wanted to do is
(i <= 12) && (edges[i] != -1)

#

especially since you wrote it that way in the break statement

#

I would try to the following code and see if that computes

for (var i: u32 = 0; (i <= 12) && (edges[i] != -1); i += 3u) {}

So use the corrected for loop with an empty loop

#

then try again with the break statement and an empty loop

#

and work your way towards your code from there 🤷

night kite
#

(btw I tried your version before and it did crash so i'm pretty sure at this point

#

And our friend rust

fn main() {
    let edges = [1, 2, 3, -1, 8, 9];
    let mut i = 0;
    while (i < 7 && edges[i] != -1) {
        println!("{}", edges[i]);
        i += 1;
    }
}

1
2
3
royal stratus
#

you are 100 percent right and I should really go to bed now. I was sure the precedence of && is higher than !=

#

I think trying out the loop without any code in it, still might be helpfull 🤷

night kite
# royal stratus I think trying out the loop without any code in it, still might be helpfull 🤷

No worries

fn main(
    @builtin(global_invocation_id) index: vec3<u32>,
) {
    let edges = triangles_table[0];
     for (var i: u32 = 0; i <= 12; i += 3u) {
        if (edges[i] != -1) {
            break;
        }
        let edge = edges[i];
        output_data[0] = vec4<f32>(0.);
    }
}

Everything in the loop is mandatory, removing only one of these makes it not crash:
if (edges[i] != -1) { break; } let edge = edges[i]; output_data[0] = vec4<f32>(0.);

#

I am much confusion

royal stratus
#

especially removing the 2nd line

#

the edge created there is not even used anywhere

#

I think most compilers would not even do anything with that line

night kite
#

I pushed the new minimal reprod if you want to run it locally

#

It has no problems running on my mac btw

#

but my main pc don't like it

royal stratus
#

I will probably try it tomorrow. The bed is the only place I should go right now 😂

night kite
#

sleep well

radiant drift
#

ah, i forgot it doesn't crash on macos

#

will run on my pc in the morning with validation layers enabled

#

probably due to not being uniform

#

non-uniform control flow is typically a problem

#

but i would just try to rewrite the loop so it's uniform

night kite
# radiant drift non-uniform control flow is typically a problem

I'm not sure what uniform control flow means, from what I could gather it is about avoiding break statements, which I tried rewriting my code without them:

@compute @workgroup_size(4, 4, 4)
fn main(
    @builtin(global_invocation_id) index: vec3<u32>,
) {
    let edges = triangles_table[0];
    var found: bool = false;

    for (var i: u32 = 0; i <= 12; i += 3u) {
        if (!found && edges[i] != -1) {
            let edge = edges[i];
            output_data[0] = vec4<f32>(0.);
            found = true;
        }
    }
}

but it also crashes

radiant drift
#

could be bounds checking too

final jacinth
#

Hi, can you try saving edges[i] in a variable rather than index it twice? I'd wager the error goes away

final jacinth
#

I think it's a naga bug

night kite
final jacinth
#

That's great to hear! I also banged my head for a long time on this same issue 6months ago. I was hoping to have a go at fixing the bug, but never came around back to it

night kite
#

My mac can run it now, but the mesh is all messed up, I'll dig more into it when I have more time

night kite
final jacinth
#

I'm on linux..

#

How does it look messed up?

night kite
#

This is the same code, running on my mac (M1)

GetAttachmentThumbnail.png