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-07 09:22:09 +0000
committerDaniel Baumann <daniel.baumann@progress-linux.org>2024-04-07 09:22:09 +0000
commit43a97878ce14b72f0981164f87f2e35e14151312 (patch)
tree620249daf56c0258faa40cbdcf9cfba06de2a846 /third_party/rust/metal/examples/compute/compute-argument-buffer.metal
parentInitial commit. (diff)
downloadfirefox-43a97878ce14b72f0981164f87f2e35e14151312.tar.xz
firefox-43a97878ce14b72f0981164f87f2e35e14151312.zip
Adding upstream version 110.0.1.upstream/110.0.1upstream
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);
+}