1 files changed, 95 insertions, 0 deletions
diff --git a/third_party/rust/metal/examples/compute/compute-argument-buffer.rs b/third_party/rust/metal/examples/compute/compute-argument-buffer.rs
new file mode 100644
index 0000000000..97527091a3
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/compute-argument-buffer.rs
@@ -0,0 +1,95 @@
+// Copyright 2017 GFX developers
+//
+// Licensed under the Apache License, Version 2.0, <LICENSE-APACHE or
+// http://apache.org/licenses/LICENSE-2.0> or the MIT license <LICENSE-MIT or
+// http://opensource.org/licenses/MIT>, at your option. This file may not be
+// copied, modified, or distributed except according to those terms.
+
+use metal::*;
+use objc::rc::autoreleasepool;
+use std::mem;
+
+static LIBRARY_SRC: &str = include_str!("compute-argument-buffer.metal");
+
+fn main() {
+    autoreleasepool(|| {
+        let device = Device::system_default().expect("no device found");
+        let command_queue = device.new_command_queue();
+
+        let data = [
+            1u32, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23,
+            24, 25, 26, 27, 28, 29, 30,
+        ];
+
+        let buffer = device.new_buffer_with_data(
+            unsafe { mem::transmute(data.as_ptr()) },
+            (data.len() * mem::size_of::<u32>()) as u64,
+            MTLResourceOptions::CPUCacheModeDefaultCache,
+        );
+
+        let sum = {
+            let data = [0u32];
+            device.new_buffer_with_data(
+                unsafe { mem::transmute(data.as_ptr()) },
+                (data.len() * mem::size_of::<u32>()) as u64,
+                MTLResourceOptions::CPUCacheModeDefaultCache,
+            )
+        };
+
+        let command_buffer = command_queue.new_command_buffer();
+        let encoder = command_buffer.new_compute_command_encoder();
+
+        let library = device
+            .new_library_with_source(LIBRARY_SRC, &CompileOptions::new())
+            .unwrap();
+        let kernel = library.get_function("sum", None).unwrap();
+
+        let argument_encoder = kernel.new_argument_encoder(0);
+        let arg_buffer = device.new_buffer(
+            argument_encoder.encoded_length(),
+            MTLResourceOptions::empty(),
+        );
+        argument_encoder.set_argument_buffer(&arg_buffer, 0);
+        argument_encoder.set_buffer(0, &buffer, 0);
+        argument_encoder.set_buffer(1, &sum, 0);
+
+        let pipeline_state_descriptor = ComputePipelineDescriptor::new();
+        pipeline_state_descriptor.set_compute_function(Some(&kernel));
+
+        let pipeline_state = device
+            .new_compute_pipeline_state_with_function(
+                pipeline_state_descriptor.compute_function().unwrap(),
+            )
+            .unwrap();
+
+        encoder.set_compute_pipeline_state(&pipeline_state);
+        encoder.set_buffer(0, Some(&arg_buffer), 0);
+
+        encoder.use_resource(&buffer, MTLResourceUsage::Read);
+        encoder.use_resource(&sum, MTLResourceUsage::Write);
+
+        let width = 16;
+
+        let thread_group_count = MTLSize {
+            width,
+            height: 1,
+            depth: 1,
+        };
+
+        let thread_group_size = MTLSize {
+            width: (data.len() as u64 + width) / width,
+            height: 1,
+            depth: 1,
+        };
+
+        encoder.dispatch_thread_groups(thread_group_count, thread_group_size);
+        encoder.end_encoding();
+        command_buffer.commit();
+        command_buffer.wait_until_completed();
+
+        let ptr = sum.contents() as *mut u32;
+        unsafe {
+            assert_eq!(465, *ptr);
+        }
+    });
+}