third_party/rust/metal/examples/compute/compute-argument-buffer.metal


1
2
3
4
5
6
7
8
9
10
11
12
13
14

#include <metal_stdlib>

using namespace metal;

struct SumInput {
    device uint *data;
    volatile device atomic_uint *sum;
};

kernel void sum(device SumInput& input [[ buffer(0) ]],
                uint gid [[ thread_position_in_grid ]])
{
    atomic_fetch_add_explicit(input.sum, input.data[gid], memory_order_relaxed);
}