1 use mesa_rust_gen::*; 2 use mesa_rust_util::bitset; 3 use mesa_rust_util::offset_of; 4 5 use std::convert::TryInto; 6 use std::ffi::CStr; 7 use std::marker::PhantomData; 8 use std::ops::Not; 9 use std::ptr; 10 use std::ptr::NonNull; 11 use std::slice; 12 13 pub struct ExecListIter<'a, T> { 14 n: &'a mut exec_node, 15 offset: usize, 16 _marker: PhantomData<T>, 17 } 18 19 impl<'a, T> ExecListIter<'a, T> { new(l: &'a mut exec_list, offset: usize) -> Self20 fn new(l: &'a mut exec_list, offset: usize) -> Self { 21 Self { 22 n: &mut l.head_sentinel, 23 offset: offset, 24 _marker: PhantomData, 25 } 26 } 27 } 28 29 impl<'a, T: 'a> Iterator for ExecListIter<'a, T> { 30 type Item = &'a mut T; 31 next(&mut self) -> Option<Self::Item>32 fn next(&mut self) -> Option<Self::Item> { 33 self.n = unsafe { &mut *self.n.next }; 34 if self.n.next.is_null() { 35 None 36 } else { 37 let t: *mut _ = self.n; 38 Some(unsafe { &mut *(t.byte_sub(self.offset).cast()) }) 39 } 40 } 41 } 42 43 #[macro_export] 44 #[cfg(debug_assertions)] 45 macro_rules! nir_pass_impl { 46 ($nir:ident, $pass:ident, $func:ident $(,$arg:expr)* $(,)?) => { 47 { 48 let func_str = ::std::stringify!($func); 49 let func_cstr = ::std::ffi::CString::new(func_str).unwrap(); 50 let res = if unsafe { should_skip_nir(func_cstr.as_ptr()) } { 51 println!("skipping {}", func_str); 52 false 53 } else { 54 $nir.metadata_set_validation_flag(); 55 if $nir.should_print() { 56 println!("{}", func_str); 57 } 58 if $nir.$pass($func $(,$arg)*) { 59 $nir.validate(&format!("after {} in {}:{}", func_str, file!(), line!())); 60 if $nir.should_print() { 61 $nir.print(); 62 } 63 $nir.metadata_check_validation_flag(); 64 true 65 } else { 66 false 67 } 68 }; 69 70 // SAFETY: mutable static can't be read safely, but this value isn't going to change 71 let ndebug = unsafe { nir_debug }; 72 if ndebug & NIR_DEBUG_CLONE != 0 { 73 $nir.validate_clone(); 74 } 75 76 if ndebug & NIR_DEBUG_SERIALIZE != 0 { 77 $nir.validate_serialize_deserialize(); 78 } 79 80 res 81 } 82 }; 83 } 84 85 #[macro_export] 86 #[cfg(not(debug_assertions))] 87 macro_rules! nir_pass_impl { 88 ($nir:ident, $pass:ident, $func:ident $(,$arg:expr)* $(,)?) => { 89 $nir.$pass($func $(,$arg)*) 90 }; 91 } 92 93 #[macro_export] 94 macro_rules! nir_pass { 95 ($nir:ident, $func:ident $(,)?) => { 96 $crate::nir_pass_impl!($nir, pass0, $func) 97 }; 98 99 ($nir:ident, $func:ident, $a:expr $(,)?) => { 100 $crate::nir_pass_impl!($nir, pass1, $func, $a) 101 }; 102 103 ($nir:ident, $func:ident, $a:expr, $b:expr $(,)?) => { 104 $crate::nir_pass_impl!($nir, pass2, $func, $a, $b) 105 }; 106 107 ($nir:ident, $func:ident, $a:expr, $b:expr, $c:expr $(,)?) => { 108 $crate::nir_pass_impl!($nir, pass3, $func, $a, $b, $c) 109 }; 110 } 111 112 pub struct NirPrintfInfo { 113 count: usize, 114 printf_info: *mut u_printf_info, 115 } 116 117 // SAFETY: `u_printf_info` is considered immutable 118 unsafe impl Send for NirPrintfInfo {} 119 unsafe impl Sync for NirPrintfInfo {} 120 121 impl NirPrintfInfo { u_printf(&self, buf: &[u8])122 pub fn u_printf(&self, buf: &[u8]) { 123 unsafe { 124 u_printf( 125 stdout_ptr(), 126 buf.as_ptr().cast(), 127 buf.len(), 128 self.printf_info.cast(), 129 self.count as u32, 130 ); 131 } 132 } 133 } 134 135 impl Drop for NirPrintfInfo { drop(&mut self)136 fn drop(&mut self) { 137 unsafe { 138 ralloc_free(self.printf_info.cast()); 139 }; 140 } 141 } 142 143 pub struct NirShader { 144 nir: NonNull<nir_shader>, 145 } 146 147 // SAFETY: It's safe to share a nir_shader between threads. 148 unsafe impl Send for NirShader {} 149 150 // SAFETY: We do not allow interior mutability with &NirShader 151 unsafe impl Sync for NirShader {} 152 153 impl NirShader { new(nir: *mut nir_shader) -> Option<Self>154 pub fn new(nir: *mut nir_shader) -> Option<Self> { 155 NonNull::new(nir).map(|nir| Self { nir: nir }) 156 } 157 deserialize( blob: &mut blob_reader, options: *const nir_shader_compiler_options, ) -> Option<Self>158 pub fn deserialize( 159 blob: &mut blob_reader, 160 options: *const nir_shader_compiler_options, 161 ) -> Option<Self> { 162 // we already create the NirShader here so it gets automatically deallocated on overrun. 163 let nir = Self::new(unsafe { nir_deserialize(ptr::null_mut(), options, blob) })?; 164 blob.overrun.not().then_some(nir) 165 } 166 serialize(&self, blob: &mut blob)167 pub fn serialize(&self, blob: &mut blob) { 168 unsafe { 169 nir_serialize(blob, self.nir.as_ptr(), false); 170 } 171 } 172 print(&self)173 pub fn print(&self) { 174 unsafe { nir_print_shader(self.nir.as_ptr(), stderr_ptr()) }; 175 } 176 get_nir(&self) -> *mut nir_shader177 pub fn get_nir(&self) -> *mut nir_shader { 178 self.nir.as_ptr() 179 } 180 dup_for_driver(&self) -> *mut nir_shader181 pub fn dup_for_driver(&self) -> *mut nir_shader { 182 unsafe { nir_shader_clone(ptr::null_mut(), self.nir.as_ptr()) } 183 } 184 sweep_mem(&mut self)185 pub fn sweep_mem(&mut self) { 186 unsafe { nir_sweep(self.nir.as_ptr()) } 187 } 188 pass0<R>(&mut self, pass: unsafe extern "C" fn(*mut nir_shader) -> R) -> R189 pub fn pass0<R>(&mut self, pass: unsafe extern "C" fn(*mut nir_shader) -> R) -> R { 190 unsafe { pass(self.nir.as_ptr()) } 191 } 192 pass1<R, A>( &mut self, pass: unsafe extern "C" fn(*mut nir_shader, a: A) -> R, a: A, ) -> R193 pub fn pass1<R, A>( 194 &mut self, 195 pass: unsafe extern "C" fn(*mut nir_shader, a: A) -> R, 196 a: A, 197 ) -> R { 198 unsafe { pass(self.nir.as_ptr(), a) } 199 } 200 pass2<R, A, B>( &mut self, pass: unsafe extern "C" fn(*mut nir_shader, a: A, b: B) -> R, a: A, b: B, ) -> R201 pub fn pass2<R, A, B>( 202 &mut self, 203 pass: unsafe extern "C" fn(*mut nir_shader, a: A, b: B) -> R, 204 a: A, 205 b: B, 206 ) -> R { 207 unsafe { pass(self.nir.as_ptr(), a, b) } 208 } 209 pass3<R, A, B, C>( &mut self, pass: unsafe extern "C" fn(*mut nir_shader, a: A, b: B, c: C) -> R, a: A, b: B, c: C, ) -> R210 pub fn pass3<R, A, B, C>( 211 &mut self, 212 pass: unsafe extern "C" fn(*mut nir_shader, a: A, b: B, c: C) -> R, 213 a: A, 214 b: B, 215 c: C, 216 ) -> R { 217 unsafe { pass(self.nir.as_ptr(), a, b, c) } 218 } 219 220 #[cfg(debug_assertions)] metadata_check_validation_flag(&self)221 pub fn metadata_check_validation_flag(&self) { 222 unsafe { nir_metadata_check_validation_flag(self.nir.as_ptr()) } 223 } 224 225 #[cfg(debug_assertions)] metadata_set_validation_flag(&mut self)226 pub fn metadata_set_validation_flag(&mut self) { 227 unsafe { nir_metadata_set_validation_flag(self.nir.as_ptr()) } 228 } 229 230 #[cfg(debug_assertions)] validate(&self, when: &str)231 pub fn validate(&self, when: &str) { 232 let cstr = std::ffi::CString::new(when).unwrap(); 233 unsafe { nir_validate_shader(self.nir.as_ptr(), cstr.as_ptr()) } 234 } 235 should_print(&self) -> bool236 pub fn should_print(&self) -> bool { 237 unsafe { should_print_nir(self.nir.as_ptr()) } 238 } 239 validate_serialize_deserialize(&mut self)240 pub fn validate_serialize_deserialize(&mut self) { 241 unsafe { nir_shader_serialize_deserialize(self.nir.as_ptr()) } 242 } 243 validate_clone(&mut self)244 pub fn validate_clone(&mut self) { 245 unsafe { 246 let nir_ptr = self.nir.as_ptr(); 247 let clone = nir_shader_clone(ralloc_parent(nir_ptr.cast()), nir_ptr); 248 nir_shader_replace(nir_ptr, clone) 249 } 250 } 251 entrypoint(&self) -> *mut nir_function_impl252 pub fn entrypoint(&self) -> *mut nir_function_impl { 253 unsafe { nir_shader_get_entrypoint(self.nir.as_ptr()) } 254 } 255 structurize(&mut self)256 pub fn structurize(&mut self) { 257 nir_pass!(self, nir_lower_goto_ifs); 258 nir_pass!(self, nir_opt_dead_cf); 259 } 260 inline(&mut self, libclc: &NirShader)261 pub fn inline(&mut self, libclc: &NirShader) { 262 nir_pass!( 263 self, 264 nir_lower_variable_initializers, 265 nir_variable_mode::nir_var_function_temp, 266 ); 267 nir_pass!(self, nir_lower_returns); 268 nir_pass!(self, nir_link_shader_functions, libclc.nir.as_ptr()); 269 nir_pass!(self, nir_inline_functions); 270 } 271 gather_info(&mut self)272 pub fn gather_info(&mut self) { 273 unsafe { nir_shader_gather_info(self.nir.as_ptr(), self.entrypoint()) } 274 } 275 remove_non_entrypoints(&mut self)276 pub fn remove_non_entrypoints(&mut self) { 277 unsafe { nir_remove_non_entrypoints(self.nir.as_ptr()) }; 278 } 279 cleanup_functions(&mut self)280 pub fn cleanup_functions(&mut self) { 281 unsafe { nir_cleanup_functions(self.nir.as_ptr()) }; 282 } 283 variables(&mut self) -> ExecListIter<nir_variable>284 pub fn variables(&mut self) -> ExecListIter<nir_variable> { 285 ExecListIter::new( 286 &mut unsafe { self.nir.as_mut() }.variables, 287 offset_of!(nir_variable, node), 288 ) 289 } 290 num_images(&self) -> u8291 pub fn num_images(&self) -> u8 { 292 unsafe { (*self.nir.as_ptr()).info.num_images } 293 } 294 num_textures(&self) -> u8295 pub fn num_textures(&self) -> u8 { 296 unsafe { (*self.nir.as_ptr()).info.num_textures } 297 } 298 reset_scratch_size(&mut self)299 pub fn reset_scratch_size(&mut self) { 300 unsafe { 301 (*self.nir.as_ptr()).scratch_size = 0; 302 } 303 } 304 scratch_size(&self) -> u32305 pub fn scratch_size(&self) -> u32 { 306 unsafe { (*self.nir.as_ptr()).scratch_size } 307 } 308 reset_shared_size(&mut self)309 pub fn reset_shared_size(&mut self) { 310 unsafe { 311 (*self.nir.as_ptr()).info.shared_size = 0; 312 } 313 } shared_size(&self) -> u32314 pub fn shared_size(&self) -> u32 { 315 unsafe { (*self.nir.as_ptr()).info.shared_size } 316 } 317 workgroup_size(&self) -> [u16; 3]318 pub fn workgroup_size(&self) -> [u16; 3] { 319 unsafe { (*self.nir.as_ptr()).info.workgroup_size } 320 } 321 subgroup_size(&self) -> u8322 pub fn subgroup_size(&self) -> u8 { 323 let subgroup_size = unsafe { (*self.nir.as_ptr()).info.subgroup_size }; 324 let valid_subgroup_sizes = [ 325 gl_subgroup_size::SUBGROUP_SIZE_REQUIRE_8, 326 gl_subgroup_size::SUBGROUP_SIZE_REQUIRE_16, 327 gl_subgroup_size::SUBGROUP_SIZE_REQUIRE_32, 328 gl_subgroup_size::SUBGROUP_SIZE_REQUIRE_64, 329 gl_subgroup_size::SUBGROUP_SIZE_REQUIRE_128, 330 ]; 331 332 if valid_subgroup_sizes.contains(&subgroup_size) { 333 subgroup_size as u8 334 } else { 335 0 336 } 337 } 338 num_subgroups(&self) -> u8339 pub fn num_subgroups(&self) -> u8 { 340 unsafe { (*self.nir.as_ptr()).info.num_subgroups } 341 } 342 set_workgroup_size_variable_if_zero(&mut self)343 pub fn set_workgroup_size_variable_if_zero(&mut self) { 344 let nir = self.nir.as_ptr(); 345 unsafe { 346 (*nir) 347 .info 348 .set_workgroup_size_variable((*nir).info.workgroup_size[0] == 0); 349 } 350 } 351 set_workgroup_size(&mut self, size: [u16; 3])352 pub fn set_workgroup_size(&mut self, size: [u16; 3]) { 353 let nir = unsafe { self.nir.as_mut() }; 354 nir.info.set_workgroup_size_variable(false); 355 nir.info.workgroup_size = size; 356 } 357 workgroup_size_variable(&self) -> bool358 pub fn workgroup_size_variable(&self) -> bool { 359 unsafe { self.nir.as_ref() }.info.workgroup_size_variable() 360 } 361 workgroup_size_hint(&self) -> [u16; 3]362 pub fn workgroup_size_hint(&self) -> [u16; 3] { 363 unsafe { self.nir.as_ref().info.anon_1.cs.workgroup_size_hint } 364 } 365 set_has_variable_shared_mem(&mut self, val: bool)366 pub fn set_has_variable_shared_mem(&mut self, val: bool) { 367 unsafe { 368 self.nir 369 .as_mut() 370 .info 371 .anon_1 372 .cs 373 .set_has_variable_shared_mem(val) 374 } 375 } 376 variables_with_mode( &mut self, mode: nir_variable_mode, ) -> impl Iterator<Item = &mut nir_variable>377 pub fn variables_with_mode( 378 &mut self, 379 mode: nir_variable_mode, 380 ) -> impl Iterator<Item = &mut nir_variable> { 381 self.variables() 382 .filter(move |v| v.data.mode() & mode.0 != 0) 383 } 384 extract_constant_initializers(&mut self)385 pub fn extract_constant_initializers(&mut self) { 386 let nir = self.nir.as_ptr(); 387 unsafe { 388 if (*nir).constant_data_size > 0 { 389 assert!((*nir).constant_data.is_null()); 390 (*nir).constant_data = rzalloc_size(nir.cast(), (*nir).constant_data_size as usize); 391 nir_gather_explicit_io_initializers( 392 nir, 393 (*nir).constant_data, 394 (*nir).constant_data_size as usize, 395 nir_variable_mode::nir_var_mem_constant, 396 ); 397 } 398 } 399 } 400 has_constant(&self) -> bool401 pub fn has_constant(&self) -> bool { 402 unsafe { 403 !self.nir.as_ref().constant_data.is_null() && self.nir.as_ref().constant_data_size > 0 404 } 405 } 406 has_printf(&self) -> bool407 pub fn has_printf(&self) -> bool { 408 unsafe { 409 !self.nir.as_ref().printf_info.is_null() && self.nir.as_ref().printf_info_count != 0 410 } 411 } 412 take_printf_info(&mut self) -> Option<NirPrintfInfo>413 pub fn take_printf_info(&mut self) -> Option<NirPrintfInfo> { 414 let nir = unsafe { self.nir.as_mut() }; 415 416 let info = nir.printf_info; 417 if info.is_null() { 418 return None; 419 } 420 let count = nir.printf_info_count as usize; 421 422 unsafe { 423 ralloc_steal(ptr::null(), info.cast()); 424 425 for i in 0..count { 426 ralloc_steal(info.cast(), (*info.add(i)).arg_sizes.cast()); 427 ralloc_steal(info.cast(), (*info.add(i)).strings.cast()); 428 } 429 }; 430 431 let result = Some(NirPrintfInfo { 432 count: count, 433 printf_info: info, 434 }); 435 436 nir.printf_info_count = 0; 437 nir.printf_info = ptr::null_mut(); 438 439 result 440 } 441 get_constant_buffer(&self) -> &[u8]442 pub fn get_constant_buffer(&self) -> &[u8] { 443 unsafe { 444 let nir = self.nir.as_ref(); 445 // Sometimes, constant_data can be a null pointer if the size is 0 446 if nir.constant_data_size == 0 { 447 &[] 448 } else { 449 slice::from_raw_parts(nir.constant_data.cast(), nir.constant_data_size as usize) 450 } 451 } 452 } 453 preserve_fp16_denorms(&mut self)454 pub fn preserve_fp16_denorms(&mut self) { 455 unsafe { 456 self.nir.as_mut().info.float_controls_execution_mode |= 457 float_controls::FLOAT_CONTROLS_DENORM_PRESERVE_FP16 as u32; 458 } 459 } 460 set_fp_rounding_mode_rtne(&mut self)461 pub fn set_fp_rounding_mode_rtne(&mut self) { 462 unsafe { 463 self.nir.as_mut().info.float_controls_execution_mode |= 464 float_controls::FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16 as u32 465 | float_controls::FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32 as u32 466 | float_controls::FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64 as u32; 467 } 468 } 469 reads_sysval(&self, sysval: gl_system_value) -> bool470 pub fn reads_sysval(&self, sysval: gl_system_value) -> bool { 471 let nir = unsafe { self.nir.as_ref() }; 472 bitset::test_bit(&nir.info.system_values_read, sysval as u32) 473 } 474 add_var( &mut self, mode: nir_variable_mode, glsl_type: *const glsl_type, loc: usize, name: &CStr, )475 pub fn add_var( 476 &mut self, 477 mode: nir_variable_mode, 478 glsl_type: *const glsl_type, 479 loc: usize, 480 name: &CStr, 481 ) { 482 unsafe { 483 let var = nir_variable_create(self.nir.as_ptr(), mode, glsl_type, name.as_ptr()); 484 (*var).data.location = loc.try_into().unwrap(); 485 } 486 } 487 } 488 489 impl Clone for NirShader { clone(&self) -> Self490 fn clone(&self) -> Self { 491 Self { 492 nir: unsafe { NonNull::new_unchecked(self.dup_for_driver()) }, 493 } 494 } 495 } 496 497 impl Drop for NirShader { drop(&mut self)498 fn drop(&mut self) { 499 unsafe { ralloc_free(self.nir.as_ptr().cast()) }; 500 } 501 } 502