Is CPU faster than GPU ? (when writing a lot on buffer)

[tatounee]

Firstly, please tell me if this isn't the good place to ask this question that is more related to WGSL than wgpu.

I try to simulate ants moving on a graph, and because each ant is independent, I through that it will a good idea to use GPU to simulate those ants in parallel.
But when comparing the execution time between my code on CPU and GPU, I find that until 150 ant simulated, the CPU was faster than the GPU.

One hypothesis I have on what is happening is that, because I write every movement of every ant on the same buffer, the GPU has some difficulty to synchronize the code on each workgroup.

Here is approximately my code :

@group(0) @binding(0) var<storage, read> graph : array<Edge>;

@group(0) @binding(1) var<storage, read_write> paths : array<u32>;

const start = 42;
const end = 666;

@compute
@workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let graph_len = arrayLength(&graph);

    var pos = start;
    while pos != end {
        let new_pos = choose_neighbour(pos);
        paths[global_id.x * graph_len + pos] = new_pos; // here I write to the buffer `paths` at each iteration
        pos = new_pos;
    }
}

Also, I don't really understand how workgroup and workgroup_size work, so maybe the issue come from here.

Do you have any idea why the GPU is slower than the GPU, and what can I do to properly use the computing power of the GPU ?

[JustAnotherCodemonkey]

I don't know why exactly the CPU is faster although I have some theories but I can tell you 1. Yes this is a great place for this (even though it's not nearly checked enough) and 2. I actually did quite a bit of research on workgroups for my PR on adding examples. You can find my example for demonstrating and explaining workgroups at https://github.com/JustAnotherCodemonkey/wgpu/tree/JustAnotherCodemonkey/add-examples/examples/hello-workgroups. I hope that helps and if that example isn't very helpful, please tell me what I can improve on before the PR gets merged.

[SludgePhD]

The total workgroup size (all dimensions of @workgroup_size multiplied) should be at least 64 in order to fully utilize the GPU, since compute units in many GPUs are typically designed to run up to 64 instances of the same program in lockstep. If the workgroup size is smaller than what the GPU is designed for, the extra GPU resources will sit around and do nothing.

But in general, if you only run a few dozen or hundred invocations of your shader, the performance will not be very impressive. There is a decent amount of driver overhead associated with every dispatch, and even mid-range GPUs have enough cores to run 2000-3000 instances of your shader at once, let alone sequentially.

You should also watch out for potential access conflicts to the paths buffer from different shader invocations: Just like how Rust prevents simultaneous read+write access to a value (as well as two or more write accesses at once), shaders also have to watch out for them. If this is your first time writing a compute shader, I'd suggest choosing a problem that does not require any buffers with read_write access, and where you can easily tell that two shader invocations will only ever write to disjoint parts of the output buffer/texture. There are generally no ordering guarantees between shader invocations, and the only synchronization mechanism GPUs have are workgroup barriers that will only synchronize the threads within a single workgroup, so if one invocation of your shader depends on data written by another, you'll likely run into trouble.