Struct MetalFunction

Source

pub struct MetalFunction { /* private fields */ }

Expand description

Apple’s id<MTLFunction> — a single compiled shader entry point.

Implementations§

Source §

impl MetalFunction

Source

pub fn new_argument_encoder( &self, buffer_index: usize, ) -> Option<ArgumentEncoder>

Create an argument encoder for the argument buffer bound at buffer_index.

Examples found in repository ?

examples/06_resources_and_archives.rs (line 26)

10fn main() {
11    let device = MetalDevice::system_default().expect("Metal device available");
12    println!("device: {}", device.name());
13
14    let queue = device
15        .new_command_queue_with_max_command_buffer_count(4)
16        .expect("bounded command queue");
17    let scratch = queue
18        .new_command_buffer_with_unretained_references()
19        .expect("bounded scratch command buffer");
20    println!("bounded queue scratch status={}", scratch.status());
21
22    let library = device
23        .new_library_with_source(common::COMPUTE_SRC)
24        .expect("compile compute library");
25    let args = library.new_function("use_args").expect("use_args function");
26    let argument_encoder = args.new_argument_encoder(0).expect("argument encoder");
27    let argument_buffer = device
28        .new_buffer(
29            argument_encoder.encoded_length(),
30            resource_options::STORAGE_MODE_SHARED,
31        )
32        .expect("argument buffer");
33    let payload = device
34        .new_buffer(16, resource_options::STORAGE_MODE_SHARED)
35        .expect("payload buffer");
36    let texture = device
37        .new_texture(TextureDescriptor::new_2d(4, 4, pixel_format::BGRA8UNORM))
38        .expect("argument texture");
39    argument_encoder.set_argument_buffer(&argument_buffer, 0);
40    argument_encoder.set_buffer(&payload, 0, 0);
41    argument_encoder.set_texture(&texture, 1);
42    println!(
43        "argument encoder length={} alignment={}",
44        argument_encoder.encoded_length(),
45        argument_encoder.alignment(),
46    );
47
48    let backing = device
49        .new_buffer(256, resource_options::STORAGE_MODE_SHARED)
50        .expect("backing buffer");
51    let buffer_texture = backing
52        .new_texture_view_2d(pixel_format::BGRA8UNORM, 16, 4, 64, 0)
53        .expect("buffer-backed texture");
54    println!(
55        "buffer-backed texture {}x{} fmt={}",
56        buffer_texture.width(),
57        buffer_texture.height(),
58        buffer_texture.pixel_format(),
59    );
60
61    if let Some(heap) = device.new_heap(1 << 20, storage_mode::SHARED) {
62        let heap_buffer = heap
63            .new_buffer(256, resource_options::STORAGE_MODE_SHARED)
64            .expect("heap buffer");
65        let heap_texture = heap
66            .new_texture(TextureDescriptor::new_2d(4, 4, pixel_format::BGRA8UNORM))
67            .expect("heap texture");
68        println!(
69            "heap size={} used={} current={} max_available={}",
70            heap.size(),
71            heap.used_size(),
72            heap.current_allocated_size(),
73            heap.max_available_size(256),
74        );
75        println!(
76            "heap buffer len={} heap texture {}x{} purgeable={}",
77            heap_buffer.length(),
78            heap_texture.width(),
79            heap_texture.height(),
80            heap.set_purgeable_state(apple_metal::purgeable_state::KEEP_CURRENT),
81        );
82    } else {
83        println!("heaps are unavailable on this device");
84    }
85
86    match device.new_log_state(log_level::INFO, 1_024) {
87        Ok(log_state) => {
88            let _ = device
89                .new_command_queue_with_log_state(4, &log_state)
90                .expect("log-state queue");
91            println!("created queue with log state");
92        }
93        Err(error) => println!("log state unavailable on this OS: {error}"),
94    }
95
96    if device.supports_dynamic_libraries() {
97        let dynamic_path = common::artifact_path("example-dylib.metallib");
98        let dynamic_library = device
99            .new_dynamic_library_with_source(
100                common::DYNAMIC_LIB_SRC,
101                dynamic_path.to_string_lossy().as_ref(),
102            )
103            .expect("dynamic library from source");
104        dynamic_library
105            .serialize_to_file(&dynamic_path)
106            .expect("serialize dynamic library");
107        let reloaded = device
108            .load_dynamic_library(&dynamic_path)
109            .expect("reload dynamic library");
110        println!("dynamic library install name: {}", reloaded.install_name());
111
112        let render_library = device
113            .new_library_with_source(common::RENDER_SRC)
114            .expect("compile render library");
115        let vertex = render_library
116            .new_function("fullscreen_vertex")
117            .expect("vertex function");
118        let fragment = render_library
119            .new_function("solid_fragment")
120            .expect("fragment function");
121        let increment = library
122            .new_function("increment")
123            .expect("increment function");
124
125        let archive_path = common::artifact_path("example-archive.metalarc");
126        let archive = device.new_binary_archive(None).expect("binary archive");
127        archive
128            .add_compute_function(&increment)
129            .expect("archive compute pipeline");
130        archive
131            .add_render_functions(&vertex, &fragment, pixel_format::BGRA8UNORM, 1)
132            .expect("archive render pipeline");
133        archive
134            .serialize_to_file(&archive_path)
135            .expect("serialize binary archive");
136        let _ = device
137            .new_binary_archive(Some(&archive_path))
138            .expect("reload binary archive");
139        println!("binary archive written to {}", archive_path.display());
140    } else {
141        println!("dynamic libraries unsupported; skipping archive serialization");
142    }
143}

