1
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
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
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
}
|
