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 28)

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

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    {
44        let mut mapping = unsafe { buffer.map_write().expect("map compute input") };
45        for (i, bytes) in mapping.chunks_exact_mut(4).take(N).enumerate() {
46            bytes.copy_from_slice(&(i as f32).to_ne_bytes());
47        }
48        let input: Vec<f32> = mapping
49            .chunks_exact(4)
50            .take(N)
51            .map(|bytes| f32::from_ne_bytes(bytes.try_into().expect("four-byte float")))
52            .collect();
53        drop(mapping);
54        println!("Input : {input:?}");
55    }
56
57    let queue = device.new_command_queue().expect("MTLCommandQueue");
58    let cb = queue.new_command_buffer().expect("MTLCommandBuffer");
59    cb.dispatch_compute_1d(&pso, &[&buffer], N, 1)
60        .expect("dispatch compute");
61    cb.commit().expect("commit compute");
62    cb.wait_until_completed().expect("complete compute");
63
64    let mapping = unsafe { buffer.map_read().expect("map compute output") };
65    let output: Vec<f32> = mapping
66        .chunks_exact(4)
67        .take(N)
68        .map(|bytes| f32::from_ne_bytes(bytes.try_into().expect("four-byte float")))
69        .collect();
70    drop(mapping);
71    println!("Output: {output:?}");
72
73    for (i, &v) in output.iter().enumerate() {
74        let expected = (i as f32) * 2.0;
75        assert_eq!(v, expected, "element {i} expected {expected} got {v}");
76    }
77    println!("✅ All {N} elements correctly doubled by the GPU kernel");
78}