summaryrefslogtreecommitdiffstats
path: root/third_party/rust/metal/examples/compute/compute-argument-buffer.metal
diff options
context:
space:
mode:
authorDaniel Baumann <daniel.baumann@progress-linux.org>2024-04-19 00:47:55 +0000
committerDaniel Baumann <daniel.baumann@progress-linux.org>2024-04-19 00:47:55 +0000
commit26a029d407be480d791972afb5975cf62c9360a6 (patch)
treef435a8308119effd964b339f76abb83a57c29483 /third_party/rust/metal/examples/compute/compute-argument-buffer.metal
parentInitial commit. (diff)
downloadfirefox-26a029d407be480d791972afb5975cf62c9360a6.tar.xz
firefox-26a029d407be480d791972afb5975cf62c9360a6.zip
Adding upstream version 124.0.1.upstream/124.0.1
Signed-off-by: Daniel Baumann <daniel.baumann@progress-linux.org>
Diffstat (limited to 'third_party/rust/metal/examples/compute/compute-argument-buffer.metal')
-rw-r--r--third_party/rust/metal/examples/compute/compute-argument-buffer.metal14
1 files changed, 14 insertions, 0 deletions
diff --git a/third_party/rust/metal/examples/compute/compute-argument-buffer.metal b/third_party/rust/metal/examples/compute/compute-argument-buffer.metal
new file mode 100644
index 0000000000..1dcc79daf5
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/compute-argument-buffer.metal
@@ -0,0 +1,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);
+}