Source §

impl MetalFunction

Source

pub const fn as_ptr(&self) -> *mut c_void

Raw id<MTLFunction> pointer.

Examples found in repository ?

examples/04_compute_shader.rs (line 31)

21fn main() {
22    let device = MetalDevice::system_default().expect("MTLCreateSystemDefaultDevice");
23    println!("Device unified={}", device.has_unified_memory());
24
25    let lib = device
26        .new_library_with_source(KERNEL_SRC)
27        .expect("compile MSL source");
28    println!("✅ Compiled library {:p}", lib.as_ptr());
29
30    let func = lib.new_function("mul2").expect("locate function 'mul2'");
31    println!("✅ Found function mul2 {:p}", func.as_ptr());
32
33    let pso = device
34        .new_compute_pipeline_state(&func)
35        .expect("build compute pipeline state");
36    println!("✅ Compute pipeline state {:p}", pso.as_ptr());
37
38    let byte_len = N * core::mem::size_of::<f32>();
39    let buffer = device
40        .new_buffer(byte_len, resource_options::STORAGE_MODE_SHARED)
41        .expect("allocate buffer");
42
43    let slice: &mut [f32] = unsafe {
44        core::slice::from_raw_parts_mut(
45            buffer.contents().expect("buffer.contents").cast::<f32>(),
46            N,
47        )
48    };
49    for (i, x) in slice.iter_mut().enumerate() {
50        *x = i as f32;
51    }
52    println!("Input : {slice:?}");
53
54    let queue = device.new_command_queue().expect("MTLCommandQueue");
55    let cb = queue.new_command_buffer().expect("MTLCommandBuffer");
56    let ok = cb.dispatch_compute_1d(&pso, &[&buffer], N, 1);
57    assert!(ok, "dispatch_compute_1d failed");
58    cb.commit();
59    cb.wait_until_completed();
60
61    let slice: &[f32] = unsafe {
62        core::slice::from_raw_parts(buffer.contents().expect("buffer.contents").cast::<f32>(), N)
63    };
64    println!("Output: {slice:?}");
65
66    for (i, &v) in slice.iter().enumerate() {
67        let expected = (i as f32) * 2.0;
68        assert_eq!(v, expected, "element {i} expected {expected} got {v}");
69    }
70    println!("✅ All {N} elements correctly doubled by the GPU kernel");
71}