Struct ComputePipelineState

Source

pub struct ComputePipelineState { /* private fields */ }

Expand description

Apple’s id<MTLComputePipelineState> — a compiled compute kernel.

Implementations§

Source §

impl ComputePipelineState

Source

pub fn thread_execution_width(&self) -> usize

Thread execution width for this compute pipeline.

Source

pub fn max_total_threads_per_threadgroup(&self) -> usize

Maximum threads per threadgroup.

Source

pub fn new_visible_function_table( &self, function_count: usize, ) -> Option<VisibleFunctionTable>

Allocate a visible function table for this pipeline.

Examples found in repository ?

examples/07_advanced_objects.rs (line 101)

11fn main() {
12    let device = MetalDevice::system_default().expect("Metal device available");
13    let queue = device.new_command_queue().expect("command queue");
14    let counter_sets = device.counter_set_names();
15    println!("counter sets: {counter_sets:?}");
16
17    if let Some(event) = device.new_shared_event() {
18        event.set_signaled_value(1);
19        println!("event signaled value={}", event.signaled_value());
20        let signal = queue
21            .new_command_buffer()
22            .expect("event signal command buffer");
23        signal.encode_signal_event(&event, 2);
24        signal.commit();
25        signal.wait_until_completed();
26        println!(
27            "event reached value 2: {}",
28            event.wait_until_signaled_value(2, 1_000),
29        );
30
31        let wait = queue
32            .new_command_buffer()
33            .expect("event wait command buffer");
34        wait.encode_wait_for_event(&event, 2);
35        wait.commit();
36        wait.wait_until_completed();
37    }
38
39    let fence_a = device.new_fence();
40    let fence_b = device.new_fence();
41    let sample_buffer = counter_sets.first().and_then(|name| {
42        if device.supports_counter_sampling(counter_sampling_point::AT_BLIT_BOUNDARY) {
43            device
44                .new_counter_sample_buffer(name, 2, storage_mode::SHARED, Some("example-samples"))
45                .ok()
46        } else {
47            None
48        }
49    });
50
51    let src = device
52        .new_buffer(64, resource_options::STORAGE_MODE_SHARED)
53        .expect("source buffer");
54    let dst = device
55        .new_buffer(64, resource_options::STORAGE_MODE_SHARED)
56        .expect("destination buffer");
57    let blit = queue.new_command_buffer().expect("blit command buffer");
58    let encoder = blit.new_blit_command_encoder().expect("blit encoder");
59    let _ = encoder.fill_buffer(&src, 0..64, b'Q');
60    if let Some(fence) = fence_a.as_ref() {
61        encoder.update_fence(fence);
62    }
63    encoder.end_encoding();
64    blit.commit();
65    blit.wait_until_completed();
66
67    let blit = queue
68        .new_command_buffer()
69        .expect("second blit command buffer");
70    let encoder = blit
71        .new_blit_command_encoder()
72        .expect("second blit encoder");
73    if let Some(fence) = fence_a.as_ref() {
74        encoder.wait_for_fence(fence);
75    }
76    if let Some(sample_buffer) = sample_buffer.as_ref() {
77        let _ = encoder.sample_counters(sample_buffer, 0, false);
78    }
79    let _ = encoder.copy_buffer(&src, 0, &dst, 0, 64);
80    encoder.end_encoding();
81    blit.commit();
82    blit.wait_until_completed();
83    if let Some(sample_buffer) = sample_buffer.as_ref() {
84        println!(
85            "resolved counter bytes={}",
86            sample_buffer
87                .resolve_range(0..1)
88                .map_or(0, |bytes| bytes.len())
89        );
90    }
91
92    let library = device
93        .new_library_with_source(common::COMPUTE_SRC)
94        .expect("compile compute library");
95    let increment = library
96        .new_function("increment")
97        .expect("increment function");
98    let pipeline = device
99        .new_compute_pipeline_state(&increment)
100        .expect("compute pipeline");
101    let visible_table = pipeline.new_visible_function_table(1);
102    let intersection_table = if device.supports_raytracing() {
103        pipeline.new_intersection_function_table(1)
104    } else {
105        None
106    };
107    if let Some(table) = intersection_table.as_ref() {
108        table.set_opaque_triangle_intersection_function(intersection_function_signature::NONE, 0);
109    }
110    let acceleration_structure = if device.supports_raytracing() {
111        device.new_acceleration_structure_with_size(256)
112    } else {
113        None
114    };
115
116    let buffer = device
117        .new_buffer(16, resource_options::STORAGE_MODE_SHARED)
118        .expect("compute buffer");
119    common::write_u32_words(&buffer, &[1, 2, 3, 4]);
120    let texture = device
121        .new_texture(apple_metal::TextureDescriptor::new_2d(
122            4,
123            4,
124            apple_metal::pixel_format::BGRA8UNORM,
125        ))
126        .expect("compute texture");
127    let compute = queue.new_command_buffer().expect("compute command buffer");
128    let encoder = compute
129        .new_compute_command_encoder()
130        .expect("compute command encoder");
131    encoder.set_compute_pipeline_state(&pipeline);
132    encoder.set_buffer(&buffer, 0, 0);
133    encoder.set_texture(&texture, 1);
134    if let Some(fence) = fence_a.as_ref() {
135        encoder.wait_for_fence(fence);
136    }
137    if let Some(table) = visible_table.as_ref() {
138        encoder.set_visible_function_table(table, 2);
139    }
140    if let Some(table) = intersection_table.as_ref() {
141        encoder.set_intersection_function_table(table, 3);
142    }
143    if let Some(acceleration_structure) = acceleration_structure.as_ref() {
144        encoder.set_acceleration_structure(acceleration_structure, 4);
145    }
146    encoder.dispatch_threadgroups((1, 1, 1), (4, 1, 1));
147    if let Some(fence) = fence_b.as_ref() {
148        encoder.update_fence(fence);
149    }
150    encoder.end_encoding();
151    compute.commit();
152    compute.wait_until_completed();
153    println!(
154        "compute buffer after dispatch: {:?}",
155        common::read_u32_words(&buffer, 4)
156    );
157
158    if let Some(indirect) = device.new_indirect_command_buffer(
159        indirect_command_type::CONCURRENT_DISPATCH,
160        1,
161        0,
162        0,
163        4,
164        resource_options::STORAGE_MODE_PRIVATE,
165    ) {
166        indirect.reset_range(0..1);
167        println!("indirect command buffer size={}", indirect.size());
168    }
169
170    if let Some(heap) = device.new_heap(1 << 20, storage_mode::SHARED) {
171        if let Ok(residency_set) = device.new_residency_set(Some("example-residency"), 4) {
172            let heap_buffer = heap
173                .new_buffer(256, resource_options::STORAGE_MODE_SHARED)
174                .expect("heap buffer");
175            residency_set.add_buffer(&heap_buffer);
176            residency_set.add_heap(&heap);
177            residency_set.commit();
178            residency_set.request_residency();
179            queue.add_residency_set(&residency_set);
180            queue.remove_residency_set(&residency_set);
181            residency_set.end_residency();
182            residency_set.remove_all_allocations();
183            residency_set.commit();
184            println!(
185                "residency allocation count={}",
186                residency_set.allocation_count()
187            );
188        } else {
189            println!("residency sets unavailable on this OS");
190        }
191    }
192
193    if let Some(capture_manager) = CaptureManager::shared() {
194        println!(
195            "capture supported for developer tools={} active={}",
196            capture_manager.supports_destination(capture_destination::DEVELOPER_TOOLS),
197            capture_manager.is_capturing(),
198        );
199        if let Some(scope) = capture_manager.new_capture_scope_with_device(&device) {
200            scope.begin();
201            scope.end();
202        }
203        if let Some(scope) = capture_manager.new_capture_scope_with_command_queue(&queue) {
204            scope.begin();
205            scope.end();
206        }
207    }
208}

