summaryrefslogtreecommitdiffstats
path: root/third_party/rust/metal/examples/compute
diff options
context:
space:
mode:
Diffstat (limited to 'third_party/rust/metal/examples/compute')
-rw-r--r--third_party/rust/metal/examples/compute/compute-argument-buffer.metal14
-rw-r--r--third_party/rust/metal/examples/compute/compute-argument-buffer.rs95
-rw-r--r--third_party/rust/metal/examples/compute/embedded-lib.rs24
-rw-r--r--third_party/rust/metal/examples/compute/main.rs91
-rw-r--r--third_party/rust/metal/examples/compute/shaders.metal10
-rw-r--r--third_party/rust/metal/examples/compute/shaders.metallibbin0 -> 3209 bytes
6 files changed, 234 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);
+}
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);
+ }
+ });
+}
diff --git a/third_party/rust/metal/examples/compute/embedded-lib.rs b/third_party/rust/metal/examples/compute/embedded-lib.rs
new file mode 100644
index 0000000000..0fd193abe3
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/embedded-lib.rs
@@ -0,0 +1,24 @@
+// 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;
+
+fn main() {
+ let library_data = include_bytes!("shaders.metallib");
+
+ autoreleasepool(|| {
+ let device = Device::system_default().expect("no device found");
+
+ let library = device.new_library_with_data(&library_data[..]).unwrap();
+ let kernel = library.get_function("sum", None).unwrap();
+
+ println!("Function name: {}", kernel.name());
+ println!("Function type: {:?}", kernel.function_type());
+ println!("OK");
+ });
+}
diff --git a/third_party/rust/metal/examples/compute/main.rs b/third_party/rust/metal/examples/compute/main.rs
new file mode 100644
index 0000000000..6497c790ae
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/main.rs
@@ -0,0 +1,91 @@
+// 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;
+
+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();
+
+ command_buffer.set_label("label");
+ let block = block::ConcreteBlock::new(move |buffer: &metal::CommandBufferRef| {
+ println!("{}", buffer.label());
+ })
+ .copy();
+
+ command_buffer.add_completed_handler(&block);
+
+ let encoder = command_buffer.new_compute_command_encoder();
+ let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR"))
+ .join("examples/compute/shaders.metallib");
+
+ let library = device.new_library_with_file(library_path).unwrap();
+ let kernel = library.get_function("sum", None).unwrap();
+
+ 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(&buffer), 0);
+ encoder.set_buffer(1, Some(&sum), 0);
+
+ 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);
+ }
+ });
+}
diff --git a/third_party/rust/metal/examples/compute/shaders.metal b/third_party/rust/metal/examples/compute/shaders.metal
new file mode 100644
index 0000000000..51363a1d36
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/shaders.metal
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+kernel void sum(device uint *data [[ buffer(0) ]],
+ volatile device atomic_uint *sum [[ buffer(1) ]],
+ uint gid [[ thread_position_in_grid ]])
+{
+ atomic_fetch_add_explicit(sum, data[gid], memory_order_relaxed);
+}
diff --git a/third_party/rust/metal/examples/compute/shaders.metallib b/third_party/rust/metal/examples/compute/shaders.metallib
new file mode 100644
index 0000000000..af7cb17240
--- /dev/null
+++ b/third_party/rust/metal/examples/compute/shaders.metallib
Binary files differ
generated by cgit v1.2.3 (git 2.39.1) at 2025-01-11 10:27:01 +0000