Revision control
Copy as Markdown
Other Tools
#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);
}