Source

pub fn new_intersection_function_table( &self, function_count: usize, ) -> Option<IntersectionFunctionTable>

Allocate an intersection function table for this pipeline.

Examples found in repository ?

examples/07_advanced_objects.rs (line 103)

11fn main() {
12    let device = MetalDevice::system_default().expect("Metal device available");
13    let queue = device.new_command_queue().expect("command queue");
14    let counter_sets = device.counter_set_names();
15    println!("counter sets: {counter_sets:?}");
16
17    if let Some(event) = device.new_shared_event() {
18        event.set_signaled_value(1);
19        println!("event signaled value={}", event.signaled_value());
20        let signal = queue
21            .new_command_buffer()
22            .expect("event signal command buffer");
23        signal.encode_signal_event(&event, 2);
24        signal.commit();
25        signal.wait_until_completed();
26        println!(
27            "event reached value 2: {}",
28            event.wait_until_signaled_value(2, 1_000),
29        );
30
31        let wait = queue
32            .new_command_buffer()
33            .expect("event wait command buffer");
34        wait.encode_wait_for_event(&event, 2);
35        wait.commit();
36        wait.wait_until_completed();
37    }
38
39    let fence_a = device.new_fence();
40    let fence_b = device.new_fence();
41    let sample_buffer = counter_sets.first().and_then(|name| {
42        if device.supports_counter_sampling(counter_sampling_point::AT_BLIT_BOUNDARY) {
43            device
44                .new_counter_sample_buffer(name, 2, storage_mode::SHARED, Some("example-samples"))
45                .ok()
46        } else {
47            None
48        }
49    });
50
51    let src = device
52        .new_buffer(64, resource_options::STORAGE_MODE_SHARED)
53        .expect("source buffer");
54    let dst = device
55        .new_buffer(64, resource_options::STORAGE_MODE_SHARED)
56        .expect("destination buffer");
57    let blit = queue.new_command_buffer().expect("blit command buffer");
58    let encoder = blit.new_blit_command_encoder().expect("blit encoder");
59    let _ = encoder.fill_buffer(&src, 0..64, b'Q');
60    if let Some(fence) = fence_a.as_ref() {
61        encoder.update_fence(fence);
62    }
63    encoder.end_encoding();
64    blit.commit();
65    blit.wait_until_completed();
66
67    let blit = queue
68        .new_command_buffer()
69        .expect("second blit command buffer");
70    let encoder = blit
71        .new_blit_command_encoder()
72        .expect("second blit encoder");
73    if let Some(fence) = fence_a.as_ref() {
74        encoder.wait_for_fence(fence);
75    }
76    if let Some(sample_buffer) = sample_buffer.as_ref() {
77        let _ = encoder.sample_counters(sample_buffer, 0, false);
78    }
79    let _ = encoder.copy_buffer(&src, 0, &dst, 0, 64);
80    encoder.end_encoding();
81    blit.commit();
82    blit.wait_until_completed();
83    if let Some(sample_buffer) = sample_buffer.as_ref() {
84        println!(
85            "resolved counter bytes={}",
86            sample_buffer
87                .resolve_range(0..1)
88                .map_or(0, |bytes| bytes.len())
89        );
90    }
91
92    let library = device
93        .new_library_with_source(common::COMPUTE_SRC)
94        .expect("compile compute library");
95    let increment = library
96        .new_function("increment")
97        .expect("increment function");
98    let pipeline = device
99        .new_compute_pipeline_state(&increment)
100        .expect("compute pipeline");
101    let visible_table = pipeline.new_visible_function_table(1);
102    let intersection_table = if device.supports_raytracing() {
103        pipeline.new_intersection_function_table(1)
104    } else {
105        None
106    };
107    if let Some(table) = intersection_table.as_ref() {
108        table.set_opaque_triangle_intersection_function(intersection_function_signature::NONE, 0);
109    }
110    let acceleration_structure = if device.supports_raytracing() {
111        device.new_acceleration_structure_with_size(256)
112    } else {
113        None
114    };
115
116    let buffer = device
117        .new_buffer(16, resource_options::STORAGE_MODE_SHARED)
118        .expect("compute buffer");
119    common::write_u32_words(&buffer, &[1, 2, 3, 4]);
120    let texture = device
121        .new_texture(apple_metal::TextureDescriptor::new_2d(
122            4,
123            4,
124            apple_metal::pixel_format::BGRA8UNORM,
125        ))
126        .expect("compute texture");
127    let compute = queue.new_command_buffer().expect("compute command buffer");
128    let encoder = compute
129        .new_compute_command_encoder()
130        .expect("compute command encoder");
131    encoder.set_compute_pipeline_state(&pipeline);
132    encoder.set_buffer(&buffer, 0, 0);
133    encoder.set_texture(&texture, 1);
134    if let Some(fence) = fence_a.as_ref() {
135        encoder.wait_for_fence(fence);
136    }
137    if let Some(table) = visible_table.as_ref() {
138        encoder.set_visible_function_table(table, 2);
139    }
140    if let Some(table) = intersection_table.as_ref() {
141        encoder.set_intersection_function_table(table, 3);
142    }
143    if let Some(acceleration_structure) = acceleration_structure.as_ref() {
144        encoder.set_acceleration_structure(acceleration_structure, 4);
145    }
146    encoder.dispatch_threadgroups((1, 1, 1), (4, 1, 1));
147    if let Some(fence) = fence_b.as_ref() {
148        encoder.update_fence(fence);
149    }
150    encoder.end_encoding();
151    compute.commit();
152    compute.wait_until_completed();
153    println!(
154        "compute buffer after dispatch: {:?}",
155        common::read_u32_words(&buffer, 4)
156    );
157
158    if let Some(indirect) = device.new_indirect_command_buffer(
159        indirect_command_type::CONCURRENT_DISPATCH,
160        1,
161        0,
162        0,
163        4,
164        resource_options::STORAGE_MODE_PRIVATE,
165    ) {
166        indirect.reset_range(0..1);
167        println!("indirect command buffer size={}", indirect.size());
168    }
169
170    if let Some(heap) = device.new_heap(1 << 20, storage_mode::SHARED) {
171        if let Ok(residency_set) = device.new_residency_set(Some("example-residency"), 4) {
172            let heap_buffer = heap
173                .new_buffer(256, resource_options::STORAGE_MODE_SHARED)
174                .expect("heap buffer");
175            residency_set.add_buffer(&heap_buffer);
176            residency_set.add_heap(&heap);
177            residency_set.commit();
178            residency_set.request_residency();
179            queue.add_residency_set(&residency_set);
180            queue.remove_residency_set(&residency_set);
181            residency_set.end_residency();
182            residency_set.remove_all_allocations();
183            residency_set.commit();
184            println!(
185                "residency allocation count={}",
186                residency_set.allocation_count()
187            );
188        } else {
189            println!("residency sets unavailable on this OS");
190        }
191    }
192
193    if let Some(capture_manager) = CaptureManager::shared() {
194        println!(
195            "capture supported for developer tools={} active={}",
196            capture_manager.supports_destination(capture_destination::DEVELOPER_TOOLS),
197            capture_manager.is_capturing(),
198        );
199        if let Some(scope) = capture_manager.new_capture_scope_with_device(&device) {
200            scope.begin();
201            scope.end();
202        }
203        if let Some(scope) = capture_manager.new_capture_scope_with_command_queue(&queue) {
204            scope.begin();
205            scope.end();
206        }
207    }
208}

Source §

impl ComputePipelineState

Source

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

Raw id<MTLComputePipelineState> pointer.

Examples found in repository ?

examples/04_compute_shader.rs (line 36)

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}