diff options
Diffstat (limited to 'third_party/rust/metal/examples')
45 files changed, 4111 insertions, 0 deletions
diff --git a/third_party/rust/metal/examples/argument-buffer/main.rs b/third_party/rust/metal/examples/argument-buffer/main.rs new file mode 100644 index 0000000000..23e88990d9 --- /dev/null +++ b/third_party/rust/metal/examples/argument-buffer/main.rs @@ -0,0 +1,88 @@ +// 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() { + autoreleasepool(|| { + let device = Device::system_default().expect("no device found"); + + /* + + // Build encoder for the following MSL argument buffer: + struct ArgumentBuffer { + texture2d<float> texture [[id(0)]]; + sampler sampler [[id(1)]]; + array<device float *, 2> buffers [[id(2)]]; + } + + */ + + let desc1 = ArgumentDescriptor::new(); + desc1.set_index(0); + desc1.set_data_type(MTLDataType::Texture); + desc1.set_texture_type(MTLTextureType::D2); + + let desc2 = ArgumentDescriptor::new(); + desc2.set_index(1); + desc2.set_data_type(MTLDataType::Sampler); + + let desc3 = ArgumentDescriptor::new(); + desc3.set_index(2); + desc3.set_data_type(MTLDataType::Pointer); + desc3.set_array_length(2); + + let encoder = device.new_argument_encoder(Array::from_slice(&[desc1, desc2, desc3])); + println!("Encoder: {:?}", encoder); + + let argument_buffer = + device.new_buffer(encoder.encoded_length(), MTLResourceOptions::empty()); + encoder.set_argument_buffer(&argument_buffer, 0); + + let sampler = { + let descriptor = SamplerDescriptor::new(); + descriptor.set_support_argument_buffers(true); + device.new_sampler(&descriptor) + }; + println!("{:?}", sampler); + + let buffer1 = device.new_buffer(1024, MTLResourceOptions::empty()); + println!("Buffer1: {:?}", buffer1); + let buffer2 = device.new_buffer(1024, MTLResourceOptions::empty()); + println!("Buffer2: {:?}", buffer2); + + encoder.set_sampler_state(1, &sampler); + encoder.set_buffer(2, &buffer1, 0); + encoder.set_buffer(3, &buffer2, 0); + + // How to use argument buffer with render encoder. + + let queue = device.new_command_queue(); + let command_buffer = queue.new_command_buffer(); + + let render_pass_descriptor = RenderPassDescriptor::new(); + let encoder = command_buffer.new_render_command_encoder(render_pass_descriptor); + + // This method makes the array of resources resident for the selected stages of the render pass. + // Call this method before issuing any draw calls that may access the array of resources. + encoder.use_resources( + &[&buffer1, &buffer2], + MTLResourceUsage::Read, + MTLRenderStages::Vertex, + ); + // Bind argument buffer to vertex stage. + encoder.set_vertex_buffer(0, Some(&argument_buffer), 0); + + // Render pass here... + + encoder.end_encoding(); + println!("Encoder: {:?}", encoder); + + command_buffer.commit(); + }); +} diff --git a/third_party/rust/metal/examples/bind/main.rs b/third_party/rust/metal/examples/bind/main.rs new file mode 100644 index 0000000000..811b1c5a17 --- /dev/null +++ b/third_party/rust/metal/examples/bind/main.rs @@ -0,0 +1,34 @@ +// Copyright 2018 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() { + autoreleasepool(|| { + let device = Device::system_default().expect("no device found"); + + let buffer = device.new_buffer(4, MTLResourceOptions::empty()); + let sampler = { + let descriptor = SamplerDescriptor::new(); + device.new_sampler(&descriptor) + }; + + let queue = device.new_command_queue(); + let cmd_buf = queue.new_command_buffer(); + + let encoder = cmd_buf.new_compute_command_encoder(); + + encoder.set_buffers(2, &[Some(&buffer), None], &[4, 0]); + encoder.set_sampler_states(1, &[Some(&sampler), None]); + + encoder.end_encoding(); + cmd_buf.commit(); + + println!("Everything is bound"); + }); +} diff --git a/third_party/rust/metal/examples/bindless/main.rs b/third_party/rust/metal/examples/bindless/main.rs new file mode 100644 index 0000000000..09c3a59ab9 --- /dev/null +++ b/third_party/rust/metal/examples/bindless/main.rs @@ -0,0 +1,149 @@ +// 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; + +const BINDLESS_TEXTURE_COUNT: NSUInteger = 100_000; // ~25Mb + +/// This example demonstrates: +/// - How to create a heap +/// - How to allocate textures from heap. +/// - How to create bindless resources via Metal's argument buffers. +/// - How to bind argument buffer to render encoder +fn main() { + autoreleasepool(|| { + let device = Device::system_default().expect("no device found"); + + /* + + MSL + + struct Textures { + texture2d<float> texture; + }; + struct BindlessTextures { + device Textures *textures; + }; + + */ + + // Tier 2 argument buffers are supported by macOS devices with a discrete GPU and by the A13 GPU. + // The maximum per-app resources available at any given time are: + // - 500,000 buffers or textures + // - 2048 unique samplers + let tier = device.argument_buffers_support(); + println!("Argument buffer support: {:?}", tier); + assert_eq!(MTLArgumentBuffersTier::Tier2, tier); + + let texture_descriptor = TextureDescriptor::new(); + texture_descriptor.set_width(1); + texture_descriptor.set_height(1); + texture_descriptor.set_depth(1); + texture_descriptor.set_texture_type(MTLTextureType::D2); + texture_descriptor.set_pixel_format(MTLPixelFormat::R8Uint); + texture_descriptor.set_storage_mode(MTLStorageMode::Private); // GPU only. + println!("Texture descriptor: {:?}", texture_descriptor); + + // Determine the size required for the heap for the given descriptor + let size_and_align = device.heap_texture_size_and_align(&texture_descriptor); + + // Align the size so that more resources will fit in the heap after this texture + // See https://developer.apple.com/documentation/metal/buffers/using_argument_buffers_with_resource_heaps + let texture_size = + (size_and_align.size & (size_and_align.align - 1)) + size_and_align.align; + let heap_size = texture_size * BINDLESS_TEXTURE_COUNT; + + let heap_descriptor = HeapDescriptor::new(); + heap_descriptor.set_storage_mode(texture_descriptor.storage_mode()); // Must be compatible + heap_descriptor.set_size(heap_size); + println!("Heap descriptor: {:?}", heap_descriptor); + + let heap = device.new_heap(&heap_descriptor); + println!("Heap: {:?}", heap); + + // Allocate textures from heap + let textures = (0..BINDLESS_TEXTURE_COUNT) + .map(|i| { + heap.new_texture(&texture_descriptor) + .expect(&format!("Failed to allocate texture {}", i)) + }) + .collect::<Vec<_>>(); + + // Crate argument encoder that knows how to encode single texture + let descriptor = ArgumentDescriptor::new(); + descriptor.set_index(0); + descriptor.set_data_type(MTLDataType::Texture); + descriptor.set_texture_type(MTLTextureType::D2); + descriptor.set_access(MTLArgumentAccess::ReadOnly); + println!("Argument descriptor: {:?}", descriptor); + + let encoder = device.new_argument_encoder(Array::from_slice(&[descriptor])); + println!("Encoder: {:?}", encoder); + + // Determinate argument buffer size to allocate. + // Size needed to encode one texture * total number of bindless textures. + let argument_buffer_size = encoder.encoded_length() * BINDLESS_TEXTURE_COUNT; + let argument_buffer = device.new_buffer(argument_buffer_size, MTLResourceOptions::empty()); + + // Encode textures to the argument buffer. + textures.iter().enumerate().for_each(|(index, texture)| { + // Offset encoder to a proper texture slot + let offset = index as NSUInteger * encoder.encoded_length(); + encoder.set_argument_buffer(&argument_buffer, offset); + encoder.set_texture(0, texture); + }); + + // How to use bindless argument buffer when drawing + + let queue = device.new_command_queue(); + let command_buffer = queue.new_command_buffer(); + + let render_pass_descriptor = RenderPassDescriptor::new(); + let encoder = command_buffer.new_render_command_encoder(render_pass_descriptor); + + // Bind argument buffer. + encoder.set_fragment_buffer(0, Some(&argument_buffer), 0); + // Make sure all textures are available to the pass. + encoder.use_heap_at(&heap, MTLRenderStages::Fragment); + + // Bind material buffer at index 1 + // Draw + + /* + + // Now instead of binding individual textures each draw call, + // you can just bind material information instead: + + MSL + + struct Material { + int diffuse_texture_index; + int normal_texture_index; + // ... + } + + fragment float4 pixel( + VertexOut v [[stage_in]], + constant const BindlessTextures * textures [[buffer(0)]], + constant Material * material [[buffer(1)]] + ) { + if (material->base_color_texture_index != -1) { + textures[material->diffuse_texture_index].texture.sampler(...) + } + if (material->normal_texture_index != -1) { + ... + } + ... + } + + */ + + encoder.end_encoding(); + command_buffer.commit(); + }); +} diff --git a/third_party/rust/metal/examples/caps/main.rs b/third_party/rust/metal/examples/caps/main.rs new file mode 100644 index 0000000000..ae8fca4f0a --- /dev/null +++ b/third_party/rust/metal/examples/caps/main.rs @@ -0,0 +1,33 @@ +// 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::*; + +fn main() { + let device = Device::system_default().expect("no device found"); + + #[cfg(feature = "private")] + { + println!("Vendor: {:?}", unsafe { device.vendor() }); + println!("Family: {:?}", unsafe { device.family_name() }); + } + println!( + "Max threads per threadgroup: {:?}", + device.max_threads_per_threadgroup() + ); + #[cfg(target_os = "macos")] + { + println!("Integrated GPU: {:?}", device.is_low_power()); + println!("Headless: {:?}", device.is_headless()); + println!("D24S8: {:?}", device.d24_s8_supported()); + } + println!("maxBufferLength: {} Mb", device.max_buffer_length() >> 20); + println!( + "Indirect argument buffer: {:?}", + device.argument_buffers_support() + ); +} diff --git a/third_party/rust/metal/examples/circle/README.md b/third_party/rust/metal/examples/circle/README.md new file mode 100644 index 0000000000..f51853ac38 --- /dev/null +++ b/third_party/rust/metal/examples/circle/README.md @@ -0,0 +1,11 @@ +## circle + +Renders a circle in a window. As metal primitive types are only limited to point, line and triangle shape, this example shows how we can form complex structures out of primitive types. + + + +## To Run + +``` +cargo run --example circle +``` diff --git a/third_party/rust/metal/examples/circle/main.rs b/third_party/rust/metal/examples/circle/main.rs new file mode 100644 index 0000000000..ed17a73402 --- /dev/null +++ b/third_party/rust/metal/examples/circle/main.rs @@ -0,0 +1,377 @@ +use metal::*; + +use winit::{ + event::{Event, WindowEvent}, + event_loop::{ControlFlow, EventLoop}, + platform::macos::WindowExtMacOS, +}; + +use cocoa::{appkit::NSView, base::id as cocoa_id}; +use core_graphics_types::geometry::CGSize; + +use objc::{rc::autoreleasepool, runtime::YES}; + +use std::mem; + +// Declare the data structures needed to carry vertex layout to +// metal shading language(MSL) program. Use #[repr(C)], to make +// the data structure compatible with C++ type data structure +// for vertex defined in MSL program as MSL program is broadly +// based on C++ +#[repr(C)] +#[derive(Debug)] +pub struct position(cty::c_float, cty::c_float); +#[repr(C)] +#[derive(Debug)] +pub struct color(cty::c_float, cty::c_float, cty::c_float); +#[repr(C)] +#[derive(Debug)] +pub struct AAPLVertex { + p: position, + c: color, +} + +fn main() { + // Create a window for viewing the content + let event_loop = EventLoop::new(); + let size = winit::dpi::LogicalSize::new(800, 600); + + let window = winit::window::WindowBuilder::new() + .with_inner_size(size) + .with_title("Metal".to_string()) + .build(&event_loop) + .unwrap(); + + // Set up the GPU device found in the system + let device = Device::system_default().expect("no device found"); + println!("Your device is: {}", device.name(),); + + // Scaffold required to sample the GPU and CPU timestamps + let mut cpu_start = 0; + let mut gpu_start = 0; + device.sample_timestamps(&mut cpu_start, &mut gpu_start); + let counter_sample_buffer = create_counter_sample_buffer(&device); + let destination_buffer = device.new_buffer( + (std::mem::size_of::<u64>() * 4 as usize) as u64, + MTLResourceOptions::StorageModeShared, + ); + let counter_sampling_point = MTLCounterSamplingPoint::AtStageBoundary; + assert!(device.supports_counter_sampling(counter_sampling_point)); + + let binary_archive_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/circle/binary_archive.metallib"); + + let binary_archive_url = + URL::new_with_string(&format!("file://{}", binary_archive_path.display())); + + let binary_archive_descriptor = BinaryArchiveDescriptor::new(); + if binary_archive_path.exists() { + binary_archive_descriptor.set_url(&binary_archive_url); + } + + // Set up a binary archive to cache compiled shaders. + let binary_archive = device + .new_binary_archive_with_descriptor(&binary_archive_descriptor) + .unwrap(); + + let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/circle/shaders.metallib"); + + // Use the metallib file generated out of .metal shader file + let library = device.new_library_with_file(library_path).unwrap(); + + // The render pipeline generated from the vertex and fragment shaders in the .metal shader file. + let pipeline_state = prepare_pipeline_state(&device, &library, &binary_archive); + + // Serialize the binary archive to disk. + binary_archive + .serialize_to_url(&binary_archive_url) + .unwrap(); + + // Set the command queue used to pass commands to the device. + let command_queue = device.new_command_queue(); + + // Currently, MetalLayer is the only interface that provide + // layers to carry drawable texture from GPU rendaring through metal + // library to viewable windows. + let layer = MetalLayer::new(); + layer.set_device(&device); + layer.set_pixel_format(MTLPixelFormat::BGRA8Unorm); + layer.set_presents_with_transaction(false); + + unsafe { + let view = window.ns_view() as cocoa_id; + view.setWantsLayer(YES); + view.setLayer(mem::transmute(layer.as_ref())); + } + + let draw_size = window.inner_size(); + layer.set_drawable_size(CGSize::new(draw_size.width as f64, draw_size.height as f64)); + + let vbuf = { + let vertex_data = create_vertex_points_for_circle(); + let vertex_data = vertex_data.as_slice(); + + device.new_buffer_with_data( + vertex_data.as_ptr() as *const _, + (vertex_data.len() * mem::size_of::<AAPLVertex>()) as u64, + MTLResourceOptions::CPUCacheModeDefaultCache | MTLResourceOptions::StorageModeManaged, + ) + }; + + event_loop.run(move |event, _, control_flow| { + autoreleasepool(|| { + // ControlFlow::Wait pauses the event loop if no events are available to process. + // This is ideal for non-game applications that only update in response to user + // input, and uses significantly less power/CPU time than ControlFlow::Poll. + *control_flow = ControlFlow::Wait; + + match event { + Event::WindowEvent { + event: WindowEvent::CloseRequested, + .. + } => { + println!("The close button was pressed; stopping"); + *control_flow = ControlFlow::Exit + } + Event::MainEventsCleared => { + // Queue a RedrawRequested event. + window.request_redraw(); + } + Event::RedrawRequested(_) => { + // It's preferrable to render in this event rather than in MainEventsCleared, since + // rendering in here allows the program to gracefully handle redraws requested + // by the OS. + let drawable = match layer.next_drawable() { + Some(drawable) => drawable, + None => return, + }; + + // Create a new command buffer for each render pass to the current drawable + let command_buffer = command_queue.new_command_buffer(); + + // Obtain a renderPassDescriptor generated from the view's drawable textures. + let render_pass_descriptor = RenderPassDescriptor::new(); + handle_render_pass_color_attachment( + &render_pass_descriptor, + drawable.texture(), + ); + handle_render_pass_sample_buffer_attachment( + &render_pass_descriptor, + &counter_sample_buffer, + ); + + // Create a render command encoder. + let encoder = + command_buffer.new_render_command_encoder(&render_pass_descriptor); + encoder.set_render_pipeline_state(&pipeline_state); + // Pass in the parameter data. + encoder.set_vertex_buffer(0, Some(&vbuf), 0); + // Draw the triangles which will eventually form the circle. + encoder.draw_primitives(MTLPrimitiveType::TriangleStrip, 0, 1080); + encoder.end_encoding(); + + resolve_samples_into_buffer( + &command_buffer, + &counter_sample_buffer, + &destination_buffer, + ); + + // Schedule a present once the framebuffer is complete using the current drawable. + command_buffer.present_drawable(&drawable); + + // Finalize rendering here & push the command buffer to the GPU. + command_buffer.commit(); + command_buffer.wait_until_completed(); + + let mut cpu_end = 0; + let mut gpu_end = 0; + device.sample_timestamps(&mut cpu_end, &mut gpu_end); + handle_timestamps(&destination_buffer, cpu_start, cpu_end, gpu_start, gpu_end); + } + _ => (), + } + }); + }); +} + +// If we want to draw a circle, we need to draw it out of the three primitive +// types available with metal framework. Triangle is used in this case to form +// the circle. If we consider a circle to be total of 360 degree at center, we +// can form small triangle with one point at origin and two points at the +// perimeter of the circle for each degree. Eventually, if we can take enough +// triangle virtices for total of 360 degree, the triangles together will +// form a circle. This function captures the triangle vertices for each degree +// and push the co-ordinates of the vertices to a rust vector +fn create_vertex_points_for_circle() -> Vec<AAPLVertex> { + let mut v: Vec<AAPLVertex> = Vec::new(); + let origin_x: f32 = 0.0; + let origin_y: f32 = 0.0; + + // Size of the circle + let circle_size = 0.8f32; + + for i in 0..720 { + let y = i as f32; + // Get the X co-ordinate of each point on the perimeter of circle + let position_x: f32 = y.to_radians().cos() * 100.0; + let position_x: f32 = position_x.trunc() / 100.0; + // Set the size of the circle + let position_x: f32 = position_x * circle_size; + // Get the Y co-ordinate of each point on the perimeter of circle + let position_y: f32 = y.to_radians().sin() * 100.0; + let position_y: f32 = position_y.trunc() / 100.0; + // Set the size of the circle + let position_y: f32 = position_y * circle_size; + + v.push(AAPLVertex { + p: position(position_x, position_y), + c: color(0.7, 0.3, 0.5), + }); + + if (i + 1) % 2 == 0 { + // For each two points on perimeter, push one point of origin + v.push(AAPLVertex { + p: position(origin_x, origin_y), + c: color(0.2, 0.7, 0.4), + }); + } + } + + v +} + +fn handle_render_pass_sample_buffer_attachment( + descriptor: &RenderPassDescriptorRef, + counter_sample_buffer: &CounterSampleBufferRef, +) { + let sample_buffer_attachment_descriptor = + descriptor.sample_buffer_attachments().object_at(0).unwrap(); + sample_buffer_attachment_descriptor.set_sample_buffer(&counter_sample_buffer); + sample_buffer_attachment_descriptor.set_start_of_vertex_sample_index(0 as NSUInteger); + sample_buffer_attachment_descriptor.set_end_of_vertex_sample_index(1 as NSUInteger); + sample_buffer_attachment_descriptor.set_start_of_fragment_sample_index(2 as NSUInteger); + sample_buffer_attachment_descriptor.set_end_of_fragment_sample_index(3 as NSUInteger); +} + +fn handle_render_pass_color_attachment(descriptor: &RenderPassDescriptorRef, texture: &TextureRef) { + let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + + color_attachment.set_texture(Some(texture)); + color_attachment.set_load_action(MTLLoadAction::Clear); + // Setting a background color + color_attachment.set_clear_color(MTLClearColor::new(0.5, 0.5, 0.8, 1.0)); + color_attachment.set_store_action(MTLStoreAction::Store); +} + +fn prepare_pipeline_state( + device: &Device, + library: &Library, + binary_archive: &BinaryArchive, +) -> RenderPipelineState { + let vert = library.get_function("vs", None).unwrap(); + let frag = library.get_function("ps", None).unwrap(); + + let pipeline_state_descriptor = RenderPipelineDescriptor::new(); + pipeline_state_descriptor.set_vertex_function(Some(&vert)); + pipeline_state_descriptor.set_fragment_function(Some(&frag)); + pipeline_state_descriptor + .color_attachments() + .object_at(0) + .unwrap() + .set_pixel_format(MTLPixelFormat::BGRA8Unorm); + // Set the binary archives to search for a cached pipeline in. + pipeline_state_descriptor.set_binary_archives(&[binary_archive]); + + // Add the pipeline descriptor to the binary archive cache. + binary_archive + .add_render_pipeline_functions_with_descriptor(&pipeline_state_descriptor) + .unwrap(); + + device + .new_render_pipeline_state(&pipeline_state_descriptor) + .unwrap() +} + +fn resolve_samples_into_buffer( + command_buffer: &CommandBufferRef, + counter_sample_buffer: &CounterSampleBufferRef, + destination_buffer: &BufferRef, +) { + let blit_encoder = command_buffer.new_blit_command_encoder(); + blit_encoder.resolve_counters( + &counter_sample_buffer, + crate::NSRange::new(0_u64, 4), + &destination_buffer, + 0_u64, + ); + blit_encoder.end_encoding(); +} + +fn handle_timestamps( + resolved_sample_buffer: &BufferRef, + cpu_start: u64, + cpu_end: u64, + gpu_start: u64, + gpu_end: u64, +) { + let samples = unsafe { + std::slice::from_raw_parts(resolved_sample_buffer.contents() as *const u64, 4 as usize) + }; + let vertex_pass_start = samples[0]; + let vertex_pass_end = samples[1]; + let fragment_pass_start = samples[2]; + let fragment_pass_end = samples[3]; + + let cpu_time_span = cpu_end - cpu_start; + let gpu_time_span = gpu_end - gpu_start; + + let vertex_micros = microseconds_between_begin( + vertex_pass_start, + vertex_pass_end, + gpu_time_span, + cpu_time_span, + ); + let fragment_micros = microseconds_between_begin( + fragment_pass_start, + fragment_pass_end, + gpu_time_span, + cpu_time_span, + ); + + println!("Vertex pass duration: {:.2} µs", vertex_micros); + println!("Fragment pass duration: {:.2} µs\n", fragment_micros); +} + +fn create_counter_sample_buffer(device: &Device) -> CounterSampleBuffer { + let counter_sample_buffer_desc = metal::CounterSampleBufferDescriptor::new(); + counter_sample_buffer_desc.set_storage_mode(metal::MTLStorageMode::Shared); + counter_sample_buffer_desc.set_sample_count(4_u64); + counter_sample_buffer_desc.set_counter_set(&fetch_timestamp_counter_set(device)); + + device + .new_counter_sample_buffer_with_descriptor(&counter_sample_buffer_desc) + .unwrap() +} + +fn fetch_timestamp_counter_set(device: &Device) -> metal::CounterSet { + let counter_sets = device.counter_sets(); + let mut timestamp_counter = None; + for cs in counter_sets.iter() { + if cs.name() == "timestamp" { + timestamp_counter = Some(cs); + break; + } + } + timestamp_counter + .expect("No timestamp counter found") + .clone() +} + +/// <https://developer.apple.com/documentation/metal/gpu_counters_and_counter_sample_buffers/converting_gpu_timestamps_into_cpu_time> +fn microseconds_between_begin(begin: u64, end: u64, gpu_time_span: u64, cpu_time_span: u64) -> f64 { + let time_span = (end as f64) - (begin as f64); + let nanoseconds = time_span / (gpu_time_span as f64) * (cpu_time_span as f64); + let microseconds = nanoseconds / 1000.0; + return microseconds; +} diff --git a/third_party/rust/metal/examples/circle/screenshot.png b/third_party/rust/metal/examples/circle/screenshot.png Binary files differnew file mode 100644 index 0000000000..38f86e733d --- /dev/null +++ b/third_party/rust/metal/examples/circle/screenshot.png diff --git a/third_party/rust/metal/examples/circle/shaders.metal b/third_party/rust/metal/examples/circle/shaders.metal new file mode 100644 index 0000000000..037af8a233 --- /dev/null +++ b/third_party/rust/metal/examples/circle/shaders.metal @@ -0,0 +1,39 @@ +#include <metal_stdlib> + +#include <simd/simd.h> + +using namespace metal; + +typedef struct { + float x; + float y; +}position; + +typedef struct { + float r; + float g; + float b; +}color; + +typedef struct { + position p; + color c; +}AAPLVertex; + +struct ColorInOut { + float4 position[[position]]; + float4 color; +}; + +vertex ColorInOut vs(constant AAPLVertex * vertex_array[[buffer(0)]], unsigned int vid[[vertex_id]]) { + ColorInOut out; + + out.position = float4(float2(vertex_array[vid].p.x, vertex_array[vid].p.y), 0.0, 1.0); + out.color = float4(float3(vertex_array[vid].c.r, vertex_array[vid].c.g, vertex_array[vid].c.b), 1.0); + + return out; +} + +fragment float4 ps(ColorInOut in [[stage_in]]) { + return in.color; +} diff --git a/third_party/rust/metal/examples/circle/shaders.metallib b/third_party/rust/metal/examples/circle/shaders.metallib Binary files differnew file mode 100644 index 0000000000..cbb9bc5e5a --- /dev/null +++ b/third_party/rust/metal/examples/circle/shaders.metallib 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..b654a666f1 --- /dev/null +++ b/third_party/rust/metal/examples/compute/main.rs @@ -0,0 +1,194 @@ +use metal::*; +use objc::rc::autoreleasepool; +use std::path::PathBuf; + +const NUM_SAMPLES: u64 = 2; + +fn main() { + let num_elements = std::env::args() + .nth(1) + .map(|s| s.parse::<u32>().unwrap()) + .unwrap_or(64 * 64); + + autoreleasepool(|| { + let device = Device::system_default().expect("No device found"); + let mut cpu_start = 0; + let mut gpu_start = 0; + device.sample_timestamps(&mut cpu_start, &mut gpu_start); + + let counter_sample_buffer = create_counter_sample_buffer(&device); + let destination_buffer = device.new_buffer( + (std::mem::size_of::<u64>() * NUM_SAMPLES as usize) as u64, + MTLResourceOptions::StorageModeShared, + ); + + let counter_sampling_point = MTLCounterSamplingPoint::AtStageBoundary; + assert!(device.supports_counter_sampling(counter_sampling_point)); + + let command_queue = device.new_command_queue(); + let command_buffer = command_queue.new_command_buffer(); + + let compute_pass_descriptor = ComputePassDescriptor::new(); + handle_compute_pass_sample_buffer_attachment( + compute_pass_descriptor, + &counter_sample_buffer, + ); + let encoder = + command_buffer.compute_command_encoder_with_descriptor(compute_pass_descriptor); + + let pipeline_state = create_pipeline_state(&device); + encoder.set_compute_pipeline_state(&pipeline_state); + + let (buffer, sum) = create_input_and_output_buffers(&device, num_elements); + encoder.set_buffer(0, Some(&buffer), 0); + encoder.set_buffer(1, Some(&sum), 0); + + let num_threads = pipeline_state.thread_execution_width(); + + let thread_group_count = MTLSize { + width: ((num_elements as NSUInteger + num_threads) / num_threads), + height: 1, + depth: 1, + }; + + let thread_group_size = MTLSize { + width: num_threads, + height: 1, + depth: 1, + }; + + encoder.dispatch_thread_groups(thread_group_count, thread_group_size); + encoder.end_encoding(); + + resolve_samples_into_buffer(command_buffer, &counter_sample_buffer, &destination_buffer); + + command_buffer.commit(); + command_buffer.wait_until_completed(); + let mut cpu_end = 0; + let mut gpu_end = 0; + device.sample_timestamps(&mut cpu_end, &mut gpu_end); + + let ptr = sum.contents() as *mut u32; + println!("Compute shader sum: {}", unsafe { *ptr }); + + unsafe { + assert_eq!(num_elements, *ptr); + } + + handle_timestamps(&destination_buffer, cpu_start, cpu_end, gpu_start, gpu_end); + }); +} + +fn create_pipeline_state(device: &Device) -> ComputePipelineState { + let library_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)); + + device + .new_compute_pipeline_state_with_function( + pipeline_state_descriptor.compute_function().unwrap(), + ) + .unwrap() +} + +fn handle_compute_pass_sample_buffer_attachment( + compute_pass_descriptor: &ComputePassDescriptorRef, + counter_sample_buffer: &CounterSampleBufferRef, +) { + let sample_buffer_attachment_descriptor = compute_pass_descriptor + .sample_buffer_attachments() + .object_at(0) + .unwrap(); + + sample_buffer_attachment_descriptor.set_sample_buffer(counter_sample_buffer); + sample_buffer_attachment_descriptor.set_start_of_encoder_sample_index(0); + sample_buffer_attachment_descriptor.set_end_of_encoder_sample_index(1); +} + +fn resolve_samples_into_buffer( + command_buffer: &CommandBufferRef, + counter_sample_buffer: &CounterSampleBufferRef, + destination_buffer: &BufferRef, +) { + let blit_encoder = command_buffer.new_blit_command_encoder(); + blit_encoder.resolve_counters( + counter_sample_buffer, + crate::NSRange::new(0_u64, NUM_SAMPLES), + destination_buffer, + 0_u64, + ); + blit_encoder.end_encoding(); +} + +fn handle_timestamps( + resolved_sample_buffer: &BufferRef, + cpu_start: u64, + cpu_end: u64, + gpu_start: u64, + gpu_end: u64, +) { + let samples = unsafe { + std::slice::from_raw_parts( + resolved_sample_buffer.contents() as *const u64, + NUM_SAMPLES as usize, + ) + }; + let pass_start = samples[0]; + let pass_end = samples[1]; + + let cpu_time_span = cpu_end - cpu_start; + let gpu_time_span = gpu_end - gpu_start; + + let micros = microseconds_between_begin(pass_start, pass_end, gpu_time_span, cpu_time_span); + println!("Compute pass duration: {} µs", micros); +} + +fn create_counter_sample_buffer(device: &Device) -> CounterSampleBuffer { + let counter_sample_buffer_desc = metal::CounterSampleBufferDescriptor::new(); + counter_sample_buffer_desc.set_storage_mode(metal::MTLStorageMode::Shared); + counter_sample_buffer_desc.set_sample_count(NUM_SAMPLES); + let counter_sets = device.counter_sets(); + + let timestamp_counter = counter_sets.iter().find(|cs| cs.name() == "timestamp"); + + counter_sample_buffer_desc + .set_counter_set(timestamp_counter.expect("No timestamp counter found")); + + device + .new_counter_sample_buffer_with_descriptor(&counter_sample_buffer_desc) + .unwrap() +} + +fn create_input_and_output_buffers( + device: &Device, + num_elements: u32, +) -> (metal::Buffer, metal::Buffer) { + let data = vec![1u32; num_elements as usize]; + + let buffer = device.new_buffer_with_data( + unsafe { std::mem::transmute(data.as_ptr()) }, + (data.len() * std::mem::size_of::<u32>()) as u64, + MTLResourceOptions::CPUCacheModeDefaultCache, + ); + + let sum = { + let data = [0u32]; + device.new_buffer_with_data( + unsafe { std::mem::transmute(data.as_ptr()) }, + (data.len() * std::mem::size_of::<u32>()) as u64, + MTLResourceOptions::CPUCacheModeDefaultCache, + ) + }; + (buffer, sum) +} + +/// <https://developer.apple.com/documentation/metal/gpu_counters_and_counter_sample_buffers/converting_gpu_timestamps_into_cpu_time> +fn microseconds_between_begin(begin: u64, end: u64, gpu_time_span: u64, cpu_time_span: u64) -> f64 { + let time_span = (end as f64) - (begin as f64); + let nanoseconds = time_span / (gpu_time_span as f64) * (cpu_time_span as f64); + nanoseconds / 1000.0 +} 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 Binary files differnew file mode 100644 index 0000000000..af7cb17240 --- /dev/null +++ b/third_party/rust/metal/examples/compute/shaders.metallib diff --git a/third_party/rust/metal/examples/events/main.rs b/third_party/rust/metal/examples/events/main.rs new file mode 100644 index 0000000000..9e4fe0e820 --- /dev/null +++ b/third_party/rust/metal/examples/events/main.rs @@ -0,0 +1,50 @@ +// Copyright 2020 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 dispatch::{Queue, QueueAttribute}; +use metal::*; + +/// This example replicates `Synchronizing Events Between a GPU and the CPU` article. +/// See https://developer.apple.com/documentation/metal/synchronization/synchronizing_events_between_a_gpu_and_the_cpu +fn main() { + let device = Device::system_default().expect("No device found"); + + let command_queue = device.new_command_queue(); + let command_buffer = command_queue.new_command_buffer(); + + // Shareable event + let shared_event = device.new_shared_event(); + + // Shareable event listener + let my_queue = Queue::create( + "com.example.apple-samplecode.MyQueue", + QueueAttribute::Serial, + ); + + // Enable `dispatch` feature to use dispatch queues, + // otherwise unsafe `from_queue_handle` is available for use with native APIs. + let shared_event_listener = SharedEventListener::from_queue(&my_queue); + + // Register CPU work + let notify_block = block::ConcreteBlock::new(move |evt: &SharedEventRef, val: u64| { + println!("Got notification from GPU: {}", val); + evt.set_signaled_value(3); + }); + + shared_event.notify(&shared_event_listener, 2, notify_block.copy()); + + // Encode GPU work + command_buffer.encode_signal_event(&shared_event, 1); + command_buffer.encode_signal_event(&shared_event, 2); + command_buffer.encode_wait_for_event(&shared_event, 3); + + command_buffer.commit(); + + command_buffer.wait_until_completed(); + + println!("Done"); +} diff --git a/third_party/rust/metal/examples/fence/main.rs b/third_party/rust/metal/examples/fence/main.rs new file mode 100644 index 0000000000..53515d39f5 --- /dev/null +++ b/third_party/rust/metal/examples/fence/main.rs @@ -0,0 +1,30 @@ +// Copyright 2020 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::*; + +fn main() { + let device = Device::system_default().expect("No device found"); + + let command_queue = device.new_command_queue(); + let command_buffer = command_queue.new_command_buffer(); + + let fence = device.new_fence(); + + let blit_encoder = command_buffer.new_blit_command_encoder(); + blit_encoder.update_fence(&fence); + blit_encoder.end_encoding(); + + let compute_encoder = command_buffer.new_compute_command_encoder(); + compute_encoder.wait_for_fence(&fence); + compute_encoder.end_encoding(); + + command_buffer.commit(); + command_buffer.wait_until_completed(); + + println!("Done"); +} diff --git a/third_party/rust/metal/examples/headless-render/README.md b/third_party/rust/metal/examples/headless-render/README.md new file mode 100644 index 0000000000..6bc434b44a --- /dev/null +++ b/third_party/rust/metal/examples/headless-render/README.md @@ -0,0 +1,11 @@ +## headless-render + +Renders the triangle from the [window example](../window) headlessly and then writes it to a PNG file. + + + +## To Run + +``` +cargo run --example headless-render +``` diff --git a/third_party/rust/metal/examples/headless-render/main.rs b/third_party/rust/metal/examples/headless-render/main.rs new file mode 100644 index 0000000000..449170251f --- /dev/null +++ b/third_party/rust/metal/examples/headless-render/main.rs @@ -0,0 +1,159 @@ +use std::mem; +use std::path::PathBuf; + +use std::fs::File; +use std::io::BufWriter; + +use metal::{ + Buffer, Device, DeviceRef, LibraryRef, MTLClearColor, MTLLoadAction, MTLOrigin, MTLPixelFormat, + MTLPrimitiveType, MTLRegion, MTLResourceOptions, MTLSize, MTLStoreAction, RenderPassDescriptor, + RenderPassDescriptorRef, RenderPipelineDescriptor, RenderPipelineState, Texture, + TextureDescriptor, TextureRef, +}; +use png::ColorType; + +const VIEW_WIDTH: u64 = 512; +const VIEW_HEIGHT: u64 = 512; +const TOTAL_BYTES: usize = (VIEW_WIDTH * VIEW_HEIGHT * 4) as usize; + +const VERTEX_SHADER: &'static str = "triangle_vertex"; +const FRAGMENT_SHADER: &'static str = "triangle_fragment"; + +// [2 bytes position, 3 bytes color] * 3 +#[rustfmt::skip] +const VERTEX_ATTRIBS: [f32; 15] = [ + 0.0, 0.5, 1.0, 0.0, 0.0, + -0.5, -0.5, 0.0, 1.0, 0.0, + 0.5, -0.5, 0.0, 0.0, 1.0, +]; + +/// This example shows how to render headlessly by: +/// +/// 1. Rendering a triangle to an MtlDrawable +/// +/// 2. Waiting for the render to complete and the color texture to be synchronized with the CPU +/// by using a blit command encoder +/// +/// 3. Reading the texture bytes from the MtlTexture +/// +/// 4. Saving the texture to a PNG file +fn main() { + let device = Device::system_default().expect("No device found"); + + let texture = create_texture(&device); + + let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/window/shaders.metallib"); + + let library = device.new_library_with_file(library_path).unwrap(); + + let pipeline_state = prepare_pipeline_state(&device, &library); + + let command_queue = device.new_command_queue(); + + let vertex_buffer = create_vertex_buffer(&device); + + let render_pass_descriptor = RenderPassDescriptor::new(); + initialize_color_attachment(&render_pass_descriptor, &texture); + + let command_buffer = command_queue.new_command_buffer(); + let rc_encoder = command_buffer.new_render_command_encoder(&render_pass_descriptor); + rc_encoder.set_render_pipeline_state(&pipeline_state); + rc_encoder.set_vertex_buffer(0, Some(&vertex_buffer), 0); + rc_encoder.draw_primitives(MTLPrimitiveType::Triangle, 0, 3); + rc_encoder.end_encoding(); + + render_pass_descriptor + .color_attachments() + .object_at(0) + .unwrap() + .set_load_action(MTLLoadAction::DontCare); + + let blit_encoder = command_buffer.new_blit_command_encoder(); + blit_encoder.synchronize_resource(&texture); + blit_encoder.end_encoding(); + + command_buffer.commit(); + + command_buffer.wait_until_completed(); + + save_image(&texture); +} + +fn save_image(texture: &TextureRef) { + let mut image = vec![0; TOTAL_BYTES]; + + texture.get_bytes( + image.as_mut_ptr() as *mut std::ffi::c_void, + VIEW_WIDTH * 4, + MTLRegion { + origin: MTLOrigin { x: 0, y: 0, z: 0 }, + size: MTLSize { + width: VIEW_WIDTH, + height: VIEW_HEIGHT, + depth: 1, + }, + }, + 0, + ); + + let out_file = + PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("examples/headless-render/out.png"); + let file = File::create(&out_file).unwrap(); + let ref mut w = BufWriter::new(file); + + let mut encoder = png::Encoder::new(w, VIEW_WIDTH as u32, VIEW_HEIGHT as u32); + encoder.set_color(ColorType::Rgba); + encoder.set_depth(png::BitDepth::Eight); + let mut writer = encoder.write_header().unwrap(); + + writer.write_image_data(&image).unwrap(); + + println!("Image saved to {:?}", out_file); +} + +fn create_texture(device: &Device) -> Texture { + let texture = TextureDescriptor::new(); + texture.set_width(VIEW_WIDTH); + texture.set_height(VIEW_HEIGHT); + texture.set_pixel_format(MTLPixelFormat::RGBA8Unorm); + + device.new_texture(&texture) +} + +fn prepare_pipeline_state(device: &DeviceRef, library: &LibraryRef) -> RenderPipelineState { + let vert = library.get_function(VERTEX_SHADER, None).unwrap(); + let frag = library.get_function(FRAGMENT_SHADER, None).unwrap(); + + let pipeline_state_descriptor = RenderPipelineDescriptor::new(); + + pipeline_state_descriptor.set_vertex_function(Some(&vert)); + pipeline_state_descriptor.set_fragment_function(Some(&frag)); + + pipeline_state_descriptor + .color_attachments() + .object_at(0) + .unwrap() + .set_pixel_format(MTLPixelFormat::RGBA8Unorm); + + device + .new_render_pipeline_state(&pipeline_state_descriptor) + .unwrap() +} + +fn create_vertex_buffer(device: &DeviceRef) -> Buffer { + device.new_buffer_with_data( + VERTEX_ATTRIBS.as_ptr() as *const _, + (VERTEX_ATTRIBS.len() * mem::size_of::<f32>()) as u64, + MTLResourceOptions::CPUCacheModeDefaultCache | MTLResourceOptions::StorageModeManaged, + ) +} + +fn initialize_color_attachment(descriptor: &RenderPassDescriptorRef, texture: &TextureRef) { + let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + + color_attachment.set_texture(Some(texture)); + color_attachment.set_load_action(MTLLoadAction::Clear); + color_attachment.set_clear_color(MTLClearColor::new(0.5, 0.2, 0.2, 1.0)); + color_attachment.set_store_action(MTLStoreAction::Store); +} diff --git a/third_party/rust/metal/examples/headless-render/screenshot.png b/third_party/rust/metal/examples/headless-render/screenshot.png Binary files differnew file mode 100644 index 0000000000..2af9c5895f --- /dev/null +++ b/third_party/rust/metal/examples/headless-render/screenshot.png diff --git a/third_party/rust/metal/examples/library/main.rs b/third_party/rust/metal/examples/library/main.rs new file mode 100644 index 0000000000..7223db89c1 --- /dev/null +++ b/third_party/rust/metal/examples/library/main.rs @@ -0,0 +1,17 @@ +// Copyright 2016 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::*; + +const PROGRAM: &'static str = ""; + +fn main() { + let device = Device::system_default().expect("no device found"); + + let options = CompileOptions::new(); + let _library = device.new_library_with_source(PROGRAM, &options); +} diff --git a/third_party/rust/metal/examples/mesh-shader/main.rs b/third_party/rust/metal/examples/mesh-shader/main.rs new file mode 100644 index 0000000000..8edb30ce1f --- /dev/null +++ b/third_party/rust/metal/examples/mesh-shader/main.rs @@ -0,0 +1,118 @@ +extern crate objc; + +use cocoa::{appkit::NSView, base::id as cocoa_id}; +use core_graphics_types::geometry::CGSize; + +use metal::*; +use objc::{rc::autoreleasepool, runtime::YES}; +use std::mem; +use winit::platform::macos::WindowExtMacOS; + +use winit::{ + event::{Event, WindowEvent}, + event_loop::ControlFlow, +}; + +fn prepare_render_pass_descriptor(descriptor: &RenderPassDescriptorRef, texture: &TextureRef) { + let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + + color_attachment.set_texture(Some(texture)); + color_attachment.set_load_action(MTLLoadAction::Clear); + color_attachment.set_clear_color(MTLClearColor::new(0.2, 0.2, 0.25, 1.0)); + color_attachment.set_store_action(MTLStoreAction::Store); +} + +fn main() { + let events_loop = winit::event_loop::EventLoop::new(); + let size = winit::dpi::LogicalSize::new(800, 600); + + let window = winit::window::WindowBuilder::new() + .with_inner_size(size) + .with_title("Metal Mesh Shader Example".to_string()) + .build(&events_loop) + .unwrap(); + + let device = Device::system_default().expect("no device found"); + + let layer = MetalLayer::new(); + layer.set_device(&device); + layer.set_pixel_format(MTLPixelFormat::BGRA8Unorm); + layer.set_presents_with_transaction(false); + + unsafe { + let view = window.ns_view() as cocoa_id; + view.setWantsLayer(YES); + view.setLayer(mem::transmute(layer.as_ref())); + } + + let draw_size = window.inner_size(); + layer.set_drawable_size(CGSize::new(draw_size.width as f64, draw_size.height as f64)); + + let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/mesh-shader/shaders.metallib"); + let library = device.new_library_with_file(library_path).unwrap(); + + let mesh = library.get_function("mesh_function", None).unwrap(); + let frag = library.get_function("fragment_function", None).unwrap(); + + let pipeline_state_desc = MeshRenderPipelineDescriptor::new(); + pipeline_state_desc + .color_attachments() + .object_at(0) + .unwrap() + .set_pixel_format(MTLPixelFormat::BGRA8Unorm); + pipeline_state_desc.set_mesh_function(Some(&mesh)); + pipeline_state_desc.set_fragment_function(Some(&frag)); + + let pipeline_state = device + .new_mesh_render_pipeline_state(&pipeline_state_desc) + .unwrap(); + + let command_queue = device.new_command_queue(); + + events_loop.run(move |event, _, control_flow| { + autoreleasepool(|| { + *control_flow = ControlFlow::Poll; + + match event { + Event::WindowEvent { event, .. } => match event { + WindowEvent::CloseRequested => *control_flow = ControlFlow::Exit, + WindowEvent::Resized(size) => { + layer.set_drawable_size(CGSize::new(size.width as f64, size.height as f64)); + } + _ => (), + }, + Event::MainEventsCleared => { + window.request_redraw(); + } + Event::RedrawRequested(_) => { + let drawable = match layer.next_drawable() { + Some(drawable) => drawable, + None => return, + }; + + let render_pass_descriptor = RenderPassDescriptor::new(); + + prepare_render_pass_descriptor(&render_pass_descriptor, drawable.texture()); + + let command_buffer = command_queue.new_command_buffer(); + let encoder = + command_buffer.new_render_command_encoder(&render_pass_descriptor); + + encoder.set_render_pipeline_state(&pipeline_state); + encoder.draw_mesh_threads( + MTLSize::new(1, 1, 1), + MTLSize::new(1, 1, 1), + MTLSize::new(1, 1, 1), + ); + + encoder.end_encoding(); + + command_buffer.present_drawable(&drawable); + command_buffer.commit(); + } + _ => {} + } + }); + }); +} diff --git a/third_party/rust/metal/examples/mesh-shader/shaders.metal b/third_party/rust/metal/examples/mesh-shader/shaders.metal new file mode 100644 index 0000000000..1a82530742 --- /dev/null +++ b/third_party/rust/metal/examples/mesh-shader/shaders.metal @@ -0,0 +1,30 @@ +#include <metal_stdlib> + +using namespace metal; + +struct VertexOut { + float4 position [[position]]; +}; + +using mesh_t = mesh<VertexOut, void, 3, 1, topology::triangle>; + +[[mesh]] void mesh_function(mesh_t m) { + VertexOut v; + v.position = float4(-1.0, -1.0, 0.0, 1.0); + + m.set_primitive_count(1); + + m.set_vertex(0, v); + v.position = float4(0.0, 1.0, 0.0, 1.0); + m.set_vertex(1, v); + v.position = float4(1.0, -1.0, 0.0, 1.0); + m.set_vertex(2, v); + + m.set_index(0, 0); + m.set_index(1, 1); + m.set_index(2, 2); +} + +fragment half4 fragment_function() { + return half4(0.1, 1.0, 0.1, 1.0); +}
\ No newline at end of file diff --git a/third_party/rust/metal/examples/mesh-shader/shaders.metallib b/third_party/rust/metal/examples/mesh-shader/shaders.metallib Binary files differnew file mode 100644 index 0000000000..4af8d60ddc --- /dev/null +++ b/third_party/rust/metal/examples/mesh-shader/shaders.metallib diff --git a/third_party/rust/metal/examples/mps/main.rs b/third_party/rust/metal/examples/mps/main.rs new file mode 100644 index 0000000000..cc01b7a40d --- /dev/null +++ b/third_party/rust/metal/examples/mps/main.rs @@ -0,0 +1,148 @@ +use metal::*; +use std::ffi::c_void; +use std::mem; + +#[repr(C)] +struct Vertex { + xyz: [f32; 3], +} + +type Ray = mps::MPSRayOriginMinDistanceDirectionMaxDistance; +type Intersection = mps::MPSIntersectionDistancePrimitiveIndexCoordinates; + +// Original example taken from https://sergeyreznik.github.io/metal-ray-tracer/part-1/index.html +fn main() { + let device = Device::system_default().expect("No device found"); + + let library_path = + std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("examples/mps/shaders.metallib"); + let library = device + .new_library_with_file(library_path) + .expect("Failed to load shader library"); + + let generate_rays_pipeline = create_pipeline("generateRays", &library, &device); + + let queue = device.new_command_queue(); + let command_buffer = queue.new_command_buffer(); + + // Simple vertex/index buffer data + + let vertices: [Vertex; 3] = [ + Vertex { + xyz: [0.25, 0.25, 0.0], + }, + Vertex { + xyz: [0.75, 0.25, 0.0], + }, + Vertex { + xyz: [0.50, 0.75, 0.0], + }, + ]; + + let vertex_stride = mem::size_of::<Vertex>(); + + let indices: [u32; 3] = [0, 1, 2]; + + // Vertex data should be stored in private or managed buffers on discrete GPU systems (AMD, NVIDIA). + // Private buffers are stored entirely in GPU memory and cannot be accessed by the CPU. Managed + // buffers maintain a copy in CPU memory and a copy in GPU memory. + let buffer_opts = MTLResourceOptions::StorageModeManaged; + + let vertex_buffer = device.new_buffer_with_data( + vertices.as_ptr() as *const c_void, + (vertex_stride * vertices.len()) as u64, + buffer_opts, + ); + + let index_buffer = device.new_buffer_with_data( + indices.as_ptr() as *const c_void, + (mem::size_of::<u32>() * indices.len()) as u64, + buffer_opts, + ); + + // Build an acceleration structure using our vertex and index buffers containing the single triangle. + let acceleration_structure = mps::TriangleAccelerationStructure::from_device(&device) + .expect("Failed to create acceleration structure"); + + acceleration_structure.set_vertex_buffer(Some(&vertex_buffer)); + acceleration_structure.set_vertex_stride(vertex_stride as u64); + acceleration_structure.set_index_buffer(Some(&index_buffer)); + acceleration_structure.set_index_type(mps::MPSDataType::UInt32); + acceleration_structure.set_triangle_count(1); + acceleration_structure.set_usage(mps::MPSAccelerationStructureUsage::None); + acceleration_structure.rebuild(); + + let ray_intersector = + mps::RayIntersector::from_device(&device).expect("Failed to create ray intersector"); + + ray_intersector.set_ray_stride(mem::size_of::<Ray>() as u64); + ray_intersector.set_ray_data_type(mps::MPSRayDataType::OriginMinDistanceDirectionMaxDistance); + ray_intersector.set_intersection_stride(mem::size_of::<Intersection>() as u64); + ray_intersector.set_intersection_data_type( + mps::MPSIntersectionDataType::DistancePrimitiveIndexCoordinates, + ); + + // Create a buffer to hold generated rays and intersection results + let ray_count = 1024; + let ray_buffer = device.new_buffer( + (mem::size_of::<Ray>() * ray_count) as u64, + MTLResourceOptions::StorageModePrivate, + ); + + let intersection_buffer = device.new_buffer( + (mem::size_of::<Intersection>() * ray_count) as u64, + MTLResourceOptions::StorageModePrivate, + ); + + // Run the compute shader to generate rays + let encoder = command_buffer.new_compute_command_encoder(); + encoder.set_buffer(0, Some(&ray_buffer), 0); + encoder.set_compute_pipeline_state(&generate_rays_pipeline); + encoder.dispatch_thread_groups( + MTLSize { + width: 4, + height: 4, + depth: 1, + }, + MTLSize { + width: 8, + height: 8, + depth: 1, + }, + ); + encoder.end_encoding(); + + // Intersect rays with triangles inside acceleration structure + ray_intersector.encode_intersection_to_command_buffer( + &command_buffer, + mps::MPSIntersectionType::Nearest, + &ray_buffer, + 0, + &intersection_buffer, + 0, + ray_count as u64, + &acceleration_structure, + ); + + command_buffer.commit(); + command_buffer.wait_until_completed(); + + println!("Done"); +} + +fn create_pipeline(func: &str, library: &LibraryRef, device: &DeviceRef) -> ComputePipelineState { + // Create compute pipelines will will execute code on the GPU + let compute_descriptor = ComputePipelineDescriptor::new(); + + // Set to YES to allow compiler to make certain optimizations + compute_descriptor.set_thread_group_size_is_multiple_of_thread_execution_width(true); + + let function = library.get_function(func, None).unwrap(); + compute_descriptor.set_compute_function(Some(&function)); + + let pipeline = device + .new_compute_pipeline_state(&compute_descriptor) + .unwrap(); + + pipeline +} diff --git a/third_party/rust/metal/examples/mps/shaders.metal b/third_party/rust/metal/examples/mps/shaders.metal new file mode 100644 index 0000000000..d824d70d1b --- /dev/null +++ b/third_party/rust/metal/examples/mps/shaders.metal @@ -0,0 +1,26 @@ +// +// Created by Sergey Reznik on 9/15/18. +// Copyright © 2018 Serhii Rieznik. All rights reserved. +// + +// Taken from https://github.com/sergeyreznik/metal-ray-tracer/tree/part-1/source/Shaders +// MIT License https://github.com/sergeyreznik/metal-ray-tracer/blob/part-1/LICENSE + +#include <MetalPerformanceShaders/MetalPerformanceShaders.h> + +using Ray = MPSRayOriginMinDistanceDirectionMaxDistance; +using Intersection = MPSIntersectionDistancePrimitiveIndexCoordinates; + +kernel void generateRays( + device Ray* rays [[buffer(0)]], + uint2 coordinates [[thread_position_in_grid]], + uint2 size [[threads_per_grid]]) +{ + float2 uv = float2(coordinates) / float2(size - 1); + + uint rayIndex = coordinates.x + coordinates.y * size.x; + rays[rayIndex].origin = MPSPackedFloat3(uv.x, uv.y, -1.0); + rays[rayIndex].direction = MPSPackedFloat3(0.0, 0.0, 1.0); + rays[rayIndex].minDistance = 0.0f; + rays[rayIndex].maxDistance = 2.0f; +} diff --git a/third_party/rust/metal/examples/mps/shaders.metallib b/third_party/rust/metal/examples/mps/shaders.metallib Binary files differnew file mode 100644 index 0000000000..2cecf7b837 --- /dev/null +++ b/third_party/rust/metal/examples/mps/shaders.metallib diff --git a/third_party/rust/metal/examples/raytracing/README.md b/third_party/rust/metal/examples/raytracing/README.md new file mode 100644 index 0000000000..0071e59747 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/README.md @@ -0,0 +1,11 @@ +## Raytracing + +A good showcase of Metal 3 raytracing features. + + + +## To Run + +``` +cargo run --example raytracing +``` diff --git a/third_party/rust/metal/examples/raytracing/camera.rs b/third_party/rust/metal/examples/raytracing/camera.rs new file mode 100644 index 0000000000..5548445c06 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/camera.rs @@ -0,0 +1,20 @@ +use glam::f32::Vec4; + +#[repr(C)] +pub struct Camera { + pub position: Vec4, + pub right: Vec4, + pub up: Vec4, + pub forward: Vec4, +} + +impl Camera { + pub fn new() -> Self { + Self { + position: Vec4::new(0.0, 3.0, 10.0, 0.0), + right: Vec4::new(1.0, 0.0, 0.0, 0.0), + up: Vec4::new(0.0, 1.0, 0.0, 0.0), + forward: Vec4::new(0.0, 0.0, -1.0, 0.0), + } + } +} diff --git a/third_party/rust/metal/examples/raytracing/geometry.rs b/third_party/rust/metal/examples/raytracing/geometry.rs new file mode 100644 index 0000000000..93fdc196d1 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/geometry.rs @@ -0,0 +1,448 @@ +use std::{ + mem::{size_of, transmute}, + sync::Arc, +}; + +use glam::{ + f32::{Mat4, Vec3, Vec4}, + Vec4Swizzles, +}; + +use metal::*; + +pub const GEOMETRY_MASK_TRIANGLE: u32 = 1; +pub const GEOMETRY_MASK_SPHERE: u32 = 2; +pub const GEOMETRY_MASK_LIGHT: u32 = 4; + +pub const FACE_MASK_NONE: u16 = 0; +pub const FACE_MASK_NEGATIVE_X: u16 = 1 << 0; +pub const FACE_MASK_POSITIVE_X: u16 = 1 << 1; +pub const FACE_MASK_NEGATIVE_Y: u16 = 1 << 2; +pub const FACE_MASK_POSITIVE_Y: u16 = 1 << 3; +pub const FACE_MASK_NEGATIVE_Z: u16 = 1 << 4; +pub const FACE_MASK_POSITIVE_Z: u16 = 1 << 5; +pub const FACE_MASK_ALL: u16 = (1 << 6) - 1; + +pub trait Geometry { + fn upload_to_buffers(&mut self) { + todo!() + } + fn clear(&mut self) { + todo!() + } + fn get_geometry_descriptor(&self) -> AccelerationStructureGeometryDescriptor { + todo!() + } + fn get_resources(&self) -> Vec<Resource> { + todo!() + } + fn get_intersection_function_name(&self) -> Option<&str> { + None + } +} + +pub fn compute_triangle_normal(v0: &Vec3, v1: &Vec3, v2: &Vec3) -> Vec3 { + let e1 = Vec3::normalize(*v1 - *v0); + let e2 = Vec3::normalize(*v2 - *v0); + return Vec3::cross(e1, e2); +} + +#[derive(Default)] +#[repr(C)] +pub struct Triangle { + pub normals: [Vec4; 3], + pub colours: [Vec4; 3], +} + +pub fn get_managed_buffer_storage_mode() -> MTLResourceOptions { + return MTLResourceOptions::StorageModeManaged; +} + +pub struct TriangleGeometry { + pub device: Device, + pub name: String, + pub index_buffer: Option<Buffer>, + pub vertex_position_buffer: Option<Buffer>, + pub vertex_normal_buffer: Option<Buffer>, + pub vertex_colour_buffer: Option<Buffer>, + pub per_primitive_data_buffer: Option<Buffer>, + pub indices: Vec<u16>, + pub vertices: Vec<Vec4>, + pub normals: Vec<Vec4>, + pub colours: Vec<Vec4>, + pub triangles: Vec<Triangle>, +} + +impl TriangleGeometry { + pub fn new(device: Device, name: String) -> Self { + Self { + device, + name, + index_buffer: None, + vertex_position_buffer: None, + vertex_normal_buffer: None, + vertex_colour_buffer: None, + per_primitive_data_buffer: None, + indices: Vec::new(), + vertices: Vec::new(), + normals: Vec::new(), + colours: Vec::new(), + triangles: Vec::new(), + } + } + + pub fn add_cube_face_with_cube_vertices( + &mut self, + cube_vertices: &[Vec3], + colour: Vec3, + i0: u16, + i1: u16, + i2: u16, + i3: u16, + inward_normals: bool, + ) { + let v0 = cube_vertices[i0 as usize]; + let v1 = cube_vertices[i1 as usize]; + let v2 = cube_vertices[i2 as usize]; + let v3 = cube_vertices[i3 as usize]; + + let n0 = compute_triangle_normal(&v0, &v1, &v2) * if inward_normals { -1f32 } else { 1f32 }; + let n1 = compute_triangle_normal(&v0, &v2, &v3) * if inward_normals { -1f32 } else { 1f32 }; + + let first_index = self.indices.len(); + let base_index = self.vertices.len() as u16; + + self.indices.push(base_index + 0); + self.indices.push(base_index + 1); + self.indices.push(base_index + 2); + self.indices.push(base_index + 0); + self.indices.push(base_index + 2); + self.indices.push(base_index + 3); + + self.vertices.push(From::from((v0, 0.0))); + self.vertices.push(From::from((v1, 0.0))); + self.vertices.push(From::from((v2, 0.0))); + self.vertices.push(From::from((v3, 0.0))); + + self.normals + .push(From::from((Vec3::normalize(n0 + n1), 0.0))); + self.normals.push(From::from((n0, 0.0))); + self.normals + .push(From::from((Vec3::normalize(n0 + n1), 0.0))); + self.normals.push(From::from((n1, 0.0))); + + for _ in 0..4 { + self.colours.push(From::from((colour, 0.0))); + } + + for triangle_index in 0..2 { + let mut triangle = Triangle::default(); + for i in 0..3 { + let index = self.indices[first_index + triangle_index * 3 + i]; + triangle.normals[i] = self.normals[index as usize]; + triangle.colours[i] = self.colours[index as usize]; + } + self.triangles.push(triangle); + } + } + + pub fn add_cube_with_faces( + &mut self, + face_mask: u16, + colour: Vec3, + transform: Mat4, + inward_normals: bool, + ) { + let mut cube_vertices = [ + Vec3::new(-0.5, -0.5, -0.5), + Vec3::new(0.5, -0.5, -0.5), + Vec3::new(-0.5, 0.5, -0.5), + Vec3::new(0.5, 0.5, -0.5), + Vec3::new(-0.5, -0.5, 0.5), + Vec3::new(0.5, -0.5, 0.5), + Vec3::new(-0.5, 0.5, 0.5), + Vec3::new(0.5, 0.5, 0.5), + ]; + + for i in 0..8 { + let transformed_vertex = Vec4::from((cube_vertices[i], 1.0)); + let transformed_vertex = transform * transformed_vertex; + cube_vertices[i] = transformed_vertex.xyz(); + } + + const CUBE_INDICES: [[u16; 4]; 6] = [ + [0, 4, 6, 2], + [1, 3, 7, 5], + [0, 1, 5, 4], + [2, 6, 7, 3], + [0, 2, 3, 1], + [4, 5, 7, 6], + ]; + + for face in 0..6 { + if face_mask & (1 << face) != 0 { + self.add_cube_face_with_cube_vertices( + &cube_vertices, + colour, + CUBE_INDICES[face][0], + CUBE_INDICES[face][1], + CUBE_INDICES[face][2], + CUBE_INDICES[face][3], + inward_normals, + ); + } + } + } +} + +impl Geometry for TriangleGeometry { + fn upload_to_buffers(&mut self) { + self.index_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.indices.as_ptr()), + (self.indices.len() * size_of::<u16>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.vertex_position_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.vertices.as_ptr()), + (self.vertices.len() * size_of::<Vec4>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.vertex_normal_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.normals.as_ptr()), + (self.normals.len() * size_of::<Vec4>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.vertex_colour_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.colours.as_ptr()), + (self.colours.len() * size_of::<Vec4>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.per_primitive_data_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.triangles.as_ptr()), + (self.triangles.len() * size_of::<Triangle>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.index_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.index_buffer.as_ref().unwrap().length(), + )); + self.vertex_position_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.vertex_position_buffer.as_ref().unwrap().length(), + )); + self.vertex_normal_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.vertex_normal_buffer.as_ref().unwrap().length(), + )); + self.vertex_colour_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.vertex_colour_buffer.as_ref().unwrap().length(), + )); + self.per_primitive_data_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.per_primitive_data_buffer.as_ref().unwrap().length(), + )); + + self.index_buffer + .as_ref() + .unwrap() + .set_label(&format!("index buffer of {}", self.name)); + self.vertex_position_buffer + .as_ref() + .unwrap() + .set_label(&format!("vertex position buffer of {}", self.name)); + self.vertex_normal_buffer + .as_ref() + .unwrap() + .set_label(&format!("vertex normal buffer of {}", self.name)); + self.vertex_colour_buffer + .as_ref() + .unwrap() + .set_label(&format!("vertex colour buffer of {}", self.name)); + self.per_primitive_data_buffer + .as_ref() + .unwrap() + .set_label(&format!("per primitive data buffer of {}", self.name)); + } + + fn clear(&mut self) { + self.indices.clear(); + self.vertices.clear(); + self.normals.clear(); + self.colours.clear(); + self.triangles.clear(); + } + + fn get_geometry_descriptor(&self) -> AccelerationStructureGeometryDescriptor { + let descriptor = AccelerationStructureTriangleGeometryDescriptor::descriptor(); + + descriptor.set_index_buffer(Some(self.index_buffer.as_ref().unwrap())); + descriptor.set_index_type(MTLIndexType::UInt16); + descriptor.set_vertex_buffer(Some(self.vertex_position_buffer.as_ref().unwrap())); + descriptor.set_vertex_stride(size_of::<Vec4>() as NSUInteger); + descriptor.set_triangle_count((self.indices.len() / 3) as NSUInteger); + descriptor + .set_primitive_data_buffer(Some(self.per_primitive_data_buffer.as_ref().unwrap())); + descriptor.set_primitive_data_stride(size_of::<Triangle>() as NSUInteger); + descriptor.set_primitive_data_element_size(size_of::<Triangle>() as NSUInteger); + From::from(descriptor) + } + + fn get_resources(&self) -> Vec<Resource> { + vec![ + From::from(self.index_buffer.as_ref().unwrap().clone()), + From::from(self.vertex_normal_buffer.as_ref().unwrap().clone()), + From::from(self.vertex_colour_buffer.as_ref().unwrap().clone()), + ] + } +} + +#[repr(C)] +pub struct BoundingBox { + pub min: Vec3, + pub max: Vec3, +} + +#[repr(C)] +pub struct Sphere { + pub origin_radius_squared: Vec4, + pub colour_radius: Vec4, +} + +pub struct SphereGeometry { + pub device: Device, + pub sphere_buffer: Option<Buffer>, + pub bounding_box_buffer: Option<Buffer>, + pub per_primitive_data_buffer: Option<Buffer>, + pub spheres: Vec<Sphere>, +} + +impl SphereGeometry { + pub fn new(device: Device) -> Self { + Self { + device, + sphere_buffer: None, + bounding_box_buffer: None, + per_primitive_data_buffer: None, + spheres: Vec::new(), + } + } + + pub fn add_sphere_with_origin(&mut self, origin: Vec3, radius: f32, colour: Vec3) { + self.spheres.push(Sphere { + origin_radius_squared: Vec4::from((origin, radius * radius)), + colour_radius: Vec4::from((colour, radius)), + }); + } +} + +impl Geometry for SphereGeometry { + fn upload_to_buffers(&mut self) { + self.sphere_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(self.spheres.as_ptr()), + (self.spheres.len() * size_of::<Sphere>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.sphere_buffer + .as_ref() + .unwrap() + .set_label("sphere buffer"); + let mut bounding_boxes = Vec::new(); + for sphere in &self.spheres { + bounding_boxes.push(BoundingBox { + min: sphere.origin_radius_squared.xyz() - sphere.colour_radius.w, + max: sphere.origin_radius_squared.xyz() + sphere.colour_radius.w, + }); + } + self.bounding_box_buffer = Some(unsafe { + self.device.new_buffer_with_data( + transmute(bounding_boxes.as_ptr()), + (bounding_boxes.len() * size_of::<BoundingBox>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ) + }); + self.bounding_box_buffer + .as_ref() + .unwrap() + .set_label("bounding box buffer"); + self.sphere_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.sphere_buffer.as_ref().unwrap().length(), + )); + self.bounding_box_buffer + .as_ref() + .unwrap() + .did_modify_range(NSRange::new( + 0, + self.bounding_box_buffer.as_ref().unwrap().length(), + )); + } + + fn clear(&mut self) { + self.spheres.clear(); + } + + fn get_geometry_descriptor(&self) -> AccelerationStructureGeometryDescriptor { + let descriptor = AccelerationStructureBoundingBoxGeometryDescriptor::descriptor(); + descriptor.set_bounding_box_buffer(Some(self.bounding_box_buffer.as_ref().unwrap())); + descriptor.set_bounding_box_count(self.spheres.len() as NSUInteger); + descriptor.set_primitive_data_buffer(Some(&self.sphere_buffer.as_ref().unwrap())); + descriptor.set_primitive_data_stride(size_of::<Sphere>() as NSUInteger); + descriptor.set_primitive_data_element_size(size_of::<Sphere>() as NSUInteger); + From::from(descriptor) + } + + fn get_resources(&self) -> Vec<Resource> { + return vec![From::from(self.sphere_buffer.as_ref().unwrap().clone())]; + } + + fn get_intersection_function_name(&self) -> Option<&str> { + Some("sphereIntersectionFunction") + } +} + +pub struct GeometryInstance { + pub geometry: Arc<dyn Geometry>, + pub transform: Mat4, + pub mask: u32, + pub index_in_scene: NSUInteger, +} + +#[repr(C)] +pub struct AreaLight { + pub position: Vec4, + pub forward: Vec4, + pub right: Vec4, + pub up: Vec4, + pub colour: Vec4, +} diff --git a/third_party/rust/metal/examples/raytracing/main.rs b/third_party/rust/metal/examples/raytracing/main.rs new file mode 100644 index 0000000000..68eaf3df59 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/main.rs @@ -0,0 +1,87 @@ +extern crate objc; + +use cocoa::{appkit::NSView, base::id as cocoa_id}; +use core_graphics_types::geometry::CGSize; +use metal::*; +use objc::{rc::autoreleasepool, runtime::YES}; +use std::mem; +use winit::{ + event::{Event, WindowEvent}, + event_loop::ControlFlow, + platform::macos::WindowExtMacOS, +}; + +pub mod camera; +pub mod geometry; +pub mod renderer; +pub mod scene; + +fn find_raytracing_supporting_device() -> Device { + for device in Device::all() { + if !device.supports_raytracing() { + continue; + } + if device.is_low_power() { + continue; + } + return device; + } + + panic!("No device in this machine supports raytracing!") +} + +fn main() { + let events_loop = winit::event_loop::EventLoop::new(); + let size = winit::dpi::LogicalSize::new(800, 600); + + let window = winit::window::WindowBuilder::new() + .with_inner_size(size) + .with_title("Metal Raytracing Example".to_string()) + .build(&events_loop) + .unwrap(); + + let device = find_raytracing_supporting_device(); + + let layer = MetalLayer::new(); + layer.set_device(&device); + layer.set_pixel_format(MTLPixelFormat::RGBA16Float); + layer.set_presents_with_transaction(false); + + unsafe { + let view = window.ns_view() as cocoa_id; + view.setWantsLayer(YES); + view.setLayer(mem::transmute(layer.as_ref())); + } + + let draw_size = window.inner_size(); + let cg_size = CGSize::new(draw_size.width as f64, draw_size.height as f64); + layer.set_drawable_size(cg_size); + + let mut renderer = renderer::Renderer::new(device); + renderer.window_resized(cg_size); + + events_loop.run(move |event, _, control_flow| { + autoreleasepool(|| { + *control_flow = ControlFlow::Poll; + + match event { + Event::WindowEvent { event, .. } => match event { + WindowEvent::CloseRequested => *control_flow = ControlFlow::Exit, + WindowEvent::Resized(size) => { + let size = CGSize::new(size.width as f64, size.height as f64); + layer.set_drawable_size(size); + renderer.window_resized(size); + } + _ => (), + }, + Event::MainEventsCleared => { + window.request_redraw(); + } + Event::RedrawRequested(_) => { + renderer.draw(&layer); + } + _ => {} + } + }); + }); +} diff --git a/third_party/rust/metal/examples/raytracing/renderer.rs b/third_party/rust/metal/examples/raytracing/renderer.rs new file mode 100644 index 0000000000..f483d3e0a8 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/renderer.rs @@ -0,0 +1,512 @@ +use core_graphics_types::{base::CGFloat, geometry::CGSize}; +use std::{ + collections::BTreeMap, + ffi::c_void, + mem::{size_of, transmute}, + ops::Index, + sync::{Arc, Condvar, Mutex}, +}; + +use glam::{Vec3, Vec4, Vec4Swizzles}; +use rand::{thread_rng, RngCore}; + +use metal::{foreign_types::ForeignType, *}; + +use crate::{camera::Camera, geometry::get_managed_buffer_storage_mode, scene::Scene}; + +#[repr(C)] +struct Uniforms { + pub width: u32, + pub height: u32, + pub frame_index: u32, + pub light_count: u32, + pub camera: Camera, +} + +pub const MAX_FRAMES_IN_FLIGHT: NSUInteger = 3; +pub const ALIGNED_UNIFORMS_SIZE: NSUInteger = (size_of::<Uniforms>() as NSUInteger + 255) & !255; +pub const UNIFORM_BUFFER_SIZE: NSUInteger = MAX_FRAMES_IN_FLIGHT * ALIGNED_UNIFORMS_SIZE; + +#[derive(Clone)] +struct Semaphore { + data: Arc<(Mutex<usize>, Condvar)>, +} + +impl Semaphore { + fn new(capacity: usize) -> Self { + Self { + data: Arc::new((Mutex::new(capacity), Condvar::new())), + } + } + + fn acquire(&self) { + let mut value = self.data.0.lock().unwrap(); + while *value == 0 { + value = self.data.1.wait(value).unwrap(); + } + *value -= 1; + } + + fn release(&self) { + let mut value = self.data.0.lock().unwrap(); + *value += 1; + self.data.1.notify_one(); + } +} + +pub struct Renderer { + pub device: Device, + pub scene: Scene, + pub uniform_buffer: Buffer, + pub resource_buffer: Buffer, + pub instance_acceleration_structure: AccelerationStructure, + pub accumulation_targets: [Texture; 2], + pub random_texture: Texture, + pub frame_index: NSUInteger, + pub uniform_buffer_index: NSUInteger, + pub uniform_buffer_offset: NSUInteger, + pub size: CGSize, + semaphore: Semaphore, + pub queue: CommandQueue, + instance_buffer: Buffer, + intersection_function_table: IntersectionFunctionTable, + primitive_acceleration_structures: Vec<AccelerationStructure>, + raytracing_pipeline: ComputePipelineState, + copy_pipeline: RenderPipelineState, +} + +impl Renderer { + pub fn new(device: Device) -> Self { + let scene = Scene::new(device.clone()); + + let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/raytracing/shaders.metallib"); + let library = device.new_library_with_file(library_path).unwrap(); + let queue = device.new_command_queue(); + + let buffer_data = [0u8; UNIFORM_BUFFER_SIZE as usize]; + let uniform_buffer = device.new_buffer_with_data( + buffer_data.as_ptr() as *const c_void, + UNIFORM_BUFFER_SIZE, + get_managed_buffer_storage_mode(), + ); + uniform_buffer.set_label("uniform buffer"); + let resources_stride = { + let mut max = 0; + for geometry in &scene.geometries { + let s = geometry.get_resources().len(); + if s > max { + max = s; + } + } + max + }; + let mut resource_buffer_data = vec![0u64; resources_stride * scene.geometries.len()]; + for geometry_index in 0..scene.geometries.len() { + let geometry = scene.geometries[geometry_index].as_ref(); + let resource_buffer_begin_index = resources_stride * geometry_index; + let resources = geometry.get_resources(); + + for argument_index in 0..resources.len() { + let resource_buffer_index = resource_buffer_begin_index + argument_index; + let resource = resources[argument_index].clone(); + resource_buffer_data[resource_buffer_index] = + if resource.conforms_to_protocol::<MTLBuffer>().unwrap() { + let buffer = unsafe { Buffer::from_ptr(transmute(resource.into_ptr())) }; + buffer.gpu_address() + } else if resource.conforms_to_protocol::<MTLTexture>().unwrap() { + let texture = unsafe { Texture::from_ptr(transmute(resource.into_ptr())) }; + texture.gpu_resource_id()._impl + } else { + panic!("Unexpected resource!") + } + } + } + let resource_buffer = device.new_buffer_with_data( + resource_buffer_data.as_ptr() as *const c_void, + (resource_buffer_data.len() * size_of::<u64>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ); + resource_buffer.set_label("resource buffer"); + resource_buffer.did_modify_range(NSRange::new(0, resource_buffer.length())); + + let mut primitive_acceleration_structures = Vec::new(); + for i in 0..scene.geometries.len() { + let mesh = scene.geometries[i].as_ref(); + let geometry_descriptor = mesh.get_geometry_descriptor(); + geometry_descriptor.set_intersection_function_table_offset(i as NSUInteger); + let geometry_descriptors = Array::from_owned_slice(&[geometry_descriptor]); + let accel_descriptor = PrimitiveAccelerationStructureDescriptor::descriptor(); + accel_descriptor.set_geometry_descriptors(&geometry_descriptors); + let accel_descriptor: AccelerationStructureDescriptor = From::from(accel_descriptor); + primitive_acceleration_structures.push( + Self::new_acceleration_structure_with_descriptor( + &device, + &queue, + &accel_descriptor, + ), + ); + } + + let mut instance_descriptors = vec![ + MTLAccelerationStructureInstanceDescriptor::default(); + scene.geometry_instances.len() + ]; + for instance_index in 0..scene.geometry_instances.len() { + let instance = scene.geometry_instances[instance_index].as_ref(); + let geometry_index = instance.index_in_scene; + instance_descriptors[instance_index].acceleration_structure_index = + geometry_index as u32; + instance_descriptors[instance_index].options = + if instance.geometry.get_intersection_function_name().is_none() { + MTLAccelerationStructureInstanceOptions::Opaque + } else { + MTLAccelerationStructureInstanceOptions::None + }; + instance_descriptors[instance_index].intersection_function_table_offset = 0; + instance_descriptors[instance_index].mask = instance.mask as u32; + for column in 0..4 { + for row in 0..3 { + instance_descriptors[instance_index].transformation_matrix[column][row] = + *instance.transform.col(column).index(row); + } + } + } + let instance_buffer = device.new_buffer_with_data( + instance_descriptors.as_ptr() as *const c_void, + (size_of::<MTLAccelerationStructureInstanceDescriptor>() + * scene.geometry_instances.len()) as NSUInteger, + get_managed_buffer_storage_mode(), + ); + instance_buffer.set_label("instance buffer"); + instance_buffer.did_modify_range(NSRange::new(0, instance_buffer.length())); + + let accel_descriptor = InstanceAccelerationStructureDescriptor::descriptor(); + accel_descriptor.set_instanced_acceleration_structures(&Array::from_owned_slice( + &primitive_acceleration_structures, + )); + accel_descriptor.set_instance_count(scene.geometry_instances.len() as NSUInteger); + accel_descriptor.set_instance_descriptor_buffer(&instance_buffer); + let accel_descriptor: AccelerationStructureDescriptor = From::from(accel_descriptor); + let instance_acceleration_structure = + Self::new_acceleration_structure_with_descriptor(&device, &queue, &accel_descriptor); + + let mut intersection_functions = BTreeMap::<String, Function>::new(); + for geometry in &scene.geometries { + if let Some(name) = geometry.get_intersection_function_name() { + if !intersection_functions.contains_key(name) { + let intersection_function = Self::new_specialised_function_with_name( + &library, + resources_stride as u32, + name, + ); + intersection_functions.insert(name.to_string(), intersection_function); + } + } + } + let raytracing_function = Self::new_specialised_function_with_name( + &library, + resources_stride as u32, + "raytracingKernel", + ); + let intersection_function_array: Vec<&FunctionRef> = intersection_functions + .values() + .map(|f| -> &FunctionRef { f }) + .collect(); + let raytracing_pipeline = Self::new_compute_pipeline_state_with_function( + &device, + &raytracing_function, + &intersection_function_array, + ); + let intersection_function_table_descriptor = IntersectionFunctionTableDescriptor::new(); + intersection_function_table_descriptor + .set_function_count(scene.geometries.len() as NSUInteger); + let intersection_function_table = raytracing_pipeline + .new_intersection_function_table_with_descriptor( + &intersection_function_table_descriptor, + ); + for geometry_index in 0..scene.geometries.len() { + let geometry = scene.geometries[geometry_index].as_ref(); + if let Some(intersection_function_name) = geometry.get_intersection_function_name() { + let intersection_function = &intersection_functions[intersection_function_name]; + let handle = raytracing_pipeline + .function_handle_with_function(intersection_function) + .unwrap(); + intersection_function_table.set_function(handle, geometry_index as NSUInteger); + } + } + let render_descriptor = RenderPipelineDescriptor::new(); + render_descriptor + .set_vertex_function(Some(&library.get_function("copyVertex", None).unwrap())); + render_descriptor + .set_fragment_function(Some(&library.get_function("copyFragment", None).unwrap())); + render_descriptor + .color_attachments() + .object_at(0) + .unwrap() + .set_pixel_format(MTLPixelFormat::RGBA16Float); + let copy_pipeline = device + .new_render_pipeline_state(&render_descriptor) + .unwrap(); + + let texture_descriptor = Self::create_target_descriptor(1024, 1024); + let accumulation_targets = [ + device.new_texture(&texture_descriptor), + device.new_texture(&texture_descriptor), + ]; + let random_texture = device.new_texture(&texture_descriptor); + + Self { + device, + scene, + uniform_buffer, + resource_buffer, + instance_acceleration_structure, + accumulation_targets, + random_texture, + frame_index: 0, + uniform_buffer_index: 0, + uniform_buffer_offset: 0, + size: CGSize::new(1024 as CGFloat, 1024 as CGFloat), + semaphore: Semaphore::new((MAX_FRAMES_IN_FLIGHT - 2) as usize), + instance_buffer, + queue, + intersection_function_table, + primitive_acceleration_structures, + raytracing_pipeline, + copy_pipeline, + } + } + + fn create_target_descriptor(width: NSUInteger, height: NSUInteger) -> TextureDescriptor { + let texture_descriptor = TextureDescriptor::new(); + texture_descriptor.set_pixel_format(MTLPixelFormat::RGBA32Float); + texture_descriptor.set_texture_type(MTLTextureType::D2); + texture_descriptor.set_width(width); + texture_descriptor.set_height(height); + texture_descriptor.set_storage_mode(MTLStorageMode::Private); + texture_descriptor.set_usage(MTLTextureUsage::ShaderRead | MTLTextureUsage::ShaderWrite); + texture_descriptor + } + + pub fn window_resized(&mut self, size: CGSize) { + self.size = size; + let texture_descriptor = + Self::create_target_descriptor(size.width as NSUInteger, size.height as NSUInteger); + self.accumulation_targets[0] = self.device.new_texture(&texture_descriptor); + self.accumulation_targets[1] = self.device.new_texture(&texture_descriptor); + texture_descriptor.set_pixel_format(MTLPixelFormat::R32Uint); + texture_descriptor.set_usage(MTLTextureUsage::ShaderRead); + texture_descriptor.set_storage_mode(MTLStorageMode::Managed); + self.random_texture = self.device.new_texture(&texture_descriptor); + let mut rng = thread_rng(); + let mut random_values = vec![0u32; (size.width * size.height) as usize]; + for v in &mut random_values { + *v = rng.next_u32(); + } + self.random_texture.replace_region( + MTLRegion::new_2d(0, 0, size.width as NSUInteger, size.height as NSUInteger), + 0, + random_values.as_ptr() as *const c_void, + size_of::<u32>() as NSUInteger * size.width as NSUInteger, + ); + self.frame_index = 0; + } + + fn update_uniforms(&mut self) { + self.uniform_buffer_offset = ALIGNED_UNIFORMS_SIZE * self.uniform_buffer_index; + + let uniforms = unsafe { + &mut *((self.uniform_buffer.contents() as *mut u8) + .add(self.uniform_buffer_offset as usize) as *mut Uniforms) + }; + + let position = self.scene.camera.position; + let target = self.scene.camera.forward; + let up = self.scene.camera.up; + + let forward = Vec3::normalize(target.xyz() - position.xyz()); + let right = Vec3::normalize(Vec3::cross(forward, up.xyz())); + let up = Vec3::normalize(Vec3::cross(right, forward)); + + uniforms.camera.position = position; + uniforms.camera.forward = Vec4::from((forward, 0.0)); + uniforms.camera.right = Vec4::from((right, 0.0)); + uniforms.camera.up = Vec4::from((up, 0.0)); + + let field_of_view = 45.0 * (std::f32::consts::PI / 180.0); + let aspect_ratio = self.size.width as f32 / self.size.height as f32; + let image_plane_height = f32::tan(field_of_view / 2.0); + let image_plane_width = aspect_ratio * image_plane_height; + + uniforms.camera.right *= image_plane_width; + uniforms.camera.up *= image_plane_height; + + uniforms.width = self.size.width as u32; + uniforms.height = self.size.height as u32; + + uniforms.frame_index = self.frame_index as u32; + self.frame_index += 1; + + uniforms.light_count = self.scene.lights.len() as u32; + + self.uniform_buffer.did_modify_range(NSRange { + location: self.uniform_buffer_offset, + length: ALIGNED_UNIFORMS_SIZE, + }); + + self.uniform_buffer_index = (self.uniform_buffer_index + 1) % MAX_FRAMES_IN_FLIGHT; + } + + pub fn draw(&mut self, layer: &MetalLayer) { + self.semaphore.acquire(); + self.update_uniforms(); + let command_buffer = self.queue.new_command_buffer(); + let sem = self.semaphore.clone(); + let block = block::ConcreteBlock::new(move |_| { + sem.release(); + }) + .copy(); + command_buffer.add_completed_handler(&block); + let width = self.size.width as NSUInteger; + let height = self.size.height as NSUInteger; + let threads_per_thread_group = MTLSize::new(8, 8, 1); + let thread_groups = MTLSize::new( + (width + threads_per_thread_group.width - 1) / threads_per_thread_group.width, + (height + threads_per_thread_group.height - 1) / threads_per_thread_group.height, + 1, + ); + let compute_encoder = command_buffer.new_compute_command_encoder(); + compute_encoder.set_buffer(0, Some(&self.uniform_buffer), self.uniform_buffer_offset); + compute_encoder.set_buffer(2, Some(&self.instance_buffer), 0); + compute_encoder.set_buffer(3, Some(&self.scene.lights_buffer), 0); + compute_encoder.set_acceleration_structure(4, Some(&self.instance_acceleration_structure)); + compute_encoder.set_intersection_function_table(5, Some(&self.intersection_function_table)); + compute_encoder.set_texture(0, Some(&self.random_texture)); + compute_encoder.set_texture(1, Some(&self.accumulation_targets[0])); + compute_encoder.set_texture(2, Some(&self.accumulation_targets[1])); + for geometry in &self.scene.geometries { + for resource in geometry.get_resources() { + compute_encoder.use_resource(&resource, MTLResourceUsage::Read); + } + } + for primitive_acceleration_structure in &self.primitive_acceleration_structures { + let resource: Resource = From::from(primitive_acceleration_structure.clone()); + compute_encoder.use_resource(&resource, MTLResourceUsage::Read); + } + compute_encoder.set_compute_pipeline_state(&self.raytracing_pipeline); + compute_encoder.dispatch_thread_groups(thread_groups, threads_per_thread_group); + compute_encoder.end_encoding(); + (self.accumulation_targets[0], self.accumulation_targets[1]) = ( + self.accumulation_targets[1].clone(), + self.accumulation_targets[0].clone(), + ); + if let Some(drawable) = layer.next_drawable() { + let render_pass_descriptor = RenderPassDescriptor::new(); + let colour_attachment = render_pass_descriptor + .color_attachments() + .object_at(0) + .unwrap(); + colour_attachment.set_texture(Some(drawable.texture())); + colour_attachment.set_load_action(MTLLoadAction::Clear); + colour_attachment.set_clear_color(MTLClearColor::new(0.0, 0.0, 0.0, 1.0)); + let render_encoder = command_buffer.new_render_command_encoder(render_pass_descriptor); + render_encoder.set_render_pipeline_state(&self.copy_pipeline); + render_encoder.set_fragment_texture(0, Some(&self.accumulation_targets[0])); + render_encoder.draw_primitives(MTLPrimitiveType::Triangle, 0, 6); + render_encoder.end_encoding(); + command_buffer.present_drawable(&drawable); + } + command_buffer.commit(); + } + + fn new_acceleration_structure_with_descriptor( + device: &Device, + queue: &CommandQueue, + descriptor: &AccelerationStructureDescriptorRef, + ) -> AccelerationStructure { + let accel_sizes = device.acceleration_structure_sizes_with_descriptor(descriptor); + let acceleration_structure = + device.new_acceleration_structure_with_size(accel_sizes.acceleration_structure_size); + let scratch_buffer = device.new_buffer( + accel_sizes.build_scratch_buffer_size, + MTLResourceOptions::StorageModePrivate, + ); + let command_buffer = queue.new_command_buffer(); + let command_encoder = command_buffer.new_acceleration_structure_command_encoder(); + let compacted_size_buffer = device.new_buffer( + size_of::<u32>() as NSUInteger, + MTLResourceOptions::StorageModeShared, + ); + command_encoder.build_acceleration_structure( + &acceleration_structure, + &descriptor, + &scratch_buffer, + 0, + ); + command_encoder.write_compacted_acceleration_structure_size( + &acceleration_structure, + &compacted_size_buffer, + 0, + ); + command_encoder.end_encoding(); + command_buffer.commit(); + command_buffer.wait_until_completed(); + let compacted_size: *const u32 = unsafe { transmute(compacted_size_buffer.contents()) }; + let compacted_size = unsafe { *compacted_size } as NSUInteger; + let compacted_acceleration_structure = + device.new_acceleration_structure_with_size(compacted_size); + let command_buffer = queue.new_command_buffer(); + let command_encoder = command_buffer.new_acceleration_structure_command_encoder(); + command_encoder.copy_and_compact_acceleration_structure( + &acceleration_structure, + &compacted_acceleration_structure, + ); + command_encoder.end_encoding(); + command_buffer.commit(); + compacted_acceleration_structure + } + + fn new_specialised_function_with_name( + library: &Library, + resources_stride: u32, + name: &str, + ) -> Function { + let constants = FunctionConstantValues::new(); + let resources_stride = resources_stride * size_of::<u64>() as u32; + constants.set_constant_value_at_index( + &resources_stride as *const u32 as *const c_void, + MTLDataType::UInt, + 0, + ); + let v = true; + constants.set_constant_value_at_index( + &v as *const bool as *const c_void, + MTLDataType::Bool, + 1, + ); + constants.set_constant_value_at_index( + &v as *const bool as *const c_void, + MTLDataType::Bool, + 2, + ); + library.get_function(name, Some(constants)).unwrap() + } + + fn new_compute_pipeline_state_with_function( + device: &Device, + function: &Function, + linked_functions: &[&FunctionRef], + ) -> ComputePipelineState { + let linked_functions = { + let lf = LinkedFunctions::new(); + lf.set_functions(linked_functions); + lf + }; + let descriptor = ComputePipelineDescriptor::new(); + descriptor.set_compute_function(Some(function)); + descriptor.set_linked_functions(linked_functions.as_ref()); + descriptor.set_thread_group_size_is_multiple_of_thread_execution_width(true); + device.new_compute_pipeline_state(&descriptor).unwrap() + } +} diff --git a/third_party/rust/metal/examples/raytracing/scene.rs b/third_party/rust/metal/examples/raytracing/scene.rs new file mode 100644 index 0000000000..8ecf8568de --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/scene.rs @@ -0,0 +1,135 @@ +use std::{ffi::c_void, mem::size_of, sync::Arc}; + +use glam::{Mat4, Vec3, Vec4}; +use rand::{thread_rng, Rng}; + +use metal::{Buffer, Device, NSRange, NSUInteger}; + +use super::{camera::Camera, geometry::*}; + +pub struct Scene { + pub device: Device, + pub camera: Camera, + pub geometries: Vec<Arc<dyn Geometry>>, + pub geometry_instances: Vec<Arc<GeometryInstance>>, + pub lights: Vec<AreaLight>, + pub lights_buffer: Buffer, +} + +impl Scene { + pub fn new(device: Device) -> Self { + let mut geometries = Vec::<Arc<dyn Geometry>>::new(); + let mut light_mesh = TriangleGeometry::new(device.clone(), "light".to_string()); + let transform = Mat4::from_translation(Vec3::new(0.0, 1.0, 0.0)) + * Mat4::from_scale(Vec3::new(0.5, 1.98, 0.5)); + light_mesh.add_cube_with_faces( + FACE_MASK_POSITIVE_Y, + Vec3::new(1.0, 1.0, 1.0), + transform, + true, + ); + light_mesh.upload_to_buffers(); + let light_mesh = Arc::new(light_mesh); + geometries.push(light_mesh.clone()); + + let mut geometry_mesh = TriangleGeometry::new(device.clone(), "geometry".to_string()); + let transform = Mat4::from_translation(Vec3::new(0.0, 1.0, 0.0)) + * Mat4::from_scale(Vec3::new(2.0, 2.0, 2.0)); + geometry_mesh.add_cube_with_faces( + FACE_MASK_NEGATIVE_Y | FACE_MASK_POSITIVE_Y | FACE_MASK_NEGATIVE_Z, + Vec3::new(0.725, 0.71, 0.68), + transform, + true, + ); + geometry_mesh.add_cube_with_faces( + FACE_MASK_NEGATIVE_X, + Vec3::new(0.63, 0.065, 0.05), + transform, + true, + ); + geometry_mesh.add_cube_with_faces( + FACE_MASK_POSITIVE_X, + Vec3::new(0.14, 0.45, 0.091), + transform, + true, + ); + let transform = Mat4::from_translation(Vec3::new(-0.335, 0.6, -0.29)) + * Mat4::from_rotation_y(0.3) + * Mat4::from_scale(Vec3::new(0.6, 1.2, 0.6)); + geometry_mesh.add_cube_with_faces( + FACE_MASK_ALL, + Vec3::new(0.725, 0.71, 0.68), + transform, + false, + ); + geometry_mesh.upload_to_buffers(); + let geometry_mesh = Arc::new(geometry_mesh); + geometries.push(geometry_mesh.clone()); + + let mut sphere_geometry = SphereGeometry::new(device.clone()); + sphere_geometry.add_sphere_with_origin( + Vec3::new(0.3275, 0.3, 0.3725), + 0.3, + Vec3::new(0.725, 0.71, 0.68), + ); + sphere_geometry.upload_to_buffers(); + let sphere_geometry = Arc::new(sphere_geometry); + geometries.push(sphere_geometry.clone()); + + let mut rng = thread_rng(); + let mut geometry_instances = Vec::new(); + let mut lights = Vec::new(); + for y in -1..2 { + for x in -1..2 { + let transform = + Mat4::from_translation(Vec3::new(x as f32 * 2.5, y as f32 * 2.5, 0.0)); + geometry_instances.push(Arc::new(GeometryInstance { + geometry: light_mesh.clone(), + transform, + mask: GEOMETRY_MASK_LIGHT, + index_in_scene: 0, + })); + geometry_instances.push(Arc::new(GeometryInstance { + geometry: geometry_mesh.clone(), + transform, + mask: GEOMETRY_MASK_TRIANGLE, + index_in_scene: 1, + })); + geometry_instances.push(Arc::new(GeometryInstance { + geometry: sphere_geometry.clone(), + transform, + mask: GEOMETRY_MASK_SPHERE, + index_in_scene: 2, + })); + lights.push(AreaLight { + position: Vec4::new(x as f32 * 2.5, y as f32 * 2.5 + 1.98, 0.0, 0.0), + forward: Vec4::new(0.0, -1.0, 0.0, 0.0), + right: Vec4::new(0.25, 0.0, 0.0, 0.0), + up: Vec4::new(0.0, 0.0, 0.25, 0.0), + colour: Vec4::new( + rng.gen_range(0f32..=1.0), + rng.gen_range(0f32..=1.0), + rng.gen_range(0f32..=1.0), + 0.0, + ), + }); + } + } + let lights_buffer = device.new_buffer_with_data( + lights.as_ptr() as *const c_void, + (lights.len() * size_of::<AreaLight>()) as NSUInteger, + get_managed_buffer_storage_mode(), + ); + lights_buffer.did_modify_range(NSRange::new(0, lights_buffer.length())); + lights_buffer.set_label("lights buffer"); + + Self { + device, + camera: Camera::new(), + geometries, + geometry_instances, + lights, + lights_buffer, + } + } +} diff --git a/third_party/rust/metal/examples/raytracing/screenshot.png b/third_party/rust/metal/examples/raytracing/screenshot.png Binary files differnew file mode 100644 index 0000000000..417a1d746d --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/screenshot.png diff --git a/third_party/rust/metal/examples/raytracing/shaders.metal b/third_party/rust/metal/examples/raytracing/shaders.metal new file mode 100644 index 0000000000..54aa2a4f47 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/shaders.metal @@ -0,0 +1,598 @@ +/* +See LICENSE folder for this sample’s licensing information. + +Abstract: +The Metal shaders used for this sample. +*/ + +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +using namespace raytracing; + + +#define GEOMETRY_MASK_TRIANGLE 1 +#define GEOMETRY_MASK_SPHERE 2 +#define GEOMETRY_MASK_LIGHT 4 + +#define GEOMETRY_MASK_GEOMETRY (GEOMETRY_MASK_TRIANGLE | GEOMETRY_MASK_SPHERE) + +#define RAY_MASK_PRIMARY (GEOMETRY_MASK_GEOMETRY | GEOMETRY_MASK_LIGHT) +#define RAY_MASK_SHADOW GEOMETRY_MASK_GEOMETRY +#define RAY_MASK_SECONDARY GEOMETRY_MASK_GEOMETRY + +struct Camera { + vector_float3 position; + vector_float3 right; + vector_float3 up; + vector_float3 forward; +}; + +struct AreaLight { + vector_float3 position; + vector_float3 forward; + vector_float3 right; + vector_float3 up; + vector_float3 color; +}; + +struct Uniforms { + unsigned int width; + unsigned int height; + unsigned int frameIndex; + unsigned int lightCount; + Camera camera; +}; + +struct Sphere { + packed_float3 origin; + float radiusSquared; + packed_float3 color; + float radius; +}; + +struct Triangle { + vector_float3 normals[3]; + vector_float3 colors[3]; +}; + +constant unsigned int resourcesStride [[function_constant(0)]]; +constant bool useIntersectionFunctions [[function_constant(1)]]; +constant bool usePerPrimitiveData [[function_constant(2)]]; +constant bool useResourcesBuffer = !usePerPrimitiveData; + +constant unsigned int primes[] = { + 2, 3, 5, 7, + 11, 13, 17, 19, + 23, 29, 31, 37, + 41, 43, 47, 53, + 59, 61, 67, 71, + 73, 79, 83, 89 +}; + +// Returns the i'th element of the Halton sequence using the d'th prime number as a +// base. The Halton sequence is a low discrepency sequence: the values appear +// random, but are more evenly distributed than a purely random sequence. Each random +// value used to render the image uses a different independent dimension, `d`, +// and each sample (frame) uses a different index `i`. To decorrelate each pixel, +// you can apply a random offset to `i`. +float halton(unsigned int i, unsigned int d) { + unsigned int b = primes[d]; + + float f = 1.0f; + float invB = 1.0f / b; + + float r = 0; + + while (i > 0) { + f = f * invB; + r = r + f * (i % b); + i = i / b; + } + + return r; +} + +// Interpolates the vertex attribute of an arbitrary type across the surface of a triangle +// given the barycentric coordinates and triangle index in an intersection structure. +template<typename T, typename IndexType> +inline T interpolateVertexAttribute(device T *attributes, + IndexType i0, + IndexType i1, + IndexType i2, + float2 uv) { + // Look up value for each vertex. + const T T0 = attributes[i0]; + const T T1 = attributes[i1]; + const T T2 = attributes[i2]; + + // Compute the sum of the vertex attributes weighted by the barycentric coordinates. + // The barycentric coordinates sum to one. + return (1.0f - uv.x - uv.y) * T0 + uv.x * T1 + uv.y * T2; +} + +template<typename T> +inline T interpolateVertexAttribute(thread T *attributes, float2 uv) { + // Look up the value for each vertex. + const T T0 = attributes[0]; + const T T1 = attributes[1]; + const T T2 = attributes[2]; + + // Compute the sum of the vertex attributes weighted by the barycentric coordinates. + // The barycentric coordinates sum to one. + return (1.0f - uv.x - uv.y) * T0 + uv.x * T1 + uv.y * T2; +} + +// Uses the inversion method to map two uniformly random numbers to a 3D +// unit hemisphere, where the probability of a given sample is proportional to the cosine +// of the angle between the sample direction and the "up" direction (0, 1, 0). +inline float3 sampleCosineWeightedHemisphere(float2 u) { + float phi = 2.0f * M_PI_F * u.x; + + float cos_phi; + float sin_phi = sincos(phi, cos_phi); + + float cos_theta = sqrt(u.y); + float sin_theta = sqrt(1.0f - cos_theta * cos_theta); + + return float3(sin_theta * cos_phi, cos_theta, sin_theta * sin_phi); +} + +// Maps two uniformly random numbers to the surface of a 2D area light +// source and returns the direction to this point, the amount of light that travels +// between the intersection point and the sample point on the light source, as well +// as the distance between these two points. + +inline void sampleAreaLight(constant AreaLight & light, + float2 u, + float3 position, + thread float3 & lightDirection, + thread float3 & lightColor, + thread float & lightDistance) +{ + // Map to -1..1 + u = u * 2.0f - 1.0f; + + // Transform into the light's coordinate system. + float3 samplePosition = light.position + + light.right * u.x + + light.up * u.y; + + // Compute the vector from sample point on the light source to intersection point. + lightDirection = samplePosition - position; + + lightDistance = length(lightDirection); + + float inverseLightDistance = 1.0f / max(lightDistance, 1e-3f); + + // Normalize the light direction. + lightDirection *= inverseLightDistance; + + // Start with the light's color. + lightColor = light.color; + + // Light falls off with the inverse square of the distance to the intersection point. + lightColor *= (inverseLightDistance * inverseLightDistance); + + // Light also falls off with the cosine of the angle between the intersection point + // and the light source. + lightColor *= saturate(dot(-lightDirection, light.forward)); +} + +// Aligns a direction on the unit hemisphere such that the hemisphere's "up" direction +// (0, 1, 0) maps to the given surface normal direction. +inline float3 alignHemisphereWithNormal(float3 sample, float3 normal) { + // Set the "up" vector to the normal + float3 up = normal; + + // Find an arbitrary direction perpendicular to the normal, which becomes the + // "right" vector. + float3 right = normalize(cross(normal, float3(0.0072f, 1.0f, 0.0034f))); + + // Find a third vector perpendicular to the previous two, which becomes the + // "forward" vector. + float3 forward = cross(right, up); + + // Map the direction on the unit hemisphere to the coordinate system aligned + // with the normal. + return sample.x * right + sample.y * up + sample.z * forward; +} + +// Return the type for a bounding box intersection function. +struct BoundingBoxIntersection { + bool accept [[accept_intersection]]; // Whether to accept or reject the intersection. + float distance [[distance]]; // Distance from the ray origin to the intersection point. +}; + +// Resources for a piece of triangle geometry. +struct TriangleResources { + device uint16_t *indices; + device float3 *vertexNormals; + device float3 *vertexColors; +}; + +// Resources for a piece of sphere geometry. +struct SphereResources { + device Sphere *spheres; +}; + +/* + Custom sphere intersection function. The [[intersection]] keyword marks this as an intersection + function. The [[bounding_box]] keyword means that this intersection function handles intersecting rays + with bounding box primitives. To create sphere primitives, the sample creates bounding boxes that + enclose the sphere primitives. + + The [[triangle_data]] and [[instancing]] keywords indicate that the intersector that calls this + intersection function returns barycentric coordinates for triangle intersections and traverses + an instance acceleration structure. These keywords must match between the intersection functions, + intersection function table, intersector, and intersection result to ensure that Metal propagates + data correctly between stages. Using fewer tags when possible may result in better performance, + as Metal may need to store less data and pass less data between stages. For example, if you do not + need barycentric coordinates, omitting [[triangle_data]] means Metal can avoid computing and storing + them. + + The arguments to the intersection function contain information about the ray, primitive to be + tested, and so on. The ray intersector provides this datas when it calls the intersection function. + Metal provides other built-in arguments, but this sample doesn't use them. + */ +[[intersection(bounding_box, triangle_data, instancing)]] +BoundingBoxIntersection sphereIntersectionFunction(// Ray parameters passed to the ray intersector below + float3 origin [[origin]], + float3 direction [[direction]], + float minDistance [[min_distance]], + float maxDistance [[max_distance]], + // Information about the primitive. + unsigned int primitiveIndex [[primitive_id]], + unsigned int geometryIndex [[geometry_intersection_function_table_offset]], + // Custom resources bound to the intersection function table. + device void *resources [[buffer(0), function_constant(useResourcesBuffer)]] + ,const device void* perPrimitiveData [[primitive_data]]) +{ + Sphere sphere; + // Look up the resources for this piece of sphere geometry. + if (usePerPrimitiveData) { + // Per-primitive data points to data from the specified buffer as was configured in the MTLAccelerationStructureBoundingBoxGeometryDescriptor. + sphere = *(const device Sphere*)perPrimitiveData; + } else + { + device SphereResources& sphereResources = *(device SphereResources *)((device char *)resources + resourcesStride * geometryIndex); + // Get the actual sphere enclosed in this bounding box. + sphere = sphereResources.spheres[primitiveIndex]; + } + + // Check for intersection between the ray and sphere mathematically. + float3 oc = origin - sphere.origin; + + float a = dot(direction, direction); + float b = 2 * dot(oc, direction); + float c = dot(oc, oc) - sphere.radiusSquared; + + float disc = b * b - 4 * a * c; + + BoundingBoxIntersection ret; + + if (disc <= 0.0f) { + // If the ray missed the sphere, return false. + ret.accept = false; + } + else { + // Otherwise, compute the intersection distance. + ret.distance = (-b - sqrt(disc)) / (2 * a); + + // The intersection function must also check whether the intersection distance is + // within the acceptable range. Intersection functions do not run in any particular order, + // so the maximum distance may be different from the one passed into the ray intersector. + ret.accept = ret.distance >= minDistance && ret.distance <= maxDistance; + } + + return ret; +} + +__attribute__((always_inline)) +float3 transformPoint(float3 p, float4x4 transform) { + return (transform * float4(p.x, p.y, p.z, 1.0f)).xyz; +} + +__attribute__((always_inline)) +float3 transformDirection(float3 p, float4x4 transform) { + return (transform * float4(p.x, p.y, p.z, 0.0f)).xyz; +} + +// Main ray tracing kernel. +kernel void raytracingKernel( + uint2 tid [[thread_position_in_grid]], + constant Uniforms & uniforms [[buffer(0)]], + texture2d<unsigned int> randomTex [[texture(0)]], + texture2d<float> prevTex [[texture(1)]], + texture2d<float, access::write> dstTex [[texture(2)]], + device void *resources [[buffer(1), function_constant(useResourcesBuffer)]], + constant MTLAccelerationStructureInstanceDescriptor *instances [[buffer(2)]], + constant AreaLight *areaLights [[buffer(3)]], + instance_acceleration_structure accelerationStructure [[buffer(4)]], + intersection_function_table<triangle_data, instancing> intersectionFunctionTable [[buffer(5)]] +) +{ + // The sample aligns the thread count to the threadgroup size, which means the thread count + // may be different than the bounds of the texture. Test to make sure this thread + // is referencing a pixel within the bounds of the texture. + if (tid.x >= uniforms.width || tid.y >= uniforms.height) return; + + // The ray to cast. + ray ray; + + // Pixel coordinates for this thread. + float2 pixel = (float2)tid; + + // Apply a random offset to the random number index to decorrelate pixels. + unsigned int offset = randomTex.read(tid).x; + + // Add a random offset to the pixel coordinates for antialiasing. + float2 r = float2(halton(offset + uniforms.frameIndex, 0), + halton(offset + uniforms.frameIndex, 1)); + + pixel += r; + + // Map pixel coordinates to -1..1. + float2 uv = (float2)pixel / float2(uniforms.width, uniforms.height); + uv = uv * 2.0f - 1.0f; + + constant Camera & camera = uniforms.camera; + + // Rays start at the camera position. + ray.origin = camera.position; + + // Map normalized pixel coordinates into camera's coordinate system. + ray.direction = normalize(uv.x * camera.right + + uv.y * camera.up + + camera.forward); + + // Don't limit intersection distance. + ray.max_distance = INFINITY; + + // Start with a fully white color. The kernel scales the light each time the + // ray bounces off of a surface, based on how much of each light component + // the surface absorbs. + float3 color = float3(1.0f, 1.0f, 1.0f); + + float3 accumulatedColor = float3(0.0f, 0.0f, 0.0f); + + // Create an intersector to test for intersection between the ray and the geometry in the scene. + intersector<triangle_data, instancing> i; + + // If the sample isn't using intersection functions, provide some hints to Metal for + // better performance. + if (!useIntersectionFunctions) { + i.assume_geometry_type(geometry_type::triangle); + i.force_opacity(forced_opacity::opaque); + } + + typename intersector<triangle_data, instancing>::result_type intersection; + + // Simulate up to three ray bounces. Each bounce propagates light backward along the + // ray's path toward the camera. + for (int bounce = 0; bounce < 3; bounce++) { + // Get the closest intersection, not the first intersection. This is the default, but + // the sample adjusts this property below when it casts shadow rays. + i.accept_any_intersection(false); + + // Check for intersection between the ray and the acceleration structure. If the sample + // isn't using intersection functions, it doesn't need to include one. + if (useIntersectionFunctions) + intersection = i.intersect(ray, accelerationStructure, bounce == 0 ? RAY_MASK_PRIMARY : RAY_MASK_SECONDARY, intersectionFunctionTable); + else + intersection = i.intersect(ray, accelerationStructure, bounce == 0 ? RAY_MASK_PRIMARY : RAY_MASK_SECONDARY); + + // Stop if the ray didn't hit anything and has bounced out of the scene. + if (intersection.type == intersection_type::none) + break; + + unsigned int instanceIndex = intersection.instance_id; + + // Look up the mask for this instance, which indicates what type of geometry the ray hit. + unsigned int mask = instances[instanceIndex].mask; + + // If the ray hit a light source, set the color to white, and stop immediately. + if (mask == GEOMETRY_MASK_LIGHT) { + accumulatedColor = float3(1.0f, 1.0f, 1.0f); + break; + } + + // The ray hit something. Look up the transformation matrix for this instance. + float4x4 objectToWorldSpaceTransform(1.0f); + + for (int column = 0; column < 4; column++) + for (int row = 0; row < 3; row++) + objectToWorldSpaceTransform[column][row] = instances[instanceIndex].transformationMatrix[column][row]; + + // Compute the intersection point in world space. + float3 worldSpaceIntersectionPoint = ray.origin + ray.direction * intersection.distance; + + unsigned primitiveIndex = intersection.primitive_id; + unsigned int geometryIndex = instances[instanceIndex].accelerationStructureIndex; + float2 barycentric_coords = intersection.triangle_barycentric_coord; + + float3 worldSpaceSurfaceNormal = 0.0f; + float3 surfaceColor = 0.0f; + + if (mask & GEOMETRY_MASK_TRIANGLE) { + Triangle triangle; + + float3 objectSpaceSurfaceNormal; + if (usePerPrimitiveData) { + // Per-primitive data points to data from the specified buffer as was configured in the MTLAccelerationStructureTriangleGeometryDescriptor. + triangle = *(const device Triangle*)intersection.primitive_data; + } else + { + // The ray hit a triangle. Look up the corresponding geometry's normal and UV buffers. + device TriangleResources & triangleResources = *(device TriangleResources *)((device char *)resources + resourcesStride * geometryIndex); + + triangle.normals[0] = triangleResources.vertexNormals[triangleResources.indices[primitiveIndex * 3 + 0]]; + triangle.normals[1] = triangleResources.vertexNormals[triangleResources.indices[primitiveIndex * 3 + 1]]; + triangle.normals[2] = triangleResources.vertexNormals[triangleResources.indices[primitiveIndex * 3 + 2]]; + + triangle.colors[0] = triangleResources.vertexColors[triangleResources.indices[primitiveIndex * 3 + 0]]; + triangle.colors[1] = triangleResources.vertexColors[triangleResources.indices[primitiveIndex * 3 + 1]]; + triangle.colors[2] = triangleResources.vertexColors[triangleResources.indices[primitiveIndex * 3 + 2]]; + } + + // Interpolate the vertex normal at the intersection point. + objectSpaceSurfaceNormal = interpolateVertexAttribute(triangle.normals, barycentric_coords); + + // Interpolate the vertex color at the intersection point. + surfaceColor = interpolateVertexAttribute(triangle.colors, barycentric_coords); + + // Transform the normal from object to world space. + worldSpaceSurfaceNormal = normalize(transformDirection(objectSpaceSurfaceNormal, objectToWorldSpaceTransform)); + } + else if (mask & GEOMETRY_MASK_SPHERE) { + Sphere sphere; + if (usePerPrimitiveData) { + // Per-primitive data points to data from the specified buffer as was configured in the MTLAccelerationStructureBoundingBoxGeometryDescriptor. + sphere = *(const device Sphere*)intersection.primitive_data; + } else + { + // The ray hit a sphere. Look up the corresponding sphere buffer. + device SphereResources & sphereResources = *(device SphereResources *)((device char *)resources + resourcesStride * geometryIndex); + sphere = sphereResources.spheres[primitiveIndex]; + } + + // Transform the sphere's origin from object space to world space. + float3 worldSpaceOrigin = transformPoint(sphere.origin, objectToWorldSpaceTransform); + + // Compute the surface normal directly in world space. + worldSpaceSurfaceNormal = normalize(worldSpaceIntersectionPoint - worldSpaceOrigin); + + // The sphere is a uniform color, so you don't need to interpolate the color across the surface. + surfaceColor = sphere.color; + } + + dstTex.write(float4(accumulatedColor, 1.0f), tid); + + // Choose a random light source to sample. + float lightSample = halton(offset + uniforms.frameIndex, 2 + bounce * 5 + 0); + unsigned int lightIndex = min((unsigned int)(lightSample * uniforms.lightCount), uniforms.lightCount - 1); + + // Choose a random point to sample on the light source. + float2 r = float2(halton(offset + uniforms.frameIndex, 2 + bounce * 5 + 1), + halton(offset + uniforms.frameIndex, 2 + bounce * 5 + 2)); + + float3 worldSpaceLightDirection; + float3 lightColor; + float lightDistance; + + // Sample the lighting between the intersection point and the point on the area light. + sampleAreaLight(areaLights[lightIndex], r, worldSpaceIntersectionPoint, worldSpaceLightDirection, + lightColor, lightDistance); + + // Scale the light color by the cosine of the angle between the light direction and + // surface normal. + lightColor *= saturate(dot(worldSpaceSurfaceNormal, worldSpaceLightDirection)); + + // Scale the light color by the number of lights to compensate for the fact that + // the sample samples only one light source at random. + lightColor *= uniforms.lightCount; + + // Scale the ray color by the color of the surface to simulate the surface absorbing light. + color *= surfaceColor; + + // Compute the shadow ray. The shadow ray checks whether the sample position on the + // light source is visible from the current intersection point. + // If it is, the kernel adds lighting to the output image. + struct ray shadowRay; + + // Add a small offset to the intersection point to avoid intersecting the same + // triangle again. + shadowRay.origin = worldSpaceIntersectionPoint + worldSpaceSurfaceNormal * 1e-3f; + + // Travel toward the light source. + shadowRay.direction = worldSpaceLightDirection; + + // Don't overshoot the light source. + shadowRay.max_distance = lightDistance - 1e-3f; + + // Shadow rays check only whether there is an object between the intersection point + // and the light source. Tell Metal to return after finding any intersection. + i.accept_any_intersection(true); + + if (useIntersectionFunctions) + intersection = i.intersect(shadowRay, accelerationStructure, RAY_MASK_SHADOW, intersectionFunctionTable); + else + intersection = i.intersect(shadowRay, accelerationStructure, RAY_MASK_SHADOW); + + // If there was no intersection, then the light source is visible from the original + // intersection point. Add the light's contribution to the image. + if (intersection.type == intersection_type::none) + accumulatedColor += lightColor * color; + + // Choose a random direction to continue the path of the ray. This causes light to + // bounce between surfaces. An app might evaluate a more complicated equation to + // calculate the amount of light that reflects between intersection points. However, + // all the math in this kernel cancels out because this app assumes a simple diffuse + // BRDF and samples the rays with a cosine distribution over the hemisphere (importance + // sampling). This requires that the kernel only multiply the colors together. This + // sampling strategy also reduces the amount of noise in the output image. + r = float2(halton(offset + uniforms.frameIndex, 2 + bounce * 5 + 3), + halton(offset + uniforms.frameIndex, 2 + bounce * 5 + 4)); + + float3 worldSpaceSampleDirection = sampleCosineWeightedHemisphere(r); + worldSpaceSampleDirection = alignHemisphereWithNormal(worldSpaceSampleDirection, worldSpaceSurfaceNormal); + + ray.origin = worldSpaceIntersectionPoint + worldSpaceSurfaceNormal * 1e-3f; + ray.direction = worldSpaceSampleDirection; + } + + // Average this frame's sample with all of the previous frames. + if (uniforms.frameIndex > 0) { + float3 prevColor = prevTex.read(tid).xyz; + prevColor *= uniforms.frameIndex; + + accumulatedColor += prevColor; + accumulatedColor /= (uniforms.frameIndex + 1); + } + + dstTex.write(float4(accumulatedColor, 1.0f), tid); +} + +// Screen filling quad in normalized device coordinates. +constant float2 quadVertices[] = { + float2(-1, -1), + float2(-1, 1), + float2( 1, 1), + float2(-1, -1), + float2( 1, 1), + float2( 1, -1) +}; + +struct CopyVertexOut { + float4 position [[position]]; + float2 uv; +}; + +// Simple vertex shader that passes through NDC quad positions. +vertex CopyVertexOut copyVertex(unsigned short vid [[vertex_id]]) { + float2 position = quadVertices[vid]; + + CopyVertexOut out; + + out.position = float4(position, 0, 1); + out.uv = position * 0.5f + 0.5f; + + return out; +} + +// Simple fragment shader that copies a texture and applies a simple tonemapping function. +fragment float4 copyFragment(CopyVertexOut in [[stage_in]], + texture2d<float> tex) +{ + constexpr sampler sam(min_filter::nearest, mag_filter::nearest, mip_filter::none); + + float3 color = tex.sample(sam, in.uv).xyz; + + // Apply a simple tonemapping function to reduce the dynamic range of the + // input image into a range which the screen can display. + color = color / (1.0f + color); + + return float4(color, 1.0f); +} diff --git a/third_party/rust/metal/examples/raytracing/shaders.metallib b/third_party/rust/metal/examples/raytracing/shaders.metallib Binary files differnew file mode 100644 index 0000000000..0965a64ff5 --- /dev/null +++ b/third_party/rust/metal/examples/raytracing/shaders.metallib diff --git a/third_party/rust/metal/examples/reflection/main.rs b/third_party/rust/metal/examples/reflection/main.rs new file mode 100644 index 0000000000..058199cc35 --- /dev/null +++ b/third_party/rust/metal/examples/reflection/main.rs @@ -0,0 +1,75 @@ +// Copyright 2016 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; + +const PROGRAM: &'static str = r" + #include <metal_stdlib> + + using namespace metal; + + typedef struct { + float2 position; + float3 color; + } vertex_t; + + struct ColorInOut { + float4 position [[position]]; + float4 color; + }; + + vertex ColorInOut vs(device vertex_t* vertex_array [[ buffer(0) ]], + unsigned int vid [[ vertex_id ]]) + { + ColorInOut out; + + out.position = float4(float2(vertex_array[vid].position), 0.0, 1.0); + out.color = float4(float3(vertex_array[vid].color), 1.0); + + return out; + } + + fragment float4 ps(ColorInOut in [[stage_in]]) + { + return in.color; + }; +"; + +fn main() { + autoreleasepool(|| { + let device = Device::system_default().expect("no device found"); + + let options = CompileOptions::new(); + let library = device.new_library_with_source(PROGRAM, &options).unwrap(); + let (vs, ps) = ( + library.get_function("vs", None).unwrap(), + library.get_function("ps", None).unwrap(), + ); + + let vertex_desc = VertexDescriptor::new(); + + let desc = RenderPipelineDescriptor::new(); + desc.set_vertex_function(Some(&vs)); + desc.set_fragment_function(Some(&ps)); + desc.set_vertex_descriptor(Some(vertex_desc)); + + println!("{:?}", desc); + + let reflect_options = MTLPipelineOption::ArgumentInfo | MTLPipelineOption::BufferTypeInfo; + let (_, reflection) = device + .new_render_pipeline_state_with_reflection(&desc, reflect_options) + .unwrap(); + + println!("Vertex arguments: "); + let vertex_arguments = reflection.vertex_arguments(); + for index in 0..vertex_arguments.count() { + let argument = vertex_arguments.object_at(index).unwrap(); + println!("{:?}", argument); + } + }); +} diff --git a/third_party/rust/metal/examples/shader-dylib/main.rs b/third_party/rust/metal/examples/shader-dylib/main.rs new file mode 100644 index 0000000000..60b6aa0d0c --- /dev/null +++ b/third_party/rust/metal/examples/shader-dylib/main.rs @@ -0,0 +1,177 @@ +use cocoa::{appkit::NSView, base::id as cocoa_id}; +use core_graphics_types::geometry::CGSize; + +use metal::*; +use objc::{rc::autoreleasepool, runtime::YES}; + +use winit::{ + event::{Event, WindowEvent}, + event_loop::ControlFlow, + platform::macos::WindowExtMacOS, +}; + +use std::mem; + +struct App { + pub _device: Device, + pub command_queue: CommandQueue, + pub layer: MetalLayer, + pub image_fill_cps: ComputePipelineState, + pub width: u32, + pub height: u32, +} + +fn select_device() -> Option<Device> { + let devices = Device::all(); + for device in devices { + if device.supports_dynamic_libraries() { + return Some(device); + } + } + + None +} + +impl App { + fn new(window: &winit::window::Window) -> Self { + let device = select_device().expect("no device found that supports dynamic libraries"); + let command_queue = device.new_command_queue(); + + let layer = MetalLayer::new(); + layer.set_device(&device); + layer.set_pixel_format(MTLPixelFormat::BGRA8Unorm); + layer.set_presents_with_transaction(false); + layer.set_framebuffer_only(false); + unsafe { + let view = window.ns_view() as cocoa_id; + view.setWantsLayer(YES); + view.setLayer(mem::transmute(layer.as_ref())); + } + let draw_size = window.inner_size(); + layer.set_drawable_size(CGSize::new(draw_size.width as f64, draw_size.height as f64)); + + // compile dynamic lib shader + let dylib_src_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/shader-dylib/test_dylib.metal"); + let install_path = + std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("target/test_dylib.metallib"); + + let dylib_src = std::fs::read_to_string(dylib_src_path).expect("bad shit"); + let opts = metal::CompileOptions::new(); + opts.set_library_type(MTLLibraryType::Dynamic); + opts.set_install_name(install_path.to_str().unwrap()); + + let lib = device + .new_library_with_source(dylib_src.as_str(), &opts) + .unwrap(); + + // create dylib + let dylib = device.new_dynamic_library(&lib).unwrap(); + dylib.set_label("test_dylib"); + + // optional: serialize binary blob that can be loaded later + let blob_url = String::from("file://") + install_path.to_str().unwrap(); + let url = URL::new_with_string(&blob_url); + dylib.serialize_to_url(&url).unwrap(); + + // create shader that links with dylib + let shader_src_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/shader-dylib/test_shader.metal"); + + let shader_src = std::fs::read_to_string(shader_src_path).expect("bad shit"); + let opts = metal::CompileOptions::new(); + // add dynamic library to link with + let libraries = [dylib.as_ref()]; + opts.set_libraries(&libraries); + + // compile + let shader_lib = device + .new_library_with_source(shader_src.as_str(), &opts) + .unwrap(); + + let func = shader_lib.get_function("test_kernel", None).unwrap(); + + // create pipeline state + // linking occurs here + let image_fill_cps = device + .new_compute_pipeline_state_with_function(&func) + .unwrap(); + + Self { + _device: device, + command_queue, + layer, + image_fill_cps, + width: draw_size.width, + height: draw_size.height, + } + } + + fn resize(&mut self, width: u32, height: u32) { + self.layer + .set_drawable_size(CGSize::new(width as f64, height as f64)); + self.width = width; + self.height = height; + } + + fn draw(&self) { + let drawable = match self.layer.next_drawable() { + Some(drawable) => drawable, + None => return, + }; + + let w = self.image_fill_cps.thread_execution_width(); + let h = self.image_fill_cps.max_total_threads_per_threadgroup() / w; + let threads_per_threadgroup = MTLSize::new(w, h, 1); + let threads_per_grid = MTLSize::new(self.width as _, self.height as _, 1); + + let command_buffer = self.command_queue.new_command_buffer(); + + { + let encoder = command_buffer.new_compute_command_encoder(); + encoder.set_compute_pipeline_state(&self.image_fill_cps); + encoder.set_texture(0, Some(&drawable.texture())); + encoder.dispatch_threads(threads_per_grid, threads_per_threadgroup); + encoder.end_encoding(); + } + + command_buffer.present_drawable(&drawable); + command_buffer.commit(); + } +} + +fn main() { + let events_loop = winit::event_loop::EventLoop::new(); + let size = winit::dpi::LogicalSize::new(800, 600); + + let window = winit::window::WindowBuilder::new() + .with_inner_size(size) + .with_title("Metal Shader Dylib Example".to_string()) + .build(&events_loop) + .unwrap(); + + let mut app = App::new(&window); + + events_loop.run(move |event, _, control_flow| { + autoreleasepool(|| { + *control_flow = ControlFlow::Poll; + + match event { + Event::WindowEvent { event, .. } => match event { + WindowEvent::CloseRequested => *control_flow = ControlFlow::Exit, + WindowEvent::Resized(size) => { + app.resize(size.width, size.height); + } + _ => (), + }, + Event::MainEventsCleared => { + window.request_redraw(); + } + Event::RedrawRequested(_) => { + app.draw(); + } + _ => {} + } + }); + }); +} diff --git a/third_party/rust/metal/examples/shader-dylib/test_dylib.metal b/third_party/rust/metal/examples/shader-dylib/test_dylib.metal new file mode 100644 index 0000000000..5faa4a803a --- /dev/null +++ b/third_party/rust/metal/examples/shader-dylib/test_dylib.metal @@ -0,0 +1,8 @@ +#include <metal_stdlib> + +using namespace metal; + +float4 get_color_test(float4 inColor) +{ + return float4(inColor.r, inColor.g, inColor.b, 0); +} diff --git a/third_party/rust/metal/examples/shader-dylib/test_shader.metal b/third_party/rust/metal/examples/shader-dylib/test_shader.metal new file mode 100644 index 0000000000..38203a64a5 --- /dev/null +++ b/third_party/rust/metal/examples/shader-dylib/test_shader.metal @@ -0,0 +1,14 @@ +#include <metal_stdlib> + +using namespace metal; + +extern float4 get_color_test(float4 inColor); + +kernel void test_kernel( + texture2d<float, access::write> image [[texture(0)]], + uint2 coordinates [[thread_position_in_grid]], + uint2 size [[threads_per_grid]]) +{ + float2 uv = float2(coordinates) / float2(size - 1); + image.write(get_color_test(float4(uv, 0.0, 1.0)), coordinates); +} diff --git a/third_party/rust/metal/examples/window/README.md b/third_party/rust/metal/examples/window/README.md new file mode 100644 index 0000000000..62233be356 --- /dev/null +++ b/third_party/rust/metal/examples/window/README.md @@ -0,0 +1,11 @@ +## window + +Renders a spinning triangle to a [winit](https://github.com/rust-windowing/winit) window. + + + +## To Run + +``` +cargo run --example window +``` diff --git a/third_party/rust/metal/examples/window/main.rs b/third_party/rust/metal/examples/window/main.rs new file mode 100644 index 0000000000..08936e82fc --- /dev/null +++ b/third_party/rust/metal/examples/window/main.rs @@ -0,0 +1,261 @@ +// Copyright 2016 metal-rs 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. + +extern crate objc; + +use cocoa::{appkit::NSView, base::id as cocoa_id}; +use core_graphics_types::geometry::CGSize; + +use metal::*; +use objc::{rc::autoreleasepool, runtime::YES}; +use std::mem; +use winit::platform::macos::WindowExtMacOS; + +use winit::{ + event::{Event, WindowEvent}, + event_loop::ControlFlow, +}; + +#[repr(C)] +struct Rect { + pub x: f32, + pub y: f32, + pub w: f32, + pub h: f32, +} + +#[repr(C)] +struct Color { + pub r: f32, + pub g: f32, + pub b: f32, + pub a: f32, +} + +#[repr(C)] +struct ClearRect { + pub rect: Rect, + pub color: Color, +} + +fn prepare_pipeline_state<'a>( + device: &DeviceRef, + library: &LibraryRef, + vertex_shader: &str, + fragment_shader: &str, +) -> RenderPipelineState { + let vert = library.get_function(vertex_shader, None).unwrap(); + let frag = library.get_function(fragment_shader, None).unwrap(); + + let pipeline_state_descriptor = RenderPipelineDescriptor::new(); + pipeline_state_descriptor.set_vertex_function(Some(&vert)); + pipeline_state_descriptor.set_fragment_function(Some(&frag)); + let attachment = pipeline_state_descriptor + .color_attachments() + .object_at(0) + .unwrap(); + attachment.set_pixel_format(MTLPixelFormat::BGRA8Unorm); + + attachment.set_blending_enabled(true); + attachment.set_rgb_blend_operation(metal::MTLBlendOperation::Add); + attachment.set_alpha_blend_operation(metal::MTLBlendOperation::Add); + attachment.set_source_rgb_blend_factor(metal::MTLBlendFactor::SourceAlpha); + attachment.set_source_alpha_blend_factor(metal::MTLBlendFactor::SourceAlpha); + attachment.set_destination_rgb_blend_factor(metal::MTLBlendFactor::OneMinusSourceAlpha); + attachment.set_destination_alpha_blend_factor(metal::MTLBlendFactor::OneMinusSourceAlpha); + + device + .new_render_pipeline_state(&pipeline_state_descriptor) + .unwrap() +} + +fn prepare_render_pass_descriptor(descriptor: &RenderPassDescriptorRef, texture: &TextureRef) { + //descriptor.color_attachments().set_object_at(0, MTLRenderPassColorAttachmentDescriptor::alloc()); + //let color_attachment: MTLRenderPassColorAttachmentDescriptor = unsafe { msg_send![descriptor.color_attachments().0, _descriptorAtIndex:0] };//descriptor.color_attachments().object_at(0); + let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + + color_attachment.set_texture(Some(texture)); + color_attachment.set_load_action(MTLLoadAction::Clear); + color_attachment.set_clear_color(MTLClearColor::new(0.2, 0.2, 0.25, 1.0)); + color_attachment.set_store_action(MTLStoreAction::Store); +} + +fn main() { + let events_loop = winit::event_loop::EventLoop::new(); + let size = winit::dpi::LogicalSize::new(800, 600); + + let window = winit::window::WindowBuilder::new() + .with_inner_size(size) + .with_title("Metal Window Example".to_string()) + .build(&events_loop) + .unwrap(); + + let device = Device::system_default().expect("no device found"); + + let layer = MetalLayer::new(); + layer.set_device(&device); + layer.set_pixel_format(MTLPixelFormat::BGRA8Unorm); + layer.set_presents_with_transaction(false); + + unsafe { + let view = window.ns_view() as cocoa_id; + view.setWantsLayer(YES); + view.setLayer(mem::transmute(layer.as_ref())); + } + + let draw_size = window.inner_size(); + layer.set_drawable_size(CGSize::new(draw_size.width as f64, draw_size.height as f64)); + + let library_path = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")) + .join("examples/window/shaders.metallib"); + + let library = device.new_library_with_file(library_path).unwrap(); + let triangle_pipeline_state = + prepare_pipeline_state(&device, &library, "triangle_vertex", "triangle_fragment"); + let clear_rect_pipeline_state = prepare_pipeline_state( + &device, + &library, + "clear_rect_vertex", + "clear_rect_fragment", + ); + + let command_queue = device.new_command_queue(); + //let nc: () = msg_send![command_queue.0, setExecutionEnabled:true]; + + let vbuf = { + let vertex_data = [ + 0.0f32, 0.5, 1.0, 0.0, 0.0, -0.5, -0.5, 0.0, 1.0, 0.0, 0.5, 0.5, 0.0, 0.0, 1.0, + ]; + + device.new_buffer_with_data( + vertex_data.as_ptr() as *const _, + (vertex_data.len() * mem::size_of::<f32>()) as u64, + MTLResourceOptions::CPUCacheModeDefaultCache | MTLResourceOptions::StorageModeManaged, + ) + }; + + let mut r = 0.0f32; + + let clear_rect = vec![ClearRect { + rect: Rect { + x: -1.0, + y: -1.0, + w: 2.0, + h: 2.0, + }, + color: Color { + r: 0.5, + g: 0.8, + b: 0.5, + a: 1.0, + }, + }]; + + let clear_rect_buffer = device.new_buffer_with_data( + clear_rect.as_ptr() as *const _, + mem::size_of::<ClearRect>() as u64, + MTLResourceOptions::CPUCacheModeDefaultCache | MTLResourceOptions::StorageModeManaged, + ); + + events_loop.run(move |event, _, control_flow| { + autoreleasepool(|| { + *control_flow = ControlFlow::Poll; + + match event { + Event::WindowEvent { event, .. } => match event { + WindowEvent::CloseRequested => *control_flow = ControlFlow::Exit, + WindowEvent::Resized(size) => { + layer.set_drawable_size(CGSize::new(size.width as f64, size.height as f64)); + } + _ => (), + }, + Event::MainEventsCleared => { + window.request_redraw(); + } + Event::RedrawRequested(_) => { + let p = vbuf.contents(); + let vertex_data = [ + 0.0f32, + 0.5, + 1.0, + 0.0, + 0.0, + -0.5 + (r.cos() / 2. + 0.5), + -0.5, + 0.0, + 1.0, + 0.0, + 0.5 - (r.cos() / 2. + 0.5), + -0.5, + 0.0, + 0.0, + 1.0, + ]; + + unsafe { + std::ptr::copy( + vertex_data.as_ptr(), + p as *mut f32, + (vertex_data.len() * mem::size_of::<f32>()) as usize, + ); + } + + vbuf.did_modify_range(crate::NSRange::new( + 0 as u64, + (vertex_data.len() * mem::size_of::<f32>()) as u64, + )); + + let drawable = match layer.next_drawable() { + Some(drawable) => drawable, + None => return, + }; + + let render_pass_descriptor = RenderPassDescriptor::new(); + + prepare_render_pass_descriptor(&render_pass_descriptor, drawable.texture()); + + let command_buffer = command_queue.new_command_buffer(); + let encoder = + command_buffer.new_render_command_encoder(&render_pass_descriptor); + + encoder.set_scissor_rect(MTLScissorRect { + x: 20, + y: 20, + width: 100, + height: 100, + }); + encoder.set_render_pipeline_state(&clear_rect_pipeline_state); + encoder.set_vertex_buffer(0, Some(&clear_rect_buffer), 0); + encoder.draw_primitives_instanced( + metal::MTLPrimitiveType::TriangleStrip, + 0, + 4, + 1, + ); + let physical_size = window.inner_size(); + encoder.set_scissor_rect(MTLScissorRect { + x: 0, + y: 0, + width: physical_size.width as _, + height: physical_size.height as _, + }); + + encoder.set_render_pipeline_state(&triangle_pipeline_state); + encoder.set_vertex_buffer(0, Some(&vbuf), 0); + encoder.draw_primitives(MTLPrimitiveType::Triangle, 0, 3); + encoder.end_encoding(); + + command_buffer.present_drawable(&drawable); + command_buffer.commit(); + + r += 0.01f32; + } + _ => {} + } + }); + }); +} diff --git a/third_party/rust/metal/examples/window/screenshot.png b/third_party/rust/metal/examples/window/screenshot.png Binary files differnew file mode 100644 index 0000000000..9f5eba8ccf --- /dev/null +++ b/third_party/rust/metal/examples/window/screenshot.png diff --git a/third_party/rust/metal/examples/window/shaders.metal b/third_party/rust/metal/examples/window/shaders.metal new file mode 100644 index 0000000000..cc05f5d57e --- /dev/null +++ b/third_party/rust/metal/examples/window/shaders.metal @@ -0,0 +1,97 @@ +#include <metal_stdlib> + +using namespace metal; + +typedef struct { + packed_float2 position; + packed_float3 color; +} vertex_t; + +struct ColorInOut { + float4 position [[position]]; + float4 color; +}; +// vertex shader function +vertex ColorInOut triangle_vertex(const device vertex_t* vertex_array [[ buffer(0) ]], + unsigned int vid [[ vertex_id ]]) +{ + ColorInOut out; + + auto device const &v = vertex_array[vid]; + out.position = float4(v.position.x, v.position.y, 0.0, 1.0); + out.color = float4(v.color.x, v.color.y, v.color.z, 0.2); + + return out; +} + +// fragment shader function +fragment float4 triangle_fragment(ColorInOut in [[stage_in]]) +{ + return in.color; +}; + + +struct Rect { + float x; + float y; + float w; + float h; +}; + +struct Color { + float r; + float g; + float b; + float a; +}; + +struct ClearRect { + Rect rect; + Color color; +}; + +float2 rect_vert( + Rect rect, + uint vid +) { + float2 pos; + + float left = rect.x; + float right = rect.x + rect.w; + float bottom = rect.y; + float top = rect.y + rect.h; + + switch (vid) { + case 0: + pos = float2(right, top); + break; + case 1: + pos = float2(left, top); + break; + case 2: + pos = float2(right, bottom); + break; + case 3: + pos = float2(left, bottom); + break; + } + return pos; +} + +vertex ColorInOut clear_rect_vertex( + const device ClearRect *clear_rect [[ buffer(0) ]], + unsigned int vid [[ vertex_id ]] +) { + ColorInOut out; + float4 pos = float4(rect_vert(clear_rect->rect, vid), 0, 1); + auto col = clear_rect->color; + + out.position = pos; + out.color = float4(col.r, col.g, col.b, col.a); + return out; +} + +fragment float4 clear_rect_fragment(ColorInOut in [[stage_in]]) +{ + return in.color; +}; diff --git a/third_party/rust/metal/examples/window/shaders.metallib b/third_party/rust/metal/examples/window/shaders.metallib Binary files differnew file mode 100644 index 0000000000..a6388fc9bc --- /dev/null +++ b/third_party/rust/metal/examples/window/shaders.metallib |
