1 // Copyright © 2022 Collabora, Ltd.
2 // SPDX-License-Identifier: MIT
3
4 #![allow(non_upper_case_globals)]
5
6 use crate::api::GetDebugFlags;
7 use crate::api::DEBUG;
8 use crate::builder::*;
9 use crate::ir::*;
10 use crate::sph::{OutputTopology, PixelImap};
11
12 use nak_bindings::*;
13
14 use compiler::bindings::*;
15 use compiler::cfg::CFGBuilder;
16 use compiler::nir::*;
17 use compiler::nir_instr_printer::NirInstrPrinter;
18 use std::cmp::max;
19 use std::collections::{HashMap, HashSet};
20 use std::ops::Index;
21
init_info_from_nir(nak: &nak_compiler, nir: &nir_shader) -> ShaderInfo22 fn init_info_from_nir(nak: &nak_compiler, nir: &nir_shader) -> ShaderInfo {
23 ShaderInfo {
24 num_gprs: 0,
25 num_instrs: 0,
26 num_control_barriers: 0,
27 slm_size: nir.scratch_size,
28 max_crs_depth: 0,
29 uses_global_mem: false,
30 writes_global_mem: false,
31 // TODO: handle this.
32 uses_fp64: false,
33 stage: match nir.info.stage() {
34 MESA_SHADER_COMPUTE => {
35 ShaderStageInfo::Compute(ComputeShaderInfo {
36 local_size: [
37 nir.info.workgroup_size[0],
38 nir.info.workgroup_size[1],
39 nir.info.workgroup_size[2],
40 ],
41 smem_size: nir.info.shared_size.try_into().unwrap(),
42 })
43 }
44 MESA_SHADER_VERTEX => ShaderStageInfo::Vertex,
45 MESA_SHADER_FRAGMENT => {
46 let info_fs = unsafe { &nir.info.__bindgen_anon_1.fs };
47 ShaderStageInfo::Fragment(FragmentShaderInfo {
48 uses_kill: false,
49 does_interlock: false,
50 post_depth_coverage: info_fs.post_depth_coverage(),
51 early_fragment_tests: info_fs.early_fragment_tests(),
52 uses_sample_shading: info_fs.uses_sample_shading(),
53 })
54 }
55 MESA_SHADER_GEOMETRY => {
56 let info_gs = unsafe { &nir.info.__bindgen_anon_1.gs };
57 let output_topology = match info_gs.output_primitive {
58 MESA_PRIM_POINTS => OutputTopology::PointList,
59 MESA_PRIM_LINE_STRIP => OutputTopology::LineStrip,
60 MESA_PRIM_TRIANGLE_STRIP => OutputTopology::TriangleStrip,
61 _ => panic!(
62 "Invalid GS input primitive {}",
63 info_gs.input_primitive
64 ),
65 };
66
67 ShaderStageInfo::Geometry(GeometryShaderInfo {
68 // TODO: Should be set if VK_NV_geometry_shader_passthrough is in use.
69 passthrough_enable: false,
70 stream_out_mask: info_gs.active_stream_mask(),
71 threads_per_input_primitive: info_gs.invocations,
72 output_topology: output_topology,
73 max_output_vertex_count: info_gs.vertices_out,
74 })
75 }
76 MESA_SHADER_TESS_CTRL => {
77 let info_tess = unsafe { &nir.info.__bindgen_anon_1.tess };
78 ShaderStageInfo::TessellationInit(TessellationInitShaderInfo {
79 per_patch_attribute_count: 6,
80 threads_per_patch: info_tess.tcs_vertices_out,
81 })
82 }
83 MESA_SHADER_TESS_EVAL => {
84 let info_tess = unsafe { &nir.info.__bindgen_anon_1.tess };
85 ShaderStageInfo::Tessellation(TessellationShaderInfo {
86 domain: match info_tess._primitive_mode {
87 TESS_PRIMITIVE_TRIANGLES => {
88 TessellationDomain::Triangle
89 }
90 TESS_PRIMITIVE_QUADS => TessellationDomain::Quad,
91 TESS_PRIMITIVE_ISOLINES => TessellationDomain::Isoline,
92 _ => panic!("Invalid tess_primitive_mode"),
93 },
94 spacing: match info_tess.spacing() {
95 TESS_SPACING_EQUAL => TessellationSpacing::Integer,
96 TESS_SPACING_FRACTIONAL_ODD => {
97 TessellationSpacing::FractionalOdd
98 }
99 TESS_SPACING_FRACTIONAL_EVEN => {
100 TessellationSpacing::FractionalEven
101 }
102 _ => panic!("Invalid gl_tess_spacing"),
103 },
104 primitives: if info_tess.point_mode() {
105 TessellationPrimitives::Points
106 } else if info_tess._primitive_mode
107 == TESS_PRIMITIVE_ISOLINES
108 {
109 TessellationPrimitives::Lines
110 } else if info_tess.ccw() {
111 TessellationPrimitives::TrianglesCCW
112 } else {
113 TessellationPrimitives::TrianglesCW
114 },
115 })
116 }
117 _ => panic!("Unknown shader stage"),
118 },
119 io: match nir.info.stage() {
120 MESA_SHADER_COMPUTE => ShaderIoInfo::None,
121 MESA_SHADER_FRAGMENT => ShaderIoInfo::Fragment(FragmentIoInfo {
122 sysvals_in: SysValInfo {
123 // Required on fragment shaders, otherwise it cause a trap.
124 ab: 1 << 31,
125 c: 0,
126 },
127 sysvals_in_d: [PixelImap::Unused; 8],
128 attr_in: [PixelImap::Unused; 128],
129 barycentric_attr_in: [0; 4],
130 reads_sample_mask: false,
131 writes_color: 0,
132 writes_sample_mask: false,
133 writes_depth: false,
134 }),
135 MESA_SHADER_VERTEX
136 | MESA_SHADER_GEOMETRY
137 | MESA_SHADER_TESS_CTRL
138 | MESA_SHADER_TESS_EVAL => {
139 let num_clip = nir.info.clip_distance_array_size();
140 let num_cull = nir.info.cull_distance_array_size();
141 let clip_enable = (1_u32 << num_clip) - 1;
142 let cull_enable = ((1_u32 << num_cull) - 1) << num_clip;
143
144 ShaderIoInfo::Vtg(VtgIoInfo {
145 sysvals_in: SysValInfo::default(),
146 sysvals_in_d: 0,
147 sysvals_out: SysValInfo::default(),
148 sysvals_out_d: 0,
149 attr_in: [0; 4],
150 attr_out: [0; 4],
151
152 // TODO: figure out how to fill this.
153 store_req_start: u8::MAX,
154 store_req_end: 0,
155
156 clip_enable: clip_enable.try_into().unwrap(),
157 cull_enable: cull_enable.try_into().unwrap(),
158 xfb: if nir.xfb_info.is_null() {
159 None
160 } else {
161 Some(Box::new(unsafe {
162 nak_xfb_from_nir(nak, nir.xfb_info)
163 }))
164 },
165 })
166 }
167 _ => panic!("Unknown shader stage"),
168 },
169 }
170 }
171
alloc_ssa_for_nir(b: &mut impl SSABuilder, ssa: &nir_def) -> Vec<SSAValue>172 fn alloc_ssa_for_nir(b: &mut impl SSABuilder, ssa: &nir_def) -> Vec<SSAValue> {
173 let (file, comps) = if ssa.bit_size == 1 {
174 (RegFile::Pred, ssa.num_components)
175 } else {
176 let bits = ssa.bit_size * ssa.num_components;
177 (RegFile::GPR, bits.div_ceil(32))
178 };
179
180 let mut vec = Vec::new();
181 for _ in 0..comps {
182 vec.push(b.alloc_ssa(file, 1)[0]);
183 }
184 vec
185 }
186
187 struct PhiAllocMap<'a> {
188 alloc: &'a mut PhiAllocator,
189 map: HashMap<(u32, u8), u32>,
190 }
191
192 impl<'a> PhiAllocMap<'a> {
new(alloc: &'a mut PhiAllocator) -> PhiAllocMap<'a>193 fn new(alloc: &'a mut PhiAllocator) -> PhiAllocMap<'a> {
194 PhiAllocMap {
195 alloc: alloc,
196 map: HashMap::new(),
197 }
198 }
199
get_phi_id(&mut self, phi: &nir_phi_instr, comp: u8) -> u32200 fn get_phi_id(&mut self, phi: &nir_phi_instr, comp: u8) -> u32 {
201 *self
202 .map
203 .entry((phi.def.index, comp))
204 .or_insert_with(|| self.alloc.alloc())
205 }
206 }
207
208 struct PerSizeFloatControls {
209 pub ftz: bool,
210 pub rnd_mode: FRndMode,
211 }
212
213 struct ShaderFloatControls {
214 pub fp16: PerSizeFloatControls,
215 pub fp32: PerSizeFloatControls,
216 pub fp64: PerSizeFloatControls,
217 }
218
219 impl Default for ShaderFloatControls {
default() -> Self220 fn default() -> Self {
221 Self {
222 fp16: PerSizeFloatControls {
223 ftz: false,
224 rnd_mode: FRndMode::NearestEven,
225 },
226 fp32: PerSizeFloatControls {
227 ftz: true, // Default FTZ on fp32
228 rnd_mode: FRndMode::NearestEven,
229 },
230 fp64: PerSizeFloatControls {
231 ftz: false,
232 rnd_mode: FRndMode::NearestEven,
233 },
234 }
235 }
236 }
237
238 impl ShaderFloatControls {
from_nir(nir: &nir_shader) -> ShaderFloatControls239 fn from_nir(nir: &nir_shader) -> ShaderFloatControls {
240 let nir_fc = nir.info.float_controls_execution_mode;
241 let mut fc: ShaderFloatControls = Default::default();
242
243 if (nir_fc & FLOAT_CONTROLS_DENORM_PRESERVE_FP16) != 0 {
244 fc.fp16.ftz = false;
245 } else if (nir_fc & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16) != 0 {
246 fc.fp16.ftz = true;
247 }
248 if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16) != 0 {
249 fc.fp16.rnd_mode = FRndMode::NearestEven;
250 } else if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16) != 0 {
251 fc.fp16.rnd_mode = FRndMode::Zero;
252 }
253
254 if (nir_fc & FLOAT_CONTROLS_DENORM_PRESERVE_FP32) != 0 {
255 fc.fp32.ftz = false;
256 } else if (nir_fc & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32) != 0 {
257 fc.fp32.ftz = true;
258 }
259 if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32) != 0 {
260 fc.fp32.rnd_mode = FRndMode::NearestEven;
261 } else if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32) != 0 {
262 fc.fp32.rnd_mode = FRndMode::Zero;
263 }
264
265 if (nir_fc & FLOAT_CONTROLS_DENORM_PRESERVE_FP64) != 0 {
266 fc.fp64.ftz = false;
267 } else if (nir_fc & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP64) != 0 {
268 fc.fp64.ftz = true;
269 }
270 if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64) != 0 {
271 fc.fp64.rnd_mode = FRndMode::NearestEven;
272 } else if (nir_fc & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64) != 0 {
273 fc.fp64.rnd_mode = FRndMode::Zero;
274 }
275
276 fc
277 }
278 }
279
280 impl Index<FloatType> for ShaderFloatControls {
281 type Output = PerSizeFloatControls;
282
index(&self, idx: FloatType) -> &PerSizeFloatControls283 fn index(&self, idx: FloatType) -> &PerSizeFloatControls {
284 match idx {
285 FloatType::F16 => &self.fp16,
286 FloatType::F32 => &self.fp32,
287 FloatType::F64 => &self.fp64,
288 }
289 }
290 }
291
292 #[derive(Clone, Copy, Eq, Hash, PartialEq)]
293 enum SyncType {
294 Sync,
295 Brk,
296 Cont,
297 }
298
299 struct ShaderFromNir<'a> {
300 nir: &'a nir_shader,
301 sm: &'a dyn ShaderModel,
302 info: ShaderInfo,
303 float_ctl: ShaderFloatControls,
304 cfg: CFGBuilder<u32, BasicBlock>,
305 label_alloc: LabelAllocator,
306 block_label: HashMap<u32, Label>,
307 bar_label: HashMap<u32, Label>,
308 sync_blocks: HashSet<u32>,
309 crs: Vec<(u32, SyncType)>,
310 fs_out_regs: [SSAValue; 34],
311 end_block_id: u32,
312 ssa_map: HashMap<u32, Vec<SSAValue>>,
313 saturated: HashSet<*const nir_def>,
314 nir_instr_printer: NirInstrPrinter,
315 }
316
317 impl<'a> ShaderFromNir<'a> {
new( nak: &nak_compiler, nir: &'a nir_shader, sm: &'a dyn ShaderModel, ) -> Self318 fn new(
319 nak: &nak_compiler,
320 nir: &'a nir_shader,
321 sm: &'a dyn ShaderModel,
322 ) -> Self {
323 Self {
324 nir: nir,
325 sm: sm,
326 info: init_info_from_nir(nak, nir),
327 float_ctl: ShaderFloatControls::from_nir(nir),
328 cfg: CFGBuilder::new(),
329 label_alloc: LabelAllocator::new(),
330 block_label: HashMap::new(),
331 bar_label: HashMap::new(),
332 sync_blocks: HashSet::new(),
333 crs: Vec::new(),
334 fs_out_regs: [SSAValue::NONE; 34],
335 end_block_id: 0,
336 ssa_map: HashMap::new(),
337 saturated: HashSet::new(),
338 nir_instr_printer: NirInstrPrinter::new().unwrap(),
339 }
340 }
341
get_block_label(&mut self, block: &nir_block) -> Label342 fn get_block_label(&mut self, block: &nir_block) -> Label {
343 *self
344 .block_label
345 .entry(block.index)
346 .or_insert_with(|| self.label_alloc.alloc())
347 }
348
push_crs(&mut self, target: &nir_block, sync_type: SyncType)349 fn push_crs(&mut self, target: &nir_block, sync_type: SyncType) {
350 self.sync_blocks.insert(target.index);
351 self.crs.push((target.index, sync_type));
352 let crs_depth = u32::try_from(self.crs.len()).unwrap();
353 self.info.max_crs_depth = max(self.info.max_crs_depth, crs_depth);
354 }
355
pop_crs(&mut self, target: &nir_block, sync_type: SyncType)356 fn pop_crs(&mut self, target: &nir_block, sync_type: SyncType) {
357 if let Some((top_index, top_sync_type)) = self.crs.pop() {
358 assert!(top_index == target.index);
359 assert!(top_sync_type == sync_type);
360 } else {
361 panic!("Tried to pop an empty stack");
362 }
363 }
364
peek_crs(&self, target: &nir_block) -> Option<SyncType>365 fn peek_crs(&self, target: &nir_block) -> Option<SyncType> {
366 for (i, (index, sync_type)) in self.crs.iter().enumerate().rev() {
367 if *index != target.index {
368 continue;
369 }
370
371 match sync_type {
372 SyncType::Sync => {
373 // Sync must always be top-of-stack
374 assert!(i == self.crs.len() - 1);
375 }
376 SyncType::Brk => {
377 // Brk cannot skip over another Brk
378 for (_, inner_sync) in &self.crs[(i + 1)..] {
379 assert!(*inner_sync != SyncType::Brk);
380 }
381 }
382 SyncType::Cont => {
383 // Cont can only skip over Sync
384 for (_, inner_sync) in &self.crs[(i + 1)..] {
385 assert!(*inner_sync == SyncType::Sync);
386 }
387 }
388 }
389
390 return Some(*sync_type);
391 }
392
393 assert!(!self.sync_blocks.contains(&target.index));
394 None
395 }
396
get_ssa(&mut self, ssa: &nir_def) -> &[SSAValue]397 fn get_ssa(&mut self, ssa: &nir_def) -> &[SSAValue] {
398 self.ssa_map.get(&ssa.index).unwrap()
399 }
400
set_ssa(&mut self, def: &nir_def, vec: Vec<SSAValue>)401 fn set_ssa(&mut self, def: &nir_def, vec: Vec<SSAValue>) {
402 if def.bit_size == 1 {
403 for s in &vec {
404 assert!(s.is_predicate());
405 }
406 } else {
407 for s in &vec {
408 assert!(!s.is_predicate());
409 }
410 let bits =
411 usize::from(def.bit_size) * usize::from(def.num_components);
412 assert!(vec.len() == bits.div_ceil(32));
413 }
414 self.ssa_map
415 .entry(def.index)
416 .and_modify(|_| panic!("Cannot set an SSA def twice"))
417 .or_insert(vec);
418 }
419
get_ssa_comp(&mut self, def: &nir_def, c: u8) -> (SSARef, u8)420 fn get_ssa_comp(&mut self, def: &nir_def, c: u8) -> (SSARef, u8) {
421 let vec = self.get_ssa(def);
422 match def.bit_size {
423 1 => (vec[usize::from(c)].into(), 0),
424 8 => (vec[usize::from(c / 4)].into(), c % 4),
425 16 => (vec[usize::from(c / 2)].into(), (c * 2) % 4),
426 32 => (vec[usize::from(c)].into(), 0),
427 64 => {
428 let comps =
429 [vec[usize::from(c) * 2 + 0], vec[usize::from(c) * 2 + 1]];
430 (comps.into(), 0)
431 }
432 _ => panic!("Unsupported bit size: {}", def.bit_size),
433 }
434 }
435
get_ssa_ref(&mut self, src: &nir_src) -> SSARef436 fn get_ssa_ref(&mut self, src: &nir_src) -> SSARef {
437 SSARef::try_from(self.get_ssa(src.as_def())).unwrap()
438 }
439
get_src(&mut self, src: &nir_src) -> Src440 fn get_src(&mut self, src: &nir_src) -> Src {
441 self.get_ssa_ref(src).into()
442 }
443
get_io_addr_offset( &mut self, addr: &nir_src, imm_bits: u8, ) -> (Src, i32)444 fn get_io_addr_offset(
445 &mut self,
446 addr: &nir_src,
447 imm_bits: u8,
448 ) -> (Src, i32) {
449 let addr = addr.as_def();
450 let addr_offset = unsafe {
451 nak_get_io_addr_offset(addr as *const _ as *mut _, imm_bits)
452 };
453
454 if let Some(base_def) = std::ptr::NonNull::new(addr_offset.base.def) {
455 let base_def = unsafe { base_def.as_ref() };
456 let base_comp = u8::try_from(addr_offset.base.comp).unwrap();
457 let (base, _) = self.get_ssa_comp(base_def, base_comp);
458 (base.into(), addr_offset.offset)
459 } else {
460 (SrcRef::Zero.into(), addr_offset.offset)
461 }
462 }
463
get_cbuf_addr_offset(&mut self, addr: &nir_src) -> (Src, u16)464 fn get_cbuf_addr_offset(&mut self, addr: &nir_src) -> (Src, u16) {
465 let (off, off_imm) = self.get_io_addr_offset(addr, 16);
466 if let Ok(off_imm_u16) = u16::try_from(off_imm) {
467 (off, off_imm_u16)
468 } else {
469 (self.get_src(addr), 0)
470 }
471 }
472
set_dst(&mut self, def: &nir_def, ssa: SSARef)473 fn set_dst(&mut self, def: &nir_def, ssa: SSARef) {
474 self.set_ssa(def, (*ssa).into());
475 }
476
try_saturate_alu_dst(&mut self, def: &nir_def) -> bool477 fn try_saturate_alu_dst(&mut self, def: &nir_def) -> bool {
478 if def.all_uses_are_fsat() {
479 self.saturated.insert(def as *const _);
480 true
481 } else {
482 false
483 }
484 }
485
alu_src_is_saturated(&self, src: &nir_alu_src) -> bool486 fn alu_src_is_saturated(&self, src: &nir_alu_src) -> bool {
487 self.saturated
488 .get(&(src.src.as_def() as *const _))
489 .is_some()
490 }
491
parse_alu(&mut self, b: &mut impl SSABuilder, alu: &nir_alu_instr)492 fn parse_alu(&mut self, b: &mut impl SSABuilder, alu: &nir_alu_instr) {
493 // Handle vectors and pack ops as a special case since they're the only
494 // ALU ops that can produce more than 16B. They are also the only ALU
495 // ops which we allow to consume small (8 and 16-bit) vector data
496 // scattered across multiple dwords
497 match alu.op {
498 nir_op_mov
499 | nir_op_pack_32_4x8_split
500 | nir_op_pack_32_2x16_split
501 | nir_op_pack_64_2x32_split
502 | nir_op_vec2
503 | nir_op_vec3
504 | nir_op_vec4
505 | nir_op_vec5
506 | nir_op_vec8
507 | nir_op_vec16 => {
508 let src_bit_size = alu.get_src(0).src.bit_size();
509 let bits = usize::from(alu.def.num_components)
510 * usize::from(alu.def.bit_size);
511
512 // Collect the sources into a vec with src_bit_size per SSA
513 // value in the vec. This implicitly makes 64-bit sources look
514 // like two 32-bit values
515 let mut srcs = Vec::new();
516 if alu.op == nir_op_mov {
517 let src = alu.get_src(0);
518 for c in 0..alu.def.num_components {
519 let s = src.swizzle[usize::from(c)];
520 let (src, byte) =
521 self.get_ssa_comp(src.src.as_def(), s);
522 for ssa in src.iter() {
523 srcs.push((*ssa, byte));
524 }
525 }
526 } else {
527 for src in alu.srcs_as_slice().iter() {
528 let s = src.swizzle[0];
529 let (src, byte) =
530 self.get_ssa_comp(src.src.as_def(), s);
531 for ssa in src.iter() {
532 srcs.push((*ssa, byte));
533 }
534 }
535 }
536
537 let mut comps = Vec::new();
538 match src_bit_size {
539 1 | 32 | 64 => {
540 for (ssa, _) in srcs {
541 comps.push(ssa);
542 }
543 }
544 8 => {
545 for dc in 0..bits.div_ceil(32) {
546 let mut psrc = [None; 4];
547 let mut psel = [0_u8; 4];
548
549 for b in 0..4 {
550 let sc = dc * 4 + b;
551 if sc < srcs.len() {
552 let (ssa, byte) = srcs[sc];
553 // Deduplicate psrc entries
554 for i in 0..4_u8 {
555 let psrc_i = &mut psrc[usize::from(i)];
556 if psrc_i.is_none() {
557 *psrc_i = Some(ssa.into());
558 } else if *psrc_i
559 != Some(Src::from(ssa))
560 {
561 continue;
562 }
563 psel[b] = i * 4 + byte;
564 break;
565 }
566 }
567 }
568
569 let psrc = {
570 let mut res = [Src::new_zero(); 4];
571
572 for (idx, src) in psrc.iter().enumerate() {
573 if let Some(src) = src {
574 res[idx] = *src;
575 }
576 }
577
578 res
579 };
580
581 comps.push(b.prmt4(psrc, psel)[0]);
582 }
583 }
584 16 => {
585 for dc in 0..bits.div_ceil(32) {
586 let mut psrc = [Src::new_zero(); 2];
587 let mut psel = [0_u8; 4];
588
589 for w in 0..2 {
590 let sc = dc * 2 + w;
591 if sc < srcs.len() {
592 let (ssa, byte) = srcs[sc];
593 psrc[w] = ssa.into();
594 psel[w * 2 + 0] = (w as u8 * 4) + byte;
595 psel[w * 2 + 1] = (w as u8 * 4) + byte + 1;
596 }
597 }
598 comps.push(b.prmt(psrc[0], psrc[1], psel)[0]);
599 }
600 }
601 _ => panic!("Unknown bit size: {src_bit_size}"),
602 }
603
604 self.set_ssa(&alu.def, comps);
605 return;
606 }
607 _ => (),
608 }
609
610 let nir_srcs = alu.srcs_as_slice();
611 let mut srcs: Vec<Src> = Vec::new();
612 for (i, alu_src) in nir_srcs.iter().enumerate() {
613 let bit_size = alu_src.src.bit_size();
614 let comps = alu.src_components(i.try_into().unwrap());
615 let ssa = self.get_ssa(alu_src.src.as_def());
616
617 match bit_size {
618 1 => {
619 assert!(comps == 1);
620 let s = usize::from(alu_src.swizzle[0]);
621 srcs.push(ssa[s].into());
622 }
623 8 | 16 => {
624 let num_bytes = usize::from(comps * (bit_size / 8));
625 assert!(num_bytes <= 4);
626
627 let mut bytes = [0_u8; 4];
628 for c in 0..usize::from(comps) {
629 let cs = alu_src.swizzle[c];
630 if bit_size == 8 {
631 bytes[c] = cs;
632 } else {
633 bytes[c * 2 + 0] = cs * 2 + 0;
634 bytes[c * 2 + 1] = cs * 2 + 1;
635 }
636 }
637
638 let mut prmt_srcs = [Src::new_zero(); 4];
639 let mut prmt = [0_u8; 4];
640 for b in 0..num_bytes {
641 for (ds, s) in prmt_srcs.iter_mut().enumerate() {
642 let dw = ssa[usize::from(bytes[b] / 4)];
643 if s.is_zero() {
644 *s = dw.into();
645 } else if *s != Src::from(dw) {
646 continue;
647 }
648 prmt[usize::from(b)] =
649 (ds as u8) * 4 + (bytes[b] % 4);
650 break;
651 }
652 }
653
654 srcs.push(b.prmt4(prmt_srcs, prmt).into());
655 }
656 32 => {
657 assert!(comps == 1);
658 let s = usize::from(alu_src.swizzle[0]);
659 srcs.push(ssa[s].into());
660 }
661 64 => {
662 assert!(comps == 1);
663 let s = usize::from(alu_src.swizzle[0]);
664 srcs.push([ssa[s * 2], ssa[s * 2 + 1]].into());
665 }
666 _ => panic!("Invalid bit size: {bit_size}"),
667 }
668 }
669
670 // Restricts an F16v2 source to just x if the ALU op is single-component. This
671 // must only be called for per-component sources (see nir_op_info::output_sizes
672 // for more details).
673 let restrict_f16v2_src = |mut src: Src| {
674 if alu.def.num_components == 1 {
675 src.src_swizzle = SrcSwizzle::Xx;
676 }
677 src
678 };
679
680 let dst: SSARef = match alu.op {
681 nir_op_b2b1 => {
682 assert!(alu.get_src(0).bit_size() == 32);
683 b.isetp(IntCmpType::I32, IntCmpOp::Ne, srcs[0], 0.into())
684 }
685 nir_op_b2b32 | nir_op_b2i8 | nir_op_b2i16 | nir_op_b2i32 => {
686 b.sel(srcs[0].bnot(), 0.into(), 1.into())
687 }
688 nir_op_b2i64 => {
689 let lo = b.sel(srcs[0].bnot(), 0.into(), 1.into());
690 let hi = b.copy(0.into());
691 [lo[0], hi[0]].into()
692 }
693 nir_op_b2f16 => b.sel(srcs[0].bnot(), 0.into(), 0x3c00.into()),
694 nir_op_b2f32 => {
695 b.sel(srcs[0].bnot(), 0.0_f32.into(), 1.0_f32.into())
696 }
697 nir_op_b2f64 => {
698 let lo = b.copy(0.into());
699 let hi = b.sel(srcs[0].bnot(), 0.into(), 0x3ff00000.into());
700 [lo[0], hi[0]].into()
701 }
702 nir_op_bcsel => b.sel(srcs[0], srcs[1], srcs[2]),
703 nir_op_bfm => {
704 let dst = b.alloc_ssa(RegFile::GPR, 1);
705 b.push_op(OpBMsk {
706 dst: dst.into(),
707 pos: srcs[1],
708 width: srcs[0],
709 wrap: true,
710 });
711 dst
712 }
713 nir_op_bit_count => {
714 let dst = b.alloc_ssa(RegFile::GPR, 1);
715 b.push_op(OpPopC {
716 dst: dst.into(),
717 src: srcs[0],
718 });
719 dst
720 }
721 nir_op_bitfield_reverse => b.brev(srcs[0]),
722 nir_op_ibitfield_extract | nir_op_ubitfield_extract => {
723 let range = b.alloc_ssa(RegFile::GPR, 1);
724 b.push_op(OpPrmt {
725 dst: range.into(),
726 srcs: [srcs[1], srcs[2]],
727 sel: 0x0040.into(),
728 mode: PrmtMode::Index,
729 });
730
731 let dst = b.alloc_ssa(RegFile::GPR, 1);
732 b.push_op(OpBfe {
733 dst: dst.into(),
734 base: srcs[0],
735 signed: !matches!(alu.op, nir_op_ubitfield_extract),
736 range: range.into(),
737 reverse: false,
738 });
739 dst
740 }
741 nir_op_extract_u8 | nir_op_extract_i8 | nir_op_extract_u16
742 | nir_op_extract_i16 => {
743 let src1 = alu.get_src(1);
744 let elem = src1.src.comp_as_uint(src1.swizzle[0]).unwrap();
745 let elem = u8::try_from(elem).unwrap();
746
747 match alu.op {
748 nir_op_extract_u8 => {
749 assert!(elem < 4);
750 let byte = elem;
751 let zero = 4;
752 b.prmt(srcs[0], 0.into(), [byte, zero, zero, zero])
753 }
754 nir_op_extract_i8 => {
755 assert!(elem < 4);
756 let byte = elem;
757 let sign = byte | 0x8;
758 b.prmt(srcs[0], 0.into(), [byte, sign, sign, sign])
759 }
760 nir_op_extract_u16 => {
761 assert!(elem < 2);
762 let byte = elem * 2;
763 let zero = 4;
764 b.prmt(srcs[0], 0.into(), [byte, byte + 1, zero, zero])
765 }
766 nir_op_extract_i16 => {
767 assert!(elem < 2);
768 let byte = elem * 2;
769 let sign = (byte + 1) | 0x8;
770 b.prmt(srcs[0], 0.into(), [byte, byte + 1, sign, sign])
771 }
772 _ => panic!("Unknown extract op: {}", alu.op),
773 }
774 }
775 nir_op_f2f16 | nir_op_f2f16_rtne | nir_op_f2f16_rtz
776 | nir_op_f2f32 | nir_op_f2f64 => {
777 let src_bits = alu.get_src(0).src.bit_size();
778 let dst_bits = alu.def.bit_size();
779 let src_type = FloatType::from_bits(src_bits.into());
780 let dst_type = FloatType::from_bits(dst_bits.into());
781
782 let dst = b.alloc_ssa(RegFile::GPR, dst_bits.div_ceil(32));
783 b.push_op(OpF2F {
784 dst: dst.into(),
785 src: srcs[0],
786 src_type: FloatType::from_bits(src_bits.into()),
787 dst_type: dst_type,
788 rnd_mode: match alu.op {
789 nir_op_f2f16_rtne => FRndMode::NearestEven,
790 nir_op_f2f16_rtz => FRndMode::Zero,
791 _ => self.float_ctl[dst_type].rnd_mode,
792 },
793 ftz: if src_bits < dst_bits {
794 self.float_ctl[src_type].ftz
795 } else {
796 self.float_ctl[dst_type].ftz
797 },
798 high: false,
799 integer_rnd: false,
800 });
801 dst
802 }
803 nir_op_find_lsb => {
804 let rev = b.brev(srcs[0]);
805 let dst = b.alloc_ssa(RegFile::GPR, 1);
806 b.push_op(OpFlo {
807 dst: dst.into(),
808 src: rev.into(),
809 signed: false,
810 return_shift_amount: true,
811 });
812 dst
813 }
814 nir_op_f2i8 | nir_op_f2i16 | nir_op_f2i32 | nir_op_f2i64
815 | nir_op_f2u8 | nir_op_f2u16 | nir_op_f2u32 | nir_op_f2u64 => {
816 let src_bits = usize::from(alu.get_src(0).bit_size());
817 let dst_bits = alu.def.bit_size();
818 let src_type = FloatType::from_bits(src_bits);
819 let dst = b.alloc_ssa(RegFile::GPR, dst_bits.div_ceil(32));
820 let dst_is_signed = alu.info().output_type & 2 != 0;
821 let dst_type =
822 IntType::from_bits(dst_bits.into(), dst_is_signed);
823 if b.sm() < 70 && dst_bits == 8 {
824 // F2I doesn't support 8-bit destinations pre-Volta
825 let tmp = b.alloc_ssa(RegFile::GPR, 1);
826 let tmp_type = IntType::from_bits(32, dst_is_signed);
827 b.push_op(OpF2I {
828 dst: tmp.into(),
829 src: srcs[0],
830 src_type,
831 dst_type: tmp_type,
832 rnd_mode: FRndMode::Zero,
833 ftz: self.float_ctl[src_type].ftz,
834 });
835 b.push_op(OpI2I {
836 dst: dst.into(),
837 src: tmp.into(),
838 src_type: tmp_type,
839 dst_type,
840 saturate: true,
841 abs: false,
842 neg: false,
843 });
844 } else {
845 b.push_op(OpF2I {
846 dst: dst.into(),
847 src: srcs[0],
848 src_type,
849 dst_type,
850 rnd_mode: FRndMode::Zero,
851 ftz: self.float_ctl[src_type].ftz,
852 });
853 }
854 dst
855 }
856 nir_op_fabs | nir_op_fadd | nir_op_fneg => {
857 let (x, y) = match alu.op {
858 nir_op_fabs => (Src::new_zero().fneg(), srcs[0].fabs()),
859 nir_op_fadd => (srcs[0], srcs[1]),
860 nir_op_fneg => (Src::new_zero().fneg(), srcs[0].fneg()),
861 _ => panic!("Unhandled case"),
862 };
863 let ftype = FloatType::from_bits(alu.def.bit_size().into());
864 let dst;
865 if alu.def.bit_size() == 64 {
866 dst = b.alloc_ssa(RegFile::GPR, 2);
867 b.push_op(OpDAdd {
868 dst: dst.into(),
869 srcs: [x, y],
870 rnd_mode: self.float_ctl[ftype].rnd_mode,
871 });
872 } else if alu.def.bit_size() == 32 {
873 dst = b.alloc_ssa(RegFile::GPR, 1);
874 b.push_op(OpFAdd {
875 dst: dst.into(),
876 srcs: [x, y],
877 saturate: self.try_saturate_alu_dst(&alu.def),
878 rnd_mode: self.float_ctl[ftype].rnd_mode,
879 ftz: self.float_ctl[ftype].ftz,
880 });
881 } else if alu.def.bit_size() == 16 {
882 assert!(
883 self.float_ctl[ftype].rnd_mode == FRndMode::NearestEven
884 );
885
886 dst = b.alloc_ssa(RegFile::GPR, 1);
887 b.push_op(OpHAdd2 {
888 dst: dst.into(),
889 srcs: [restrict_f16v2_src(x), restrict_f16v2_src(y)],
890 saturate: self.try_saturate_alu_dst(&alu.def),
891 ftz: self.float_ctl[ftype].ftz,
892 f32: false,
893 });
894 } else {
895 panic!("Unsupported float type: f{}", alu.def.bit_size());
896 }
897 dst
898 }
899 nir_op_fceil | nir_op_ffloor | nir_op_fround_even
900 | nir_op_ftrunc => {
901 let dst = b.alloc_ssa(RegFile::GPR, 1);
902 let ty = FloatType::from_bits(alu.def.bit_size().into());
903 let rnd_mode = match alu.op {
904 nir_op_fceil => FRndMode::PosInf,
905 nir_op_ffloor => FRndMode::NegInf,
906 nir_op_ftrunc => FRndMode::Zero,
907 nir_op_fround_even => FRndMode::NearestEven,
908 _ => unreachable!(),
909 };
910 let ftz = self.float_ctl[ty].ftz;
911 if b.sm() >= 70 {
912 assert!(
913 alu.def.bit_size() == 32 || alu.def.bit_size() == 16
914 );
915 b.push_op(OpFRnd {
916 dst: dst.into(),
917 src: srcs[0],
918 src_type: ty,
919 dst_type: ty,
920 rnd_mode,
921 ftz,
922 });
923 } else {
924 assert!(alu.def.bit_size() == 32);
925 b.push_op(OpF2F {
926 dst: dst.into(),
927 src: srcs[0],
928 src_type: ty,
929 dst_type: ty,
930 rnd_mode,
931 ftz,
932 integer_rnd: true,
933 high: false,
934 });
935 }
936 dst
937 }
938 nir_op_fcos => b.fcos(srcs[0]),
939 nir_op_feq | nir_op_fge | nir_op_flt | nir_op_fneu => {
940 let src_type =
941 FloatType::from_bits(alu.get_src(0).bit_size().into());
942 let cmp_op = match alu.op {
943 nir_op_feq => FloatCmpOp::OrdEq,
944 nir_op_fge => FloatCmpOp::OrdGe,
945 nir_op_flt => FloatCmpOp::OrdLt,
946 nir_op_fneu => FloatCmpOp::UnordNe,
947 _ => panic!("Usupported float comparison"),
948 };
949
950 let dst = b.alloc_ssa(RegFile::Pred, alu.def.num_components);
951 if alu.get_src(0).bit_size() == 64 {
952 assert!(alu.def.num_components == 1);
953 b.push_op(OpDSetP {
954 dst: dst.into(),
955 set_op: PredSetOp::And,
956 cmp_op: cmp_op,
957 srcs: [srcs[0], srcs[1]],
958 accum: SrcRef::True.into(),
959 });
960 } else if alu.get_src(0).bit_size() == 32 {
961 assert!(alu.def.num_components == 1);
962 b.push_op(OpFSetP {
963 dst: dst.into(),
964 set_op: PredSetOp::And,
965 cmp_op: cmp_op,
966 srcs: [srcs[0], srcs[1]],
967 accum: SrcRef::True.into(),
968 ftz: self.float_ctl[src_type].ftz,
969 });
970 } else if alu.get_src(0).bit_size() == 16 {
971 assert!(
972 alu.def.num_components == 1
973 || alu.def.num_components == 2
974 );
975
976 let dsts = if alu.def.num_components == 2 {
977 [dst[0].into(), dst[1].into()]
978 } else {
979 [dst[0].into(), Dst::None]
980 };
981
982 b.push_op(OpHSetP2 {
983 dsts,
984 set_op: PredSetOp::And,
985 cmp_op: cmp_op,
986 srcs: [
987 restrict_f16v2_src(srcs[0]),
988 restrict_f16v2_src(srcs[1]),
989 ],
990 accum: SrcRef::True.into(),
991 ftz: self.float_ctl[src_type].ftz,
992 horizontal: false,
993 });
994 } else {
995 panic!(
996 "Unsupported float type: f{}",
997 alu.get_src(0).bit_size()
998 );
999 }
1000 dst
1001 }
1002 nir_op_fexp2 => b.fexp2(srcs[0]),
1003 nir_op_ffma => {
1004 let ftype = FloatType::from_bits(alu.def.bit_size().into());
1005 let dst;
1006 if alu.def.bit_size() == 64 {
1007 debug_assert!(!self.float_ctl[ftype].ftz);
1008 dst = b.alloc_ssa(RegFile::GPR, 2);
1009 b.push_op(OpDFma {
1010 dst: dst.into(),
1011 srcs: [srcs[0], srcs[1], srcs[2]],
1012 rnd_mode: self.float_ctl[ftype].rnd_mode,
1013 });
1014 } else if alu.def.bit_size() == 32 {
1015 dst = b.alloc_ssa(RegFile::GPR, 1);
1016 b.push_op(OpFFma {
1017 dst: dst.into(),
1018 srcs: [srcs[0], srcs[1], srcs[2]],
1019 saturate: self.try_saturate_alu_dst(&alu.def),
1020 rnd_mode: self.float_ctl[ftype].rnd_mode,
1021 // The hardware doesn't like FTZ+DNZ and DNZ implies FTZ
1022 // anyway so only set one of the two bits.
1023 ftz: self.float_ctl[ftype].ftz,
1024 dnz: false,
1025 });
1026 } else if alu.def.bit_size() == 16 {
1027 assert!(
1028 self.float_ctl[ftype].rnd_mode == FRndMode::NearestEven
1029 );
1030
1031 dst = b.alloc_ssa(RegFile::GPR, 1);
1032 b.push_op(OpHFma2 {
1033 dst: dst.into(),
1034 srcs: [
1035 restrict_f16v2_src(srcs[0]),
1036 restrict_f16v2_src(srcs[1]),
1037 restrict_f16v2_src(srcs[2]),
1038 ],
1039 saturate: self.try_saturate_alu_dst(&alu.def),
1040 ftz: self.float_ctl[ftype].ftz,
1041 dnz: false,
1042 f32: false,
1043 });
1044 } else {
1045 panic!("Unsupported float type: f{}", alu.def.bit_size());
1046 }
1047 dst
1048 }
1049 nir_op_ffmaz => {
1050 assert!(alu.def.bit_size() == 32);
1051 // DNZ implies FTZ so we need FTZ set or this is invalid
1052 assert!(self.float_ctl.fp32.ftz);
1053 let dst = b.alloc_ssa(RegFile::GPR, 1);
1054 b.push_op(OpFFma {
1055 dst: dst.into(),
1056 srcs: [srcs[0], srcs[1], srcs[2]],
1057 saturate: self.try_saturate_alu_dst(&alu.def),
1058 rnd_mode: self.float_ctl.fp32.rnd_mode,
1059 // The hardware doesn't like FTZ+DNZ and DNZ implies FTZ
1060 // anyway so only set one of the two bits.
1061 ftz: false,
1062 dnz: true,
1063 });
1064 dst
1065 }
1066 nir_op_flog2 => {
1067 assert!(alu.def.bit_size() == 32);
1068 b.mufu(MuFuOp::Log2, srcs[0])
1069 }
1070 nir_op_fmax | nir_op_fmin => {
1071 let dst;
1072 if alu.def.bit_size() == 64 {
1073 dst = b.alloc_ssa(RegFile::GPR, 2);
1074 b.push_op(OpDMnMx {
1075 dst: dst.into(),
1076 srcs: [srcs[0], srcs[1]],
1077 min: (alu.op == nir_op_fmin).into(),
1078 });
1079 } else if alu.def.bit_size() == 32 {
1080 dst = b.alloc_ssa(RegFile::GPR, 1);
1081 b.push_op(OpFMnMx {
1082 dst: dst.into(),
1083 srcs: [srcs[0], srcs[1]],
1084 min: (alu.op == nir_op_fmin).into(),
1085 ftz: self.float_ctl.fp32.ftz,
1086 });
1087 } else if alu.def.bit_size() == 16 {
1088 dst = b.alloc_ssa(RegFile::GPR, 1);
1089 b.push_op(OpHMnMx2 {
1090 dst: dst.into(),
1091 srcs: [
1092 restrict_f16v2_src(srcs[0]),
1093 restrict_f16v2_src(srcs[1]),
1094 ],
1095 min: (alu.op == nir_op_fmin).into(),
1096 ftz: self.float_ctl.fp16.ftz,
1097 });
1098 } else {
1099 panic!("Unsupported float type: f{}", alu.def.bit_size());
1100 }
1101 dst
1102 }
1103 nir_op_fmul => {
1104 let ftype = FloatType::from_bits(alu.def.bit_size().into());
1105 let dst;
1106 if alu.def.bit_size() == 64 {
1107 debug_assert!(!self.float_ctl[ftype].ftz);
1108 dst = b.alloc_ssa(RegFile::GPR, 2);
1109 b.push_op(OpDMul {
1110 dst: dst.into(),
1111 srcs: [srcs[0], srcs[1]],
1112 rnd_mode: self.float_ctl[ftype].rnd_mode,
1113 });
1114 } else if alu.def.bit_size() == 32 {
1115 dst = b.alloc_ssa(RegFile::GPR, 1);
1116 b.push_op(OpFMul {
1117 dst: dst.into(),
1118 srcs: [srcs[0], srcs[1]],
1119 saturate: self.try_saturate_alu_dst(&alu.def),
1120 rnd_mode: self.float_ctl[ftype].rnd_mode,
1121 ftz: self.float_ctl[ftype].ftz,
1122 dnz: false,
1123 });
1124 } else if alu.def.bit_size() == 16 {
1125 assert!(
1126 self.float_ctl[ftype].rnd_mode == FRndMode::NearestEven
1127 );
1128
1129 dst = b.alloc_ssa(RegFile::GPR, 1);
1130 b.push_op(OpHMul2 {
1131 dst: dst.into(),
1132 srcs: [
1133 restrict_f16v2_src(srcs[0]),
1134 restrict_f16v2_src(srcs[1]),
1135 ],
1136 saturate: self.try_saturate_alu_dst(&alu.def),
1137 ftz: self.float_ctl[ftype].ftz,
1138 dnz: false,
1139 });
1140 } else {
1141 panic!("Unsupported float type: f{}", alu.def.bit_size());
1142 }
1143 dst
1144 }
1145 nir_op_fmulz => {
1146 assert!(alu.def.bit_size() == 32);
1147 // DNZ implies FTZ so we need FTZ set or this is invalid
1148 assert!(self.float_ctl.fp32.ftz);
1149 let dst = b.alloc_ssa(RegFile::GPR, 1);
1150 b.push_op(OpFMul {
1151 dst: dst.into(),
1152 srcs: [srcs[0], srcs[1]],
1153 saturate: self.try_saturate_alu_dst(&alu.def),
1154 rnd_mode: self.float_ctl.fp32.rnd_mode,
1155 // The hardware doesn't like FTZ+DNZ and DNZ implies FTZ
1156 // anyway so only set one of the two bits.
1157 ftz: false,
1158 dnz: true,
1159 });
1160 dst
1161 }
1162 nir_op_fquantize2f16 => {
1163 let tmp = b.alloc_ssa(RegFile::GPR, 1);
1164 b.push_op(OpF2F {
1165 dst: tmp.into(),
1166 src: srcs[0],
1167 src_type: FloatType::F32,
1168 dst_type: FloatType::F16,
1169 rnd_mode: FRndMode::NearestEven,
1170 ftz: true,
1171 high: false,
1172 integer_rnd: false,
1173 });
1174 assert!(alu.def.bit_size() == 32);
1175 let dst = b.alloc_ssa(RegFile::GPR, 1);
1176 b.push_op(OpF2F {
1177 dst: dst.into(),
1178 src: tmp.into(),
1179 src_type: FloatType::F16,
1180 dst_type: FloatType::F32,
1181 rnd_mode: FRndMode::NearestEven,
1182 ftz: true,
1183 high: false,
1184 integer_rnd: false,
1185 });
1186 if b.sm() < 70 {
1187 // Pre-Volta, F2F.ftz doesn't flush denorms so we need to do
1188 // that manually
1189 let denorm = b.fsetp(
1190 FloatCmpOp::OrdLt,
1191 srcs[0].fabs(),
1192 0x38800000.into(),
1193 );
1194 // Get the correctly signed zero
1195 let zero =
1196 b.lop2(LogicOp2::And, srcs[0], 0x80000000.into());
1197 b.sel(denorm.into(), zero.into(), dst.into())
1198 } else {
1199 dst
1200 }
1201 }
1202 nir_op_frcp => {
1203 assert!(alu.def.bit_size() == 32);
1204 b.mufu(MuFuOp::Rcp, srcs[0])
1205 }
1206 nir_op_frsq => {
1207 assert!(alu.def.bit_size() == 32);
1208 b.mufu(MuFuOp::Rsq, srcs[0])
1209 }
1210 nir_op_fsat => {
1211 let ftype = FloatType::from_bits(alu.def.bit_size().into());
1212
1213 if self.alu_src_is_saturated(&alu.srcs_as_slice()[0]) {
1214 b.copy(srcs[0])
1215 } else if alu.def.bit_size() == 32 {
1216 let dst = b.alloc_ssa(RegFile::GPR, 1);
1217 b.push_op(OpFAdd {
1218 dst: dst.into(),
1219 srcs: [srcs[0], 0.into()],
1220 saturate: true,
1221 rnd_mode: self.float_ctl[ftype].rnd_mode,
1222 ftz: self.float_ctl[ftype].ftz,
1223 });
1224 dst
1225 } else if alu.def.bit_size() == 16 {
1226 assert!(
1227 self.float_ctl[ftype].rnd_mode == FRndMode::NearestEven
1228 );
1229
1230 let dst = b.alloc_ssa(RegFile::GPR, 1);
1231 b.push_op(OpHAdd2 {
1232 dst: dst.into(),
1233 srcs: [restrict_f16v2_src(srcs[0]), 0.into()],
1234 saturate: true,
1235 ftz: self.float_ctl[ftype].ftz,
1236 f32: false,
1237 });
1238 dst
1239 } else {
1240 panic!("Unsupported float type: f{}", alu.def.bit_size());
1241 }
1242 }
1243 nir_op_fsign => {
1244 if alu.def.bit_size() == 64 {
1245 let lz = b.dsetp(FloatCmpOp::OrdLt, srcs[0], 0.into());
1246 let gz = b.dsetp(FloatCmpOp::OrdGt, srcs[0], 0.into());
1247 let hi = b.sel(lz.into(), 0xbff00000.into(), 0.into());
1248 let hi = b.sel(gz.into(), 0x3ff00000.into(), hi.into());
1249 let lo = b.copy(0.into());
1250 [lo[0], hi[0]].into()
1251 } else if alu.def.bit_size() == 32 {
1252 let lz = b.fset(FloatCmpOp::OrdLt, srcs[0], 0.into());
1253 let gz = b.fset(FloatCmpOp::OrdGt, srcs[0], 0.into());
1254 b.fadd(gz.into(), Src::from(lz).fneg())
1255 } else if alu.def.bit_size() == 16 {
1256 let x = restrict_f16v2_src(srcs[0]);
1257
1258 let lz = restrict_f16v2_src(
1259 b.hset2(FloatCmpOp::OrdLt, x, 0.into()).into(),
1260 );
1261 let gz = restrict_f16v2_src(
1262 b.hset2(FloatCmpOp::OrdGt, x, 0.into()).into(),
1263 );
1264
1265 b.hadd2(gz, lz.fneg())
1266 } else {
1267 panic!("Unsupported float type: f{}", alu.def.bit_size());
1268 }
1269 }
1270 nir_op_fsin => b.fsin(srcs[0]),
1271 nir_op_fsqrt => b.mufu(MuFuOp::Sqrt, srcs[0]),
1272 nir_op_i2f16 | nir_op_i2f32 | nir_op_i2f64 => {
1273 let src_bits = alu.get_src(0).src.bit_size();
1274 let dst_bits = alu.def.bit_size();
1275 let dst_type = FloatType::from_bits(dst_bits.into());
1276 let dst = b.alloc_ssa(RegFile::GPR, dst_bits.div_ceil(32));
1277 b.push_op(OpI2F {
1278 dst: dst.into(),
1279 src: srcs[0],
1280 dst_type: dst_type,
1281 src_type: IntType::from_bits(src_bits.into(), true),
1282 rnd_mode: self.float_ctl[dst_type].rnd_mode,
1283 });
1284 dst
1285 }
1286 nir_op_i2i8 | nir_op_i2i16 | nir_op_i2i32 | nir_op_i2i64
1287 | nir_op_u2u8 | nir_op_u2u16 | nir_op_u2u32 | nir_op_u2u64 => {
1288 let src_bits = alu.get_src(0).src.bit_size();
1289 let dst_bits = alu.def.bit_size();
1290
1291 let mut prmt = [0_u8; 8];
1292 match alu.op {
1293 nir_op_i2i8 | nir_op_i2i16 | nir_op_i2i32
1294 | nir_op_i2i64 => {
1295 let sign = ((src_bits / 8) - 1) | 0x8;
1296 for i in 0..8 {
1297 if i < (src_bits / 8) {
1298 prmt[usize::from(i)] = i;
1299 } else {
1300 prmt[usize::from(i)] = sign;
1301 }
1302 }
1303 }
1304 nir_op_u2u8 | nir_op_u2u16 | nir_op_u2u32
1305 | nir_op_u2u64 => {
1306 for i in 0..8 {
1307 if i < (src_bits / 8) {
1308 prmt[usize::from(i)] = i;
1309 } else {
1310 prmt[usize::from(i)] = 4;
1311 }
1312 }
1313 }
1314 _ => panic!("Invalid integer conversion: {}", alu.op),
1315 }
1316 let prmt_lo: [u8; 4] = prmt[0..4].try_into().unwrap();
1317 let prmt_hi: [u8; 4] = prmt[4..8].try_into().unwrap();
1318
1319 let src = srcs[0].as_ssa().unwrap();
1320 if src_bits == 64 {
1321 if dst_bits == 64 {
1322 *src
1323 } else {
1324 b.prmt(src[0].into(), src[1].into(), prmt_lo)
1325 }
1326 } else {
1327 if dst_bits == 64 {
1328 let lo = b.prmt(src[0].into(), 0.into(), prmt_lo);
1329 let hi = b.prmt(src[0].into(), 0.into(), prmt_hi);
1330 [lo[0], hi[0]].into()
1331 } else {
1332 b.prmt(src[0].into(), 0.into(), prmt_lo)
1333 }
1334 }
1335 }
1336 nir_op_iabs => b.iabs(srcs[0]),
1337 nir_op_iadd => match alu.def.bit_size {
1338 32 => b.iadd(srcs[0], srcs[1], 0.into()),
1339 64 => b.iadd64(srcs[0], srcs[1], 0.into()),
1340 x => panic!("unsupported bit size for nir_op_iadd: {x}"),
1341 },
1342 nir_op_iadd3 => match alu.def.bit_size {
1343 32 => b.iadd(srcs[0], srcs[1], srcs[2]),
1344 64 => b.iadd64(srcs[0], srcs[1], srcs[2]),
1345 x => panic!("unsupported bit size for nir_op_iadd3: {x}"),
1346 },
1347 nir_op_iand => b.lop2(LogicOp2::And, srcs[0], srcs[1]),
1348 nir_op_ieq => {
1349 if alu.get_src(0).bit_size() == 1 {
1350 b.lop2(LogicOp2::Xor, srcs[0], srcs[1].bnot())
1351 } else if alu.get_src(0).bit_size() == 64 {
1352 b.isetp64(IntCmpType::I32, IntCmpOp::Eq, srcs[0], srcs[1])
1353 } else {
1354 assert!(alu.get_src(0).bit_size() == 32);
1355 b.isetp(IntCmpType::I32, IntCmpOp::Eq, srcs[0], srcs[1])
1356 }
1357 }
1358 nir_op_ifind_msb | nir_op_ifind_msb_rev | nir_op_ufind_msb
1359 | nir_op_ufind_msb_rev => {
1360 let dst = b.alloc_ssa(RegFile::GPR, 1);
1361 b.push_op(OpFlo {
1362 dst: dst.into(),
1363 src: srcs[0],
1364 signed: match alu.op {
1365 nir_op_ifind_msb | nir_op_ifind_msb_rev => true,
1366 nir_op_ufind_msb | nir_op_ufind_msb_rev => false,
1367 _ => panic!("Not a find_msb op"),
1368 },
1369 return_shift_amount: match alu.op {
1370 nir_op_ifind_msb | nir_op_ufind_msb => false,
1371 nir_op_ifind_msb_rev | nir_op_ufind_msb_rev => true,
1372 _ => panic!("Not a find_msb op"),
1373 },
1374 });
1375 dst
1376 }
1377 nir_op_ige | nir_op_ilt | nir_op_uge | nir_op_ult => {
1378 let x = *srcs[0].as_ssa().unwrap();
1379 let y = *srcs[1].as_ssa().unwrap();
1380 let (cmp_type, cmp_op) = match alu.op {
1381 nir_op_ige => (IntCmpType::I32, IntCmpOp::Ge),
1382 nir_op_ilt => (IntCmpType::I32, IntCmpOp::Lt),
1383 nir_op_uge => (IntCmpType::U32, IntCmpOp::Ge),
1384 nir_op_ult => (IntCmpType::U32, IntCmpOp::Lt),
1385 _ => panic!("Not an integer comparison"),
1386 };
1387 if alu.get_src(0).bit_size() == 64 {
1388 b.isetp64(cmp_type, cmp_op, x.into(), y.into())
1389 } else {
1390 assert!(alu.get_src(0).bit_size() == 32);
1391 b.isetp(cmp_type, cmp_op, x.into(), y.into())
1392 }
1393 }
1394 nir_op_imad => {
1395 assert!(alu.def.bit_size() == 32);
1396 let dst = b.alloc_ssa(RegFile::GPR, 1);
1397 b.push_op(OpIMad {
1398 dst: dst.into(),
1399 srcs: [srcs[0], srcs[1], srcs[2]],
1400 signed: false,
1401 });
1402 dst
1403 }
1404 nir_op_imax | nir_op_imin | nir_op_umax | nir_op_umin => {
1405 let (tp, min) = match alu.op {
1406 nir_op_imax => (IntCmpType::I32, SrcRef::False),
1407 nir_op_imin => (IntCmpType::I32, SrcRef::True),
1408 nir_op_umax => (IntCmpType::U32, SrcRef::False),
1409 nir_op_umin => (IntCmpType::U32, SrcRef::True),
1410 _ => panic!("Not an integer min/max"),
1411 };
1412 assert!(alu.def.bit_size() == 32);
1413 b.imnmx(tp, srcs[0], srcs[1], min.into())
1414 }
1415 nir_op_imul => {
1416 assert!(alu.def.bit_size() == 32);
1417 b.imul(srcs[0], srcs[1])
1418 }
1419 nir_op_imul_2x32_64 | nir_op_umul_2x32_64 => {
1420 let signed = alu.op == nir_op_imul_2x32_64;
1421 b.imul_2x32_64(srcs[0], srcs[1], signed)
1422 }
1423 nir_op_imul_high | nir_op_umul_high => {
1424 let signed = alu.op == nir_op_imul_high;
1425 let dst64 = b.imul_2x32_64(srcs[0], srcs[1], signed);
1426 dst64[1].into()
1427 }
1428 nir_op_ine => {
1429 if alu.get_src(0).bit_size() == 1 {
1430 b.lop2(LogicOp2::Xor, srcs[0], srcs[1])
1431 } else if alu.get_src(0).bit_size() == 64 {
1432 b.isetp64(IntCmpType::I32, IntCmpOp::Ne, srcs[0], srcs[1])
1433 } else {
1434 assert!(alu.get_src(0).bit_size() == 32);
1435 b.isetp(IntCmpType::I32, IntCmpOp::Ne, srcs[0], srcs[1])
1436 }
1437 }
1438 nir_op_ineg => {
1439 if alu.def.bit_size == 64 {
1440 b.ineg64(srcs[0])
1441 } else {
1442 assert!(alu.def.bit_size() == 32);
1443 b.ineg(srcs[0])
1444 }
1445 }
1446 nir_op_inot => {
1447 if alu.def.bit_size() == 1 {
1448 b.lop2(LogicOp2::PassB, true.into(), srcs[0].bnot())
1449 } else {
1450 assert!(alu.def.bit_size() == 32);
1451 b.lop2(LogicOp2::PassB, 0.into(), srcs[0].bnot())
1452 }
1453 }
1454 nir_op_ior => b.lop2(LogicOp2::Or, srcs[0], srcs[1]),
1455 nir_op_ishl => {
1456 if alu.def.bit_size() == 64 {
1457 let shift = if let Some(s) = nir_srcs[1].comp_as_uint(0) {
1458 (s as u32).into()
1459 } else {
1460 srcs[1]
1461 };
1462 b.shl64(srcs[0], shift)
1463 } else {
1464 assert!(alu.def.bit_size() == 32);
1465 b.shl(srcs[0], srcs[1])
1466 }
1467 }
1468 nir_op_ishr => {
1469 if alu.def.bit_size() == 64 {
1470 let shift = if let Some(s) = nir_srcs[1].comp_as_uint(0) {
1471 (s as u32).into()
1472 } else {
1473 srcs[1]
1474 };
1475 b.shr64(srcs[0], shift, true)
1476 } else {
1477 assert!(alu.def.bit_size() == 32);
1478 b.shr(srcs[0], srcs[1], true)
1479 }
1480 }
1481 nir_op_isub => match alu.def.bit_size {
1482 32 => b.iadd(srcs[0], srcs[1].ineg(), 0.into()),
1483 64 => b.iadd64(srcs[0], srcs[1].ineg(), 0.into()),
1484 x => panic!("unsupported bit size for nir_op_iadd: {x}"),
1485 },
1486 nir_op_ixor => b.lop2(LogicOp2::Xor, srcs[0], srcs[1]),
1487 nir_op_pack_half_2x16_split | nir_op_pack_half_2x16_rtz_split => {
1488 assert!(alu.get_src(0).bit_size() == 32);
1489
1490 let rnd_mode = match alu.op {
1491 nir_op_pack_half_2x16_split => FRndMode::NearestEven,
1492 nir_op_pack_half_2x16_rtz_split => FRndMode::Zero,
1493 _ => panic!("Unhandled fp16 pack op"),
1494 };
1495
1496 if self.sm.sm() >= 86 {
1497 let result: SSARef = b.alloc_ssa(RegFile::GPR, 1);
1498 b.push_op(OpF2FP {
1499 dst: result.into(),
1500 srcs: [srcs[1], srcs[0]],
1501 rnd_mode: rnd_mode,
1502 });
1503
1504 result
1505 } else {
1506 let low = b.alloc_ssa(RegFile::GPR, 1);
1507 let high = b.alloc_ssa(RegFile::GPR, 1);
1508
1509 b.push_op(OpF2F {
1510 dst: low.into(),
1511 src: srcs[0],
1512 src_type: FloatType::F32,
1513 dst_type: FloatType::F16,
1514 rnd_mode: rnd_mode,
1515 ftz: false,
1516 high: false,
1517 integer_rnd: false,
1518 });
1519
1520 let src_bits = usize::from(alu.get_src(1).bit_size());
1521 let src_type = FloatType::from_bits(src_bits);
1522 assert!(matches!(src_type, FloatType::F32));
1523 b.push_op(OpF2F {
1524 dst: high.into(),
1525 src: srcs[1],
1526 src_type: FloatType::F32,
1527 dst_type: FloatType::F16,
1528 rnd_mode: rnd_mode,
1529 ftz: false,
1530 high: false,
1531 integer_rnd: false,
1532 });
1533
1534 b.prmt(low.into(), high.into(), [0, 1, 4, 5])
1535 }
1536 }
1537 nir_op_prmt_nv => {
1538 let dst = b.alloc_ssa(RegFile::GPR, 1);
1539 b.push_op(OpPrmt {
1540 dst: dst.into(),
1541 srcs: [srcs[1], srcs[2]],
1542 sel: srcs[0],
1543 mode: PrmtMode::Index,
1544 });
1545 dst
1546 }
1547 nir_op_sdot_4x8_iadd => {
1548 let dst = b.alloc_ssa(RegFile::GPR, 1);
1549 b.push_op(OpIDp4 {
1550 dst: dst.into(),
1551 src_types: [IntType::I8, IntType::I8],
1552 srcs: [srcs[0], srcs[1], srcs[2]],
1553 });
1554 dst
1555 }
1556 nir_op_sudot_4x8_iadd => {
1557 let dst = b.alloc_ssa(RegFile::GPR, 1);
1558 b.push_op(OpIDp4 {
1559 dst: dst.into(),
1560 src_types: [IntType::I8, IntType::U8],
1561 srcs: [srcs[0], srcs[1], srcs[2]],
1562 });
1563 dst
1564 }
1565 nir_op_udot_4x8_uadd => {
1566 let dst = b.alloc_ssa(RegFile::GPR, 1);
1567 b.push_op(OpIDp4 {
1568 dst: dst.into(),
1569 src_types: [IntType::U8, IntType::U8],
1570 srcs: [srcs[0], srcs[1], srcs[2]],
1571 });
1572 dst
1573 }
1574 nir_op_u2f16 | nir_op_u2f32 | nir_op_u2f64 => {
1575 let src_bits = alu.get_src(0).src.bit_size();
1576 let dst_bits = alu.def.bit_size();
1577 let dst_type = FloatType::from_bits(dst_bits.into());
1578 let dst = b.alloc_ssa(RegFile::GPR, dst_bits.div_ceil(32));
1579 b.push_op(OpI2F {
1580 dst: dst.into(),
1581 src: srcs[0],
1582 dst_type: dst_type,
1583 src_type: IntType::from_bits(src_bits.into(), false),
1584 rnd_mode: self.float_ctl[dst_type].rnd_mode,
1585 });
1586 dst
1587 }
1588 nir_op_uadd_sat => {
1589 let x = srcs[0].as_ssa().unwrap();
1590 let y = srcs[1].as_ssa().unwrap();
1591 let sum_lo = b.alloc_ssa(RegFile::GPR, 1);
1592 let ovf_lo = b.alloc_ssa(RegFile::Pred, 1);
1593 b.push_op(OpIAdd3 {
1594 dst: sum_lo.into(),
1595 overflow: [ovf_lo.into(), Dst::None],
1596 srcs: [0.into(), x[0].into(), y[0].into()],
1597 });
1598 if alu.def.bit_size() == 64 {
1599 let sum_hi = b.alloc_ssa(RegFile::GPR, 1);
1600 let ovf_hi = b.alloc_ssa(RegFile::Pred, 1);
1601 b.push_op(OpIAdd3X {
1602 dst: sum_hi.into(),
1603 overflow: [ovf_hi.into(), Dst::None],
1604 srcs: [0.into(), x[1].into(), y[1].into()],
1605 carry: [ovf_lo.into(), false.into()],
1606 });
1607 let lo =
1608 b.sel(ovf_hi.into(), u32::MAX.into(), sum_lo.into());
1609 let hi =
1610 b.sel(ovf_hi.into(), u32::MAX.into(), sum_hi.into());
1611 [lo[0], hi[0]].into()
1612 } else {
1613 assert!(alu.def.bit_size() == 32);
1614 b.sel(ovf_lo.into(), u32::MAX.into(), sum_lo.into())
1615 }
1616 }
1617 nir_op_usub_sat => {
1618 let x = srcs[0].as_ssa().unwrap();
1619 let y = srcs[1].as_ssa().unwrap();
1620 let sum_lo = b.alloc_ssa(RegFile::GPR, 1);
1621 let ovf_lo = b.alloc_ssa(RegFile::Pred, 1);
1622 // The result of OpIAdd3X is the 33-bit value
1623 //
1624 // s|o = x + !y + 1
1625 //
1626 // The overflow bit of this result is true if and only if the
1627 // subtract did NOT overflow.
1628 b.push_op(OpIAdd3 {
1629 dst: sum_lo.into(),
1630 overflow: [ovf_lo.into(), Dst::None],
1631 srcs: [0.into(), x[0].into(), Src::from(y[0]).ineg()],
1632 });
1633 if alu.def.bit_size() == 64 {
1634 let sum_hi = b.alloc_ssa(RegFile::GPR, 1);
1635 let ovf_hi = b.alloc_ssa(RegFile::Pred, 1);
1636 b.push_op(OpIAdd3X {
1637 dst: sum_hi.into(),
1638 overflow: [ovf_hi.into(), Dst::None],
1639 srcs: [0.into(), x[1].into(), Src::from(y[1]).bnot()],
1640 carry: [ovf_lo.into(), false.into()],
1641 });
1642 let lo = b.sel(ovf_hi.into(), sum_lo.into(), 0.into());
1643 let hi = b.sel(ovf_hi.into(), sum_hi.into(), 0.into());
1644 [lo[0], hi[0]].into()
1645 } else {
1646 assert!(alu.def.bit_size() == 32);
1647 b.sel(ovf_lo.into(), sum_lo.into(), 0.into())
1648 }
1649 }
1650 nir_op_unpack_32_2x16_split_x => {
1651 b.prmt(srcs[0], 0.into(), [0, 1, 4, 4])
1652 }
1653 nir_op_unpack_32_2x16_split_y => {
1654 b.prmt(srcs[0], 0.into(), [2, 3, 4, 4])
1655 }
1656 nir_op_unpack_64_2x32_split_x => {
1657 let src0_x = srcs[0].as_ssa().unwrap()[0];
1658 b.copy(src0_x.into())
1659 }
1660 nir_op_unpack_64_2x32_split_y => {
1661 let src0_y = srcs[0].as_ssa().unwrap()[1];
1662 b.copy(src0_y.into())
1663 }
1664 nir_op_unpack_half_2x16_split_x
1665 | nir_op_unpack_half_2x16_split_y => {
1666 assert!(alu.def.bit_size() == 32);
1667 let dst = b.alloc_ssa(RegFile::GPR, 1);
1668
1669 b.push_op(OpF2F {
1670 dst: dst[0].into(),
1671 src: srcs[0],
1672 src_type: FloatType::F16,
1673 dst_type: FloatType::F32,
1674 rnd_mode: FRndMode::NearestEven,
1675 ftz: false,
1676 high: alu.op == nir_op_unpack_half_2x16_split_y,
1677 integer_rnd: false,
1678 });
1679
1680 dst
1681 }
1682 nir_op_ushr => {
1683 if alu.def.bit_size() == 64 {
1684 let shift = if let Some(s) = nir_srcs[1].comp_as_uint(0) {
1685 (s as u32).into()
1686 } else {
1687 srcs[1]
1688 };
1689 b.shr64(srcs[0], shift, false)
1690 } else {
1691 assert!(alu.def.bit_size() == 32);
1692 b.shr(srcs[0], srcs[1], false)
1693 }
1694 }
1695 _ => panic!("Unsupported ALU instruction: {}", alu.info().name()),
1696 };
1697 self.set_dst(&alu.def, dst);
1698 }
1699
parse_tex(&mut self, b: &mut impl SSABuilder, tex: &nir_tex_instr)1700 fn parse_tex(&mut self, b: &mut impl SSABuilder, tex: &nir_tex_instr) {
1701 let dim = match tex.sampler_dim {
1702 GLSL_SAMPLER_DIM_1D => {
1703 if tex.is_array {
1704 TexDim::Array1D
1705 } else {
1706 TexDim::_1D
1707 }
1708 }
1709 GLSL_SAMPLER_DIM_2D => {
1710 if tex.is_array {
1711 TexDim::Array2D
1712 } else {
1713 TexDim::_2D
1714 }
1715 }
1716 GLSL_SAMPLER_DIM_3D => {
1717 assert!(!tex.is_array);
1718 TexDim::_3D
1719 }
1720 GLSL_SAMPLER_DIM_CUBE => {
1721 if tex.is_array {
1722 TexDim::ArrayCube
1723 } else {
1724 TexDim::Cube
1725 }
1726 }
1727 GLSL_SAMPLER_DIM_BUF => TexDim::_1D,
1728 GLSL_SAMPLER_DIM_MS => {
1729 if tex.is_array {
1730 TexDim::Array2D
1731 } else {
1732 TexDim::_2D
1733 }
1734 }
1735 _ => panic!("Unsupported texture dimension: {}", tex.sampler_dim),
1736 };
1737
1738 let srcs = tex.srcs_as_slice();
1739 assert!(srcs[0].src_type == nir_tex_src_backend1);
1740 if srcs.len() > 1 {
1741 assert!(srcs.len() == 2);
1742 assert!(srcs[1].src_type == nir_tex_src_backend2);
1743 }
1744
1745 let flags: nak_nir_tex_flags =
1746 unsafe { std::mem::transmute_copy(&tex.backend_flags) };
1747
1748 let mask = tex.def.components_read();
1749 let mut mask = u8::try_from(mask).unwrap();
1750 if flags.is_sparse() {
1751 mask &= !(1 << (tex.def.num_components - 1));
1752 }
1753
1754 let dst_comps = u8::try_from(mask.count_ones()).unwrap();
1755 let dst = b.alloc_ssa(RegFile::GPR, dst_comps);
1756
1757 // On Volta and later, the destination is split in two
1758 let mut dsts = [Dst::None; 2];
1759 if dst_comps > 2 && b.sm() >= 70 {
1760 dsts[0] = SSARef::try_from(&dst[0..2]).unwrap().into();
1761 dsts[1] = SSARef::try_from(&dst[2..]).unwrap().into();
1762 } else {
1763 dsts[0] = dst.into();
1764 }
1765
1766 let fault = if flags.is_sparse() {
1767 b.alloc_ssa(RegFile::Pred, 1).into()
1768 } else {
1769 Dst::None
1770 };
1771
1772 if tex.op == nir_texop_hdr_dim_nv {
1773 let src = self.get_src(&srcs[0].src);
1774 assert!(fault.is_none());
1775 b.push_op(OpTxq {
1776 dsts: dsts,
1777 src: src,
1778 query: TexQuery::Dimension,
1779 mask: mask,
1780 });
1781 } else if tex.op == nir_texop_tex_type_nv {
1782 let src = self.get_src(&srcs[0].src);
1783 assert!(fault.is_none());
1784 b.push_op(OpTxq {
1785 dsts: dsts,
1786 src: src,
1787 query: TexQuery::TextureType,
1788 mask: mask,
1789 });
1790 } else {
1791 let lod_mode = match flags.lod_mode() {
1792 NAK_NIR_LOD_MODE_AUTO => TexLodMode::Auto,
1793 NAK_NIR_LOD_MODE_ZERO => TexLodMode::Zero,
1794 NAK_NIR_LOD_MODE_BIAS => TexLodMode::Bias,
1795 NAK_NIR_LOD_MODE_LOD => TexLodMode::Lod,
1796 NAK_NIR_LOD_MODE_CLAMP => TexLodMode::Clamp,
1797 NAK_NIR_LOD_MODE_BIAS_CLAMP => TexLodMode::BiasClamp,
1798 _ => panic!("Invalid LOD mode"),
1799 };
1800
1801 let offset_mode = match flags.offset_mode() {
1802 NAK_NIR_OFFSET_MODE_NONE => Tld4OffsetMode::None,
1803 NAK_NIR_OFFSET_MODE_AOFFI => Tld4OffsetMode::AddOffI,
1804 NAK_NIR_OFFSET_MODE_PER_PX => Tld4OffsetMode::PerPx,
1805 _ => panic!("Invalid offset mode"),
1806 };
1807
1808 let srcs = [self.get_src(&srcs[0].src), self.get_src(&srcs[1].src)];
1809
1810 if tex.op == nir_texop_txd {
1811 assert!(lod_mode == TexLodMode::Auto);
1812 assert!(offset_mode != Tld4OffsetMode::PerPx);
1813 assert!(!flags.has_z_cmpr());
1814 b.push_op(OpTxd {
1815 dsts: dsts,
1816 fault,
1817 srcs: srcs,
1818 dim: dim,
1819 offset: offset_mode == Tld4OffsetMode::AddOffI,
1820 mask: mask,
1821 });
1822 } else if tex.op == nir_texop_lod {
1823 assert!(offset_mode == Tld4OffsetMode::None);
1824 b.push_op(OpTmml {
1825 dsts: dsts,
1826 srcs: srcs,
1827 dim: dim,
1828 mask: mask,
1829 });
1830 } else if tex.op == nir_texop_txf || tex.op == nir_texop_txf_ms {
1831 assert!(offset_mode != Tld4OffsetMode::PerPx);
1832 b.push_op(OpTld {
1833 dsts: dsts,
1834 fault,
1835 srcs: srcs,
1836 dim: dim,
1837 lod_mode: lod_mode,
1838 is_ms: tex.op == nir_texop_txf_ms,
1839 offset: offset_mode == Tld4OffsetMode::AddOffI,
1840 mask: mask,
1841 });
1842 } else if tex.op == nir_texop_tg4 {
1843 b.push_op(OpTld4 {
1844 dsts: dsts,
1845 fault,
1846 srcs: srcs,
1847 dim: dim,
1848 comp: tex.component().try_into().unwrap(),
1849 offset_mode: offset_mode,
1850 z_cmpr: flags.has_z_cmpr(),
1851 mask: mask,
1852 });
1853 } else {
1854 assert!(offset_mode != Tld4OffsetMode::PerPx);
1855 b.push_op(OpTex {
1856 dsts: dsts,
1857 fault,
1858 srcs: srcs,
1859 dim: dim,
1860 lod_mode: lod_mode,
1861 z_cmpr: flags.has_z_cmpr(),
1862 offset: offset_mode == Tld4OffsetMode::AddOffI,
1863 mask: mask,
1864 });
1865 }
1866 }
1867
1868 let mut di = 0_usize;
1869 let mut nir_dst = Vec::new();
1870 for i in 0..tex.def.num_components() {
1871 if flags.is_sparse() && i == tex.def.num_components - 1 {
1872 let Dst::SSA(fault) = fault else {
1873 panic!("No fault value for sparse op");
1874 };
1875 nir_dst.push(b.sel(fault.into(), 0.into(), 1.into())[0]);
1876 } else if mask & (1 << i) == 0 {
1877 nir_dst.push(b.copy(0.into())[0]);
1878 } else {
1879 nir_dst.push(dst[di]);
1880 di += 1;
1881 }
1882 }
1883 self.set_ssa(tex.def.as_def(), nir_dst);
1884 }
1885
get_atomic_type(&self, intrin: &nir_intrinsic_instr) -> AtomType1886 fn get_atomic_type(&self, intrin: &nir_intrinsic_instr) -> AtomType {
1887 let bit_size = intrin.def.bit_size();
1888 match intrin.atomic_op() {
1889 nir_atomic_op_iadd => AtomType::U(bit_size),
1890 nir_atomic_op_imin => AtomType::I(bit_size),
1891 nir_atomic_op_umin => AtomType::U(bit_size),
1892 nir_atomic_op_imax => AtomType::I(bit_size),
1893 nir_atomic_op_umax => AtomType::U(bit_size),
1894 nir_atomic_op_iand => AtomType::U(bit_size),
1895 nir_atomic_op_ior => AtomType::U(bit_size),
1896 nir_atomic_op_ixor => AtomType::U(bit_size),
1897 nir_atomic_op_xchg => AtomType::U(bit_size),
1898 nir_atomic_op_fadd => AtomType::F(bit_size),
1899 nir_atomic_op_fmin => AtomType::F(bit_size),
1900 nir_atomic_op_fmax => AtomType::F(bit_size),
1901 nir_atomic_op_cmpxchg => AtomType::U(bit_size),
1902 _ => panic!("Unsupported NIR atomic op"),
1903 }
1904 }
1905
get_atomic_op( &self, intrin: &nir_intrinsic_instr, cmp_src: AtomCmpSrc, ) -> AtomOp1906 fn get_atomic_op(
1907 &self,
1908 intrin: &nir_intrinsic_instr,
1909 cmp_src: AtomCmpSrc,
1910 ) -> AtomOp {
1911 match intrin.atomic_op() {
1912 nir_atomic_op_iadd => AtomOp::Add,
1913 nir_atomic_op_imin => AtomOp::Min,
1914 nir_atomic_op_umin => AtomOp::Min,
1915 nir_atomic_op_imax => AtomOp::Max,
1916 nir_atomic_op_umax => AtomOp::Max,
1917 nir_atomic_op_iand => AtomOp::And,
1918 nir_atomic_op_ior => AtomOp::Or,
1919 nir_atomic_op_ixor => AtomOp::Xor,
1920 nir_atomic_op_xchg => AtomOp::Exch,
1921 nir_atomic_op_fadd => AtomOp::Add,
1922 nir_atomic_op_fmin => AtomOp::Min,
1923 nir_atomic_op_fmax => AtomOp::Max,
1924 nir_atomic_op_cmpxchg => AtomOp::CmpExch(cmp_src),
1925 _ => panic!("Unsupported NIR atomic op"),
1926 }
1927 }
1928
get_eviction_priority( &mut self, access: gl_access_qualifier, ) -> MemEvictionPriority1929 fn get_eviction_priority(
1930 &mut self,
1931 access: gl_access_qualifier,
1932 ) -> MemEvictionPriority {
1933 if self.sm.sm() >= 70 && access & ACCESS_NON_TEMPORAL != 0 {
1934 MemEvictionPriority::First
1935 } else {
1936 MemEvictionPriority::Normal
1937 }
1938 }
1939
get_image_dim(&mut self, intrin: &nir_intrinsic_instr) -> ImageDim1940 fn get_image_dim(&mut self, intrin: &nir_intrinsic_instr) -> ImageDim {
1941 let is_array = intrin.image_array();
1942 let image_dim = intrin.image_dim();
1943 match intrin.image_dim() {
1944 GLSL_SAMPLER_DIM_1D => {
1945 if is_array {
1946 ImageDim::_1DArray
1947 } else {
1948 ImageDim::_1D
1949 }
1950 }
1951 GLSL_SAMPLER_DIM_2D => {
1952 if is_array {
1953 ImageDim::_2DArray
1954 } else {
1955 ImageDim::_2D
1956 }
1957 }
1958 GLSL_SAMPLER_DIM_3D => {
1959 assert!(!is_array);
1960 ImageDim::_3D
1961 }
1962 GLSL_SAMPLER_DIM_CUBE => ImageDim::_2DArray,
1963 GLSL_SAMPLER_DIM_BUF => {
1964 assert!(!is_array);
1965 ImageDim::_1DBuffer
1966 }
1967 _ => panic!("Unsupported image dimension: {}", image_dim),
1968 }
1969 }
1970
get_image_coord( &mut self, intrin: &nir_intrinsic_instr, dim: ImageDim, ) -> Src1971 fn get_image_coord(
1972 &mut self,
1973 intrin: &nir_intrinsic_instr,
1974 dim: ImageDim,
1975 ) -> Src {
1976 let vec = self.get_ssa(intrin.get_src(1).as_def());
1977 // let sample = self.get_src(&srcs[2]);
1978 let comps = usize::from(dim.coord_comps());
1979 SSARef::try_from(&vec[0..comps]).unwrap().into()
1980 }
1981
parse_intrinsic( &mut self, b: &mut impl SSABuilder, intrin: &nir_intrinsic_instr, )1982 fn parse_intrinsic(
1983 &mut self,
1984 b: &mut impl SSABuilder,
1985 intrin: &nir_intrinsic_instr,
1986 ) {
1987 let srcs = intrin.srcs_as_slice();
1988 match intrin.intrinsic {
1989 nir_intrinsic_al2p_nv => {
1990 let offset = self.get_src(&srcs[0]);
1991 let addr = u16::try_from(intrin.base()).unwrap();
1992
1993 let flags = intrin.flags();
1994 let flags: nak_nir_attr_io_flags =
1995 unsafe { std::mem::transmute_copy(&flags) };
1996
1997 let access = AttrAccess {
1998 addr: addr,
1999 comps: 1,
2000 patch: flags.patch(),
2001 output: flags.output(),
2002 phys: false,
2003 };
2004
2005 let dst = b.alloc_ssa(RegFile::GPR, 1);
2006 b.push_op(OpAL2P {
2007 dst: dst.into(),
2008 offset: offset,
2009 access: access,
2010 });
2011 self.set_dst(&intrin.def, dst);
2012 }
2013 nir_intrinsic_ald_nv | nir_intrinsic_ast_nv => {
2014 let addr = u16::try_from(intrin.base()).unwrap();
2015 let base = u16::try_from(intrin.range_base()).unwrap();
2016 let range = u16::try_from(intrin.range()).unwrap();
2017 let range = base..(base + range);
2018
2019 let flags = intrin.flags();
2020 let flags: nak_nir_attr_io_flags =
2021 unsafe { std::mem::transmute_copy(&flags) };
2022 assert!(!flags.patch() || !flags.phys());
2023
2024 if let ShaderIoInfo::Vtg(io) = &mut self.info.io {
2025 if flags.patch() {
2026 match &mut self.info.stage {
2027 ShaderStageInfo::TessellationInit(stage) => {
2028 assert!(flags.output());
2029 stage.per_patch_attribute_count = max(
2030 stage.per_patch_attribute_count,
2031 (range.end / 4).try_into().unwrap(),
2032 );
2033 }
2034 ShaderStageInfo::Tessellation(_) => (),
2035 _ => panic!("Patch I/O not supported"),
2036 }
2037 } else {
2038 if flags.output() {
2039 if intrin.intrinsic == nir_intrinsic_ast_nv {
2040 io.mark_store_req(range.clone());
2041 }
2042 io.mark_attrs_written(range);
2043 } else {
2044 io.mark_attrs_read(range);
2045 }
2046 }
2047 } else {
2048 panic!("Must be a VTG stage");
2049 }
2050
2051 let access = AttrAccess {
2052 addr: addr,
2053 comps: intrin.num_components,
2054 patch: flags.patch(),
2055 output: flags.output(),
2056 phys: flags.phys(),
2057 };
2058
2059 if intrin.intrinsic == nir_intrinsic_ald_nv {
2060 let vtx = self.get_src(&srcs[0]);
2061 let offset = self.get_src(&srcs[1]);
2062
2063 assert!(intrin.def.bit_size() == 32);
2064 let dst = b.alloc_ssa(RegFile::GPR, access.comps);
2065 b.push_op(OpALd {
2066 dst: dst.into(),
2067 vtx: vtx,
2068 offset: offset,
2069 access: access,
2070 });
2071 self.set_dst(&intrin.def, dst);
2072 } else if intrin.intrinsic == nir_intrinsic_ast_nv {
2073 assert!(srcs[0].bit_size() == 32);
2074 let data = self.get_src(&srcs[0]);
2075 let vtx = self.get_src(&srcs[1]);
2076 let offset = self.get_src(&srcs[2]);
2077
2078 b.push_op(OpASt {
2079 data: data,
2080 vtx: vtx,
2081 offset: offset,
2082 access: access,
2083 });
2084 } else {
2085 panic!("Invalid VTG I/O intrinsic");
2086 }
2087 }
2088 nir_intrinsic_as_uniform => {
2089 let src = self.get_ssa(srcs[0].as_def());
2090 let mut dst = Vec::new();
2091 for comp in src {
2092 let u = b.alloc_ssa(RegFile::UGPR, 1);
2093 b.push_op(OpR2UR {
2094 src: [*comp].into(),
2095 dst: u.into(),
2096 });
2097 dst.push(u[0]);
2098 }
2099 self.set_ssa(&intrin.def, dst);
2100 }
2101 nir_intrinsic_ddx
2102 | nir_intrinsic_ddx_coarse
2103 | nir_intrinsic_ddx_fine => {
2104 // TODO: Real coarse derivatives
2105
2106 assert!(intrin.def.bit_size() == 32);
2107 let ftype = FloatType::F32;
2108 let scratch = b.alloc_ssa(RegFile::GPR, 1);
2109
2110 b.push_op(OpShfl {
2111 dst: scratch[0].into(),
2112 in_bounds: Dst::None,
2113 src: self.get_src(&srcs[0]),
2114 lane: 1_u32.into(),
2115 c: (0x3_u32 | 0x1c_u32 << 8).into(),
2116 op: ShflOp::Bfly,
2117 });
2118
2119 let dst = b.alloc_ssa(RegFile::GPR, 1);
2120
2121 b.push_op(OpFSwzAdd {
2122 dst: dst[0].into(),
2123 srcs: [scratch[0].into(), self.get_src(&srcs[0])],
2124 ops: [
2125 FSwzAddOp::SubLeft,
2126 FSwzAddOp::SubRight,
2127 FSwzAddOp::SubLeft,
2128 FSwzAddOp::SubRight,
2129 ],
2130 rnd_mode: self.float_ctl[ftype].rnd_mode,
2131 ftz: self.float_ctl[ftype].ftz,
2132 });
2133
2134 self.set_dst(&intrin.def, dst);
2135 }
2136 nir_intrinsic_ddy
2137 | nir_intrinsic_ddy_coarse
2138 | nir_intrinsic_ddy_fine => {
2139 // TODO: Real coarse derivatives
2140
2141 assert!(intrin.def.bit_size() == 32);
2142 let ftype = FloatType::F32;
2143 let scratch = b.alloc_ssa(RegFile::GPR, 1);
2144
2145 b.push_op(OpShfl {
2146 dst: scratch[0].into(),
2147 in_bounds: Dst::None,
2148 src: self.get_src(&srcs[0]),
2149 lane: 2_u32.into(),
2150 c: (0x3_u32 | 0x1c_u32 << 8).into(),
2151 op: ShflOp::Bfly,
2152 });
2153
2154 let dst = b.alloc_ssa(RegFile::GPR, 1);
2155
2156 b.push_op(OpFSwzAdd {
2157 dst: dst[0].into(),
2158 srcs: [scratch[0].into(), self.get_src(&srcs[0])],
2159 ops: [
2160 FSwzAddOp::SubLeft,
2161 FSwzAddOp::SubLeft,
2162 FSwzAddOp::SubRight,
2163 FSwzAddOp::SubRight,
2164 ],
2165 rnd_mode: self.float_ctl[ftype].rnd_mode,
2166 ftz: self.float_ctl[ftype].ftz,
2167 });
2168
2169 self.set_dst(&intrin.def, dst);
2170 }
2171 nir_intrinsic_ballot => {
2172 assert!(srcs[0].bit_size() == 1);
2173 let src = self.get_src(&srcs[0]);
2174
2175 assert!(intrin.def.bit_size() == 32);
2176 let dst = b.alloc_ssa(RegFile::GPR, 1);
2177
2178 b.push_op(OpVote {
2179 op: VoteOp::Any,
2180 ballot: dst.into(),
2181 vote: Dst::None,
2182 pred: src,
2183 });
2184 self.set_dst(&intrin.def, dst);
2185 }
2186 nir_intrinsic_bar_break_nv => {
2187 let src = self.get_src(&srcs[0]);
2188 let bar_in = b.bmov_to_bar(src);
2189 let cond = self.get_src(&srcs[1]);
2190
2191 let bar_out = b.alloc_ssa(RegFile::Bar, 1);
2192 b.push_op(OpBreak {
2193 bar_out: bar_out.into(),
2194 bar_in: bar_in.into(),
2195 cond: cond.into(),
2196 });
2197
2198 self.set_dst(&intrin.def, b.bmov_to_gpr(bar_out.into()));
2199 }
2200 nir_intrinsic_bar_set_nv => {
2201 let label = self.label_alloc.alloc();
2202 let old = self.bar_label.insert(intrin.def.index, label);
2203 assert!(old.is_none());
2204
2205 let bar_clear = b.alloc_ssa(RegFile::Bar, 1);
2206 b.push_op(OpBClear {
2207 dst: bar_clear.into(),
2208 });
2209
2210 let bar_out = b.alloc_ssa(RegFile::Bar, 1);
2211 b.push_op(OpBSSy {
2212 bar_out: bar_out.into(),
2213 bar_in: bar_clear.into(),
2214 cond: SrcRef::True.into(),
2215 target: label,
2216 });
2217
2218 self.set_dst(&intrin.def, b.bmov_to_gpr(bar_out.into()));
2219 }
2220 nir_intrinsic_bar_sync_nv => {
2221 let src = self.get_src(&srcs[0]);
2222
2223 let bar = b.bmov_to_bar(src);
2224 b.push_op(OpBSync {
2225 bar: bar.into(),
2226 cond: SrcRef::True.into(),
2227 });
2228
2229 let bar_set_idx = &srcs[1].as_def().index;
2230 if let Some(label) = self.bar_label.get(bar_set_idx) {
2231 b.push_op(OpNop {
2232 label: Some(*label),
2233 });
2234 }
2235 }
2236 nir_intrinsic_bindless_image_atomic
2237 | nir_intrinsic_bindless_image_atomic_swap => {
2238 let handle = self.get_src(&srcs[0]);
2239 let dim = self.get_image_dim(intrin);
2240 let coord = self.get_image_coord(intrin, dim);
2241 // let sample = self.get_src(&srcs[2]);
2242 let atom_type = self.get_atomic_type(intrin);
2243 let atom_op = self.get_atomic_op(intrin, AtomCmpSrc::Packed);
2244
2245 assert!(
2246 intrin.def.bit_size() == 32 || intrin.def.bit_size() == 64
2247 );
2248 assert!(intrin.def.num_components() == 1);
2249 let dst = b.alloc_ssa(RegFile::GPR, intrin.def.bit_size() / 32);
2250
2251 let data = if intrin.intrinsic
2252 == nir_intrinsic_bindless_image_atomic_swap
2253 {
2254 if intrin.def.bit_size() == 64 {
2255 SSARef::from([
2256 self.get_ssa(srcs[3].as_def())[0],
2257 self.get_ssa(srcs[3].as_def())[1],
2258 self.get_ssa(srcs[4].as_def())[0],
2259 self.get_ssa(srcs[4].as_def())[1],
2260 ])
2261 .into()
2262 } else {
2263 SSARef::from([
2264 self.get_ssa(srcs[3].as_def())[0],
2265 self.get_ssa(srcs[4].as_def())[0],
2266 ])
2267 .into()
2268 }
2269 } else {
2270 self.get_src(&srcs[3])
2271 };
2272
2273 let is_reduction =
2274 atom_op.is_reduction() && intrin.def.components_read() == 0;
2275
2276 b.push_op(OpSuAtom {
2277 dst: if self.sm.sm() >= 70 && is_reduction {
2278 Dst::None
2279 } else {
2280 dst.into()
2281 },
2282 fault: Dst::None,
2283 handle: handle,
2284 coord: coord,
2285 data: data,
2286 atom_op: atom_op,
2287 atom_type: atom_type,
2288 image_dim: dim,
2289 mem_order: MemOrder::Strong(MemScope::System),
2290 mem_eviction_priority: self
2291 .get_eviction_priority(intrin.access()),
2292 });
2293 self.set_dst(&intrin.def, dst);
2294 }
2295 nir_intrinsic_bindless_image_load => {
2296 let handle = self.get_src(&srcs[0]);
2297 let dim = self.get_image_dim(intrin);
2298 let coord = self.get_image_coord(intrin, dim);
2299 // let sample = self.get_src(&srcs[2]);
2300
2301 let comps = intrin.num_components;
2302 assert!(intrin.def.bit_size() == 32);
2303 assert!(comps == 1 || comps == 2 || comps == 4);
2304
2305 let dst = b.alloc_ssa(RegFile::GPR, comps);
2306
2307 b.push_op(OpSuLd {
2308 dst: dst.into(),
2309 fault: Dst::None,
2310 image_dim: dim,
2311 mem_order: MemOrder::Strong(MemScope::System),
2312 mem_eviction_priority: self
2313 .get_eviction_priority(intrin.access()),
2314 mask: (1 << comps) - 1,
2315 handle: handle,
2316 coord: coord,
2317 });
2318 self.set_dst(&intrin.def, dst);
2319 }
2320 nir_intrinsic_bindless_image_sparse_load => {
2321 let handle = self.get_src(&srcs[0]);
2322 let dim = self.get_image_dim(intrin);
2323 let coord = self.get_image_coord(intrin, dim);
2324 // let sample = self.get_src(&srcs[2]);
2325
2326 let comps = intrin.num_components;
2327 assert!(intrin.def.bit_size() == 32);
2328 assert!(comps == 5);
2329
2330 let dst = b.alloc_ssa(RegFile::GPR, comps - 1);
2331 let fault = b.alloc_ssa(RegFile::Pred, 1);
2332
2333 b.push_op(OpSuLd {
2334 dst: dst.into(),
2335 fault: fault.into(),
2336 image_dim: dim,
2337 mem_order: MemOrder::Strong(MemScope::System),
2338 mem_eviction_priority: self
2339 .get_eviction_priority(intrin.access()),
2340 mask: (1 << (comps - 1)) - 1,
2341 handle: handle,
2342 coord: coord,
2343 });
2344
2345 let mut final_dst = Vec::new();
2346 for i in 0..usize::from(comps) - 1 {
2347 final_dst.push(dst[i]);
2348 }
2349 final_dst.push(b.sel(fault.into(), 0.into(), 1.into())[0]);
2350
2351 self.set_ssa(&intrin.def, final_dst);
2352 }
2353 nir_intrinsic_bindless_image_store => {
2354 let handle = self.get_src(&srcs[0]);
2355 let dim = self.get_image_dim(intrin);
2356 let coord = self.get_image_coord(intrin, dim);
2357 // let sample = self.get_src(&srcs[2]);
2358 let data = self.get_src(&srcs[3]);
2359
2360 let comps = intrin.num_components;
2361 assert!(srcs[3].bit_size() == 32);
2362 assert!(comps == 1 || comps == 2 || comps == 4);
2363
2364 b.push_op(OpSuSt {
2365 image_dim: dim,
2366 mem_order: MemOrder::Strong(MemScope::System),
2367 mem_eviction_priority: self
2368 .get_eviction_priority(intrin.access()),
2369 mask: (1 << comps) - 1,
2370 handle: handle,
2371 coord: coord,
2372 data: data,
2373 });
2374 }
2375 nir_intrinsic_copy_fs_outputs_nv => {
2376 let ShaderIoInfo::Fragment(info) = &mut self.info.io else {
2377 panic!(
2378 "copy_fs_outputs_nv is only allowed in fragment shaders"
2379 );
2380 };
2381
2382 for i in 0..32 {
2383 if !self.fs_out_regs[i].is_none() {
2384 info.writes_color |= 1 << i;
2385 }
2386 }
2387 let mask_idx = (NAK_FS_OUT_SAMPLE_MASK / 4) as usize;
2388 info.writes_sample_mask = !self.fs_out_regs[mask_idx].is_none();
2389 let depth_idx = (NAK_FS_OUT_DEPTH / 4) as usize;
2390 info.writes_depth = !self.fs_out_regs[depth_idx].is_none();
2391
2392 let mut srcs = Vec::new();
2393 for i in 0..8 {
2394 // Even though the mask is per-component, the actual output
2395 // space is per-output vec4s.
2396 if info.writes_color & (0xf << (i * 4)) != 0 {
2397 for c in 0..4 {
2398 let reg = self.fs_out_regs[i * 4 + c];
2399 if reg.is_none() {
2400 srcs.push(b.undef().into());
2401 } else {
2402 srcs.push(reg.into());
2403 }
2404 }
2405 }
2406 }
2407
2408 // These always come together for some reason
2409 if info.writes_sample_mask || info.writes_depth {
2410 if info.writes_sample_mask {
2411 srcs.push(self.fs_out_regs[mask_idx].into());
2412 } else {
2413 srcs.push(b.undef().into());
2414 }
2415 if info.writes_depth {
2416 srcs.push(self.fs_out_regs[depth_idx].into());
2417 }
2418 }
2419
2420 b.push_op(OpRegOut { srcs: srcs });
2421 }
2422 nir_intrinsic_demote => {
2423 if let ShaderStageInfo::Fragment(info) = &mut self.info.stage {
2424 info.uses_kill = true;
2425 } else {
2426 panic!("OpKill is only available in fragment shaders");
2427 }
2428 b.push_op(OpKill {});
2429 }
2430 nir_intrinsic_demote_if => {
2431 if let ShaderStageInfo::Fragment(info) = &mut self.info.stage {
2432 info.uses_kill = true;
2433 } else {
2434 panic!("OpKill is only available in fragment shaders");
2435 }
2436 let cond = self.get_ssa(srcs[0].as_def())[0];
2437 b.predicate(cond.into()).push_op(OpKill {});
2438 }
2439 nir_intrinsic_global_atomic => {
2440 let bit_size = intrin.def.bit_size();
2441 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2442 let data = self.get_src(&srcs[1]);
2443 let atom_type = self.get_atomic_type(intrin);
2444 let atom_op = self.get_atomic_op(intrin, AtomCmpSrc::Separate);
2445
2446 assert!(intrin.def.num_components() == 1);
2447 let dst = b.alloc_ssa(RegFile::GPR, bit_size.div_ceil(32));
2448
2449 let is_reduction =
2450 atom_op.is_reduction() && intrin.def.components_read() == 0;
2451
2452 b.push_op(OpAtom {
2453 dst: if is_reduction { Dst::None } else { dst.into() },
2454 addr: addr,
2455 cmpr: 0.into(),
2456 data: data,
2457 atom_op: atom_op,
2458 atom_type: atom_type,
2459 addr_offset: offset,
2460 mem_space: MemSpace::Global(MemAddrType::A64),
2461 mem_order: MemOrder::Strong(MemScope::System),
2462 mem_eviction_priority: MemEvictionPriority::Normal, // Note: no intrinic access
2463 });
2464 self.set_dst(&intrin.def, dst);
2465 }
2466 nir_intrinsic_global_atomic_swap => {
2467 assert!(intrin.atomic_op() == nir_atomic_op_cmpxchg);
2468 let bit_size = intrin.def.bit_size();
2469 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2470 let cmpr = self.get_src(&srcs[1]);
2471 let data = self.get_src(&srcs[2]);
2472 let atom_type = AtomType::U(bit_size);
2473
2474 assert!(intrin.def.num_components() == 1);
2475 let dst = b.alloc_ssa(RegFile::GPR, bit_size.div_ceil(32));
2476
2477 b.push_op(OpAtom {
2478 dst: dst.into(),
2479 addr: addr,
2480 cmpr: cmpr,
2481 data: data,
2482 atom_op: AtomOp::CmpExch(AtomCmpSrc::Separate),
2483 atom_type: atom_type,
2484 addr_offset: offset,
2485 mem_space: MemSpace::Global(MemAddrType::A64),
2486 mem_order: MemOrder::Strong(MemScope::System),
2487 mem_eviction_priority: MemEvictionPriority::Normal, // Note: no intrinic access
2488 });
2489 self.set_dst(&intrin.def, dst);
2490 }
2491 nir_intrinsic_ipa_nv => {
2492 let addr = u16::try_from(intrin.base()).unwrap();
2493
2494 let flags = intrin.flags();
2495 let flags: nak_nir_ipa_flags =
2496 unsafe { std::mem::transmute_copy(&flags) };
2497
2498 let mode = match flags.interp_mode() {
2499 NAK_INTERP_MODE_PERSPECTIVE => PixelImap::Perspective,
2500 NAK_INTERP_MODE_SCREEN_LINEAR => PixelImap::ScreenLinear,
2501 NAK_INTERP_MODE_CONSTANT => PixelImap::Constant,
2502 _ => panic!("Unsupported interp mode"),
2503 };
2504
2505 let freq = match flags.interp_freq() {
2506 NAK_INTERP_FREQ_PASS => InterpFreq::Pass,
2507 NAK_INTERP_FREQ_PASS_MUL_W => InterpFreq::PassMulW,
2508 NAK_INTERP_FREQ_CONSTANT => InterpFreq::Constant,
2509 NAK_INTERP_FREQ_STATE => InterpFreq::State,
2510 _ => panic!("Invalid interp freq"),
2511 };
2512
2513 let loc = match flags.interp_loc() {
2514 NAK_INTERP_LOC_DEFAULT => InterpLoc::Default,
2515 NAK_INTERP_LOC_CENTROID => InterpLoc::Centroid,
2516 NAK_INTERP_LOC_OFFSET => InterpLoc::Offset,
2517 _ => panic!("Invalid interp loc"),
2518 };
2519
2520 let inv_w = if freq == InterpFreq::PassMulW {
2521 self.get_src(&srcs[0])
2522 } else {
2523 0.into()
2524 };
2525
2526 let offset = if loc == InterpLoc::Offset {
2527 self.get_src(&srcs[1])
2528 } else {
2529 0.into()
2530 };
2531
2532 let ShaderIoInfo::Fragment(io) = &mut self.info.io else {
2533 panic!("OpIpa is only used for fragment shaders");
2534 };
2535
2536 io.mark_attr_read(addr, mode);
2537
2538 let dst = b.alloc_ssa(RegFile::GPR, 1);
2539 b.push_op(OpIpa {
2540 dst: dst.into(),
2541 addr: addr,
2542 freq: freq,
2543 loc: loc,
2544 inv_w: inv_w,
2545 offset: offset,
2546 });
2547 self.set_dst(&intrin.def, dst);
2548 }
2549 nir_intrinsic_isberd_nv => {
2550 let dst = b.alloc_ssa(RegFile::GPR, 1);
2551 b.push_op(OpIsberd {
2552 dst: dst.into(),
2553 idx: self.get_src(&srcs[0]),
2554 });
2555 self.set_dst(&intrin.def, dst);
2556 }
2557 nir_intrinsic_load_barycentric_at_offset_nv => (),
2558 nir_intrinsic_load_barycentric_centroid => (),
2559 nir_intrinsic_load_barycentric_pixel => (),
2560 nir_intrinsic_load_barycentric_sample => (),
2561 nir_intrinsic_load_global | nir_intrinsic_load_global_constant => {
2562 let size_B =
2563 (intrin.def.bit_size() / 8) * intrin.def.num_components();
2564 assert!(u32::from(size_B) <= intrin.align());
2565 let order =
2566 if intrin.intrinsic == nir_intrinsic_load_global_constant {
2567 MemOrder::Constant
2568 } else {
2569 MemOrder::Strong(MemScope::System)
2570 };
2571 let access = MemAccess {
2572 mem_type: MemType::from_size(size_B, false),
2573 space: MemSpace::Global(MemAddrType::A64),
2574 order: order,
2575 eviction_priority: self
2576 .get_eviction_priority(intrin.access()),
2577 };
2578 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2579 let dst = b.alloc_ssa(RegFile::GPR, size_B.div_ceil(4));
2580
2581 b.push_op(OpLd {
2582 dst: dst.into(),
2583 addr: addr,
2584 offset: offset,
2585 access: access,
2586 });
2587 self.set_dst(&intrin.def, dst);
2588 }
2589 nir_intrinsic_ldtram_nv => {
2590 let ShaderIoInfo::Fragment(io) = &mut self.info.io else {
2591 panic!("ldtram_nv is only used for fragment shaders");
2592 };
2593
2594 assert!(
2595 intrin.def.bit_size() == 32
2596 && intrin.def.num_components == 2
2597 );
2598
2599 let flags = intrin.flags();
2600 let use_c = flags != 0;
2601
2602 let addr = u16::try_from(intrin.base()).unwrap();
2603
2604 io.mark_barycentric_attr_in(addr);
2605
2606 let dst = b.alloc_ssa(RegFile::GPR, 2);
2607 b.push_op(OpLdTram {
2608 dst: dst.into(),
2609 addr,
2610 use_c,
2611 });
2612 self.set_dst(&intrin.def, dst);
2613 }
2614 nir_intrinsic_load_sample_id => {
2615 let dst = b.alloc_ssa(RegFile::GPR, 1);
2616 b.push_op(OpPixLd {
2617 dst: dst.into(),
2618 val: PixVal::MyIndex,
2619 });
2620 self.set_dst(&intrin.def, dst);
2621 }
2622 nir_intrinsic_load_sample_mask_in => {
2623 if let ShaderIoInfo::Fragment(info) = &mut self.info.io {
2624 info.reads_sample_mask = true;
2625 } else {
2626 panic!(
2627 "sample_mask_in is only available in fragment shaders"
2628 );
2629 }
2630
2631 let dst = b.alloc_ssa(RegFile::GPR, 1);
2632 b.push_op(OpPixLd {
2633 dst: dst.into(),
2634 val: PixVal::CovMask,
2635 });
2636 self.set_dst(&intrin.def, dst);
2637 }
2638 nir_intrinsic_load_tess_coord_xy => {
2639 // Loading gl_TessCoord in tessellation evaluation shaders is
2640 // weird. It's treated as a per-vertex output which is indexed
2641 // by LANEID.
2642 match &self.info.stage {
2643 ShaderStageInfo::Tessellation(_) => (),
2644 _ => panic!(
2645 "load_tess_coord is only available in tessellation \
2646 shaders"
2647 ),
2648 };
2649
2650 assert!(intrin.def.bit_size() == 32);
2651 assert!(intrin.def.num_components() == 2);
2652
2653 let vtx = b.alloc_ssa(RegFile::GPR, 1);
2654 b.push_op(OpS2R {
2655 dst: vtx.into(),
2656 idx: 0,
2657 });
2658
2659 let access = AttrAccess {
2660 addr: NAK_ATTR_TESS_COORD,
2661 comps: 2,
2662 patch: false,
2663 output: true,
2664 phys: false,
2665 };
2666
2667 // This is recorded as a patch output in parse_shader() because
2668 // the hardware requires it be in the SPH, whether we use it or
2669 // not.
2670
2671 let dst = b.alloc_ssa(RegFile::GPR, access.comps);
2672 b.push_op(OpALd {
2673 dst: dst.into(),
2674 vtx: vtx.into(),
2675 offset: 0.into(),
2676 access: access,
2677 });
2678 self.set_dst(&intrin.def, dst);
2679 }
2680 nir_intrinsic_load_scratch => {
2681 let size_B =
2682 (intrin.def.bit_size() / 8) * intrin.def.num_components();
2683 assert!(u32::from(size_B) <= intrin.align());
2684 let access = MemAccess {
2685 mem_type: MemType::from_size(size_B, false),
2686 space: MemSpace::Local,
2687 order: MemOrder::Strong(MemScope::CTA),
2688 eviction_priority: MemEvictionPriority::Normal,
2689 };
2690 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2691 let dst = b.alloc_ssa(RegFile::GPR, size_B.div_ceil(4));
2692
2693 b.push_op(OpLd {
2694 dst: dst.into(),
2695 addr: addr,
2696 offset: offset,
2697 access: access,
2698 });
2699 self.set_dst(&intrin.def, dst);
2700 }
2701 nir_intrinsic_load_shared => {
2702 let size_B =
2703 (intrin.def.bit_size() / 8) * intrin.def.num_components();
2704 assert!(u32::from(size_B) <= intrin.align());
2705 let access = MemAccess {
2706 mem_type: MemType::from_size(size_B, false),
2707 space: MemSpace::Shared,
2708 order: MemOrder::Strong(MemScope::CTA),
2709 eviction_priority: MemEvictionPriority::Normal,
2710 };
2711 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2712 let offset = offset + intrin.base();
2713 let dst = b.alloc_ssa(RegFile::GPR, size_B.div_ceil(4));
2714
2715 b.push_op(OpLd {
2716 dst: dst.into(),
2717 addr: addr,
2718 offset: offset,
2719 access: access,
2720 });
2721 self.set_dst(&intrin.def, dst);
2722 }
2723 nir_intrinsic_load_sysval_nv => {
2724 let idx = u8::try_from(intrin.base()).unwrap();
2725 debug_assert!(intrin.def.num_components == 1);
2726 debug_assert!(
2727 intrin.def.bit_size == 32 || intrin.def.bit_size == 64
2728 );
2729 let comps = intrin.def.bit_size / 32;
2730 let dst = b.alloc_ssa(RegFile::GPR, comps);
2731 if idx == NAK_SV_CLOCK || idx == NAK_SV_CLOCK + 1 {
2732 debug_assert!(idx + comps <= NAK_SV_CLOCK + 2);
2733 b.push_op(OpCS2R {
2734 dst: dst.into(),
2735 idx: idx,
2736 });
2737 } else {
2738 debug_assert!(intrin.def.bit_size == 32);
2739 b.push_op(OpS2R {
2740 dst: dst.into(),
2741 idx: idx,
2742 });
2743 }
2744 self.set_dst(&intrin.def, dst);
2745 }
2746 nir_intrinsic_ldc_nv => {
2747 let size_B =
2748 (intrin.def.bit_size() / 8) * intrin.def.num_components();
2749 let idx = &srcs[0];
2750
2751 let (off, off_imm) = self.get_cbuf_addr_offset(&srcs[1]);
2752
2753 let dst = b.alloc_ssa(RegFile::GPR, size_B.div_ceil(4));
2754
2755 if let Some(idx_imm) = idx.as_uint() {
2756 let idx_imm: u8 = idx_imm.try_into().unwrap();
2757 let cb = CBufRef {
2758 buf: CBuf::Binding(idx_imm),
2759 offset: off_imm,
2760 };
2761 if off.is_zero() {
2762 for (i, comp) in dst.iter().enumerate() {
2763 let i = u16::try_from(i).unwrap();
2764 b.copy_to((*comp).into(), cb.offset(i * 4).into());
2765 }
2766 } else {
2767 b.push_op(OpLdc {
2768 dst: dst.into(),
2769 cb: cb.into(),
2770 offset: off,
2771 mode: LdcMode::Indexed,
2772 mem_type: MemType::from_size(size_B, false),
2773 });
2774 }
2775 } else {
2776 // In the IndexedSegmented mode, the hardware computes the
2777 // actual index and offset as follows:
2778 //
2779 // idx = imm_idx + reg[31:16]
2780 // offset = imm_offset + reg[15:0]
2781 // ldc c[idx][offset]
2782 //
2783 // So pack the index and offset accordingly
2784 let idx = self.get_src(idx);
2785 let off_idx = b.prmt(off, idx, [0, 1, 4, 5]);
2786 let cb = CBufRef {
2787 buf: CBuf::Binding(0),
2788 offset: off_imm,
2789 };
2790 b.push_op(OpLdc {
2791 dst: dst.into(),
2792 cb: cb.into(),
2793 offset: off_idx.into(),
2794 mode: LdcMode::IndexedSegmented,
2795 mem_type: MemType::from_size(size_B, false),
2796 });
2797 }
2798 self.set_dst(&intrin.def, dst);
2799 }
2800 nir_intrinsic_ldcx_nv => {
2801 let size_B =
2802 (intrin.def.bit_size() / 8) * intrin.def.num_components();
2803
2804 let handle = self.get_ssa_ref(&srcs[0]);
2805 let (off, off_imm) = self.get_cbuf_addr_offset(&srcs[1]);
2806
2807 let cb = CBufRef {
2808 buf: CBuf::BindlessSSA(handle),
2809 offset: off_imm,
2810 };
2811
2812 let dst = b.alloc_ssa(RegFile::GPR, size_B.div_ceil(4));
2813 if off.is_zero() {
2814 for (i, comp) in dst.iter().enumerate() {
2815 let i = u16::try_from(i).unwrap();
2816 b.copy_to((*comp).into(), cb.offset(i * 4).into());
2817 }
2818 } else {
2819 b.push_op(OpLdc {
2820 dst: dst.into(),
2821 cb: cb.into(),
2822 offset: off,
2823 mode: LdcMode::Indexed,
2824 mem_type: MemType::from_size(size_B, false),
2825 });
2826 }
2827 self.set_dst(&intrin.def, dst);
2828 }
2829 nir_intrinsic_pin_cx_handle_nv => {
2830 let handle = self.get_ssa_ref(&srcs[0]);
2831 b.push_op(OpPin {
2832 src: handle.into(),
2833 dst: handle.into(),
2834 });
2835 }
2836 nir_intrinsic_unpin_cx_handle_nv => {
2837 let handle = self.get_ssa_ref(&srcs[0]);
2838 b.push_op(OpUnpin {
2839 src: handle.into(),
2840 dst: handle.into(),
2841 });
2842 }
2843 nir_intrinsic_barrier => {
2844 let modes = intrin.memory_modes();
2845 let semantics = intrin.memory_semantics();
2846 if (modes & nir_var_mem_global) != 0
2847 && (semantics & NIR_MEMORY_RELEASE) != 0
2848 {
2849 // Pre-Volta doesn't have WBAll but it also seems that we
2850 // don't need it.
2851 if self.sm.sm() >= 70 {
2852 b.push_op(OpCCtl {
2853 op: CCtlOp::WBAll,
2854 mem_space: MemSpace::Global(MemAddrType::A64),
2855 addr: 0.into(),
2856 addr_offset: 0,
2857 });
2858 }
2859 }
2860 match intrin.execution_scope() {
2861 SCOPE_NONE => (),
2862 SCOPE_WORKGROUP => {
2863 assert!(
2864 self.nir.info.stage() == MESA_SHADER_COMPUTE
2865 || self.nir.info.stage() == MESA_SHADER_KERNEL
2866 );
2867 self.info.num_control_barriers = 1;
2868 b.push_op(OpBar {});
2869 }
2870 _ => panic!("Unhandled execution scope"),
2871 }
2872 if intrin.memory_scope() != SCOPE_NONE {
2873 let mem_scope = match intrin.memory_scope() {
2874 SCOPE_INVOCATION | SCOPE_SUBGROUP => MemScope::CTA,
2875 SCOPE_WORKGROUP | SCOPE_QUEUE_FAMILY | SCOPE_DEVICE => {
2876 MemScope::GPU
2877 }
2878 _ => panic!("Unhandled memory scope"),
2879 };
2880 b.push_op(OpMemBar { scope: mem_scope });
2881 }
2882 if (modes & nir_var_mem_global) != 0
2883 && (semantics & NIR_MEMORY_ACQUIRE) != 0
2884 {
2885 b.push_op(OpCCtl {
2886 op: CCtlOp::IVAll,
2887 mem_space: MemSpace::Global(MemAddrType::A64),
2888 addr: 0.into(),
2889 addr_offset: 0,
2890 });
2891 }
2892 }
2893 nir_intrinsic_quad_broadcast
2894 | nir_intrinsic_read_invocation
2895 | nir_intrinsic_shuffle
2896 | nir_intrinsic_shuffle_down
2897 | nir_intrinsic_shuffle_up
2898 | nir_intrinsic_shuffle_xor => {
2899 assert!(srcs[0].bit_size() == 32);
2900 assert!(srcs[0].num_components() == 1);
2901 let data = self.get_src(&srcs[0]);
2902
2903 assert!(srcs[1].bit_size() == 32);
2904 let idx = self.get_src(&srcs[1]);
2905
2906 assert!(intrin.def.bit_size() == 32);
2907 let dst = b.alloc_ssa(RegFile::GPR, 1);
2908
2909 b.push_op(OpShfl {
2910 dst: dst.into(),
2911 in_bounds: Dst::None,
2912 src: data,
2913 lane: idx,
2914 c: match intrin.intrinsic {
2915 nir_intrinsic_quad_broadcast => 0x1c_03.into(),
2916 nir_intrinsic_shuffle_up => 0.into(),
2917 _ => 0x1f.into(),
2918 },
2919 op: match intrin.intrinsic {
2920 nir_intrinsic_shuffle_down => ShflOp::Down,
2921 nir_intrinsic_shuffle_up => ShflOp::Up,
2922 nir_intrinsic_shuffle_xor => ShflOp::Bfly,
2923 _ => ShflOp::Idx,
2924 },
2925 });
2926 self.set_dst(&intrin.def, dst);
2927 }
2928 nir_intrinsic_quad_swap_horizontal
2929 | nir_intrinsic_quad_swap_vertical
2930 | nir_intrinsic_quad_swap_diagonal => {
2931 assert!(srcs[0].bit_size() == 32);
2932 assert!(srcs[0].num_components() == 1);
2933 let data = self.get_src(&srcs[0]);
2934
2935 assert!(intrin.def.bit_size() == 32);
2936 let dst = b.alloc_ssa(RegFile::GPR, 1);
2937 b.push_op(OpShfl {
2938 dst: dst.into(),
2939 in_bounds: Dst::None,
2940 src: data,
2941 lane: match intrin.intrinsic {
2942 nir_intrinsic_quad_swap_horizontal => 1_u32.into(),
2943 nir_intrinsic_quad_swap_vertical => 2_u32.into(),
2944 nir_intrinsic_quad_swap_diagonal => 3_u32.into(),
2945 op => panic!("Unknown quad intrinsic {}", op),
2946 },
2947 c: 0x1c_03.into(),
2948 op: ShflOp::Bfly,
2949 });
2950 self.set_dst(&intrin.def, dst);
2951 }
2952 nir_intrinsic_shared_atomic => {
2953 let bit_size = intrin.def.bit_size();
2954 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2955 let data = self.get_src(&srcs[1]);
2956 let atom_type = self.get_atomic_type(intrin);
2957 let atom_op = self.get_atomic_op(intrin, AtomCmpSrc::Separate);
2958
2959 assert!(intrin.def.num_components() == 1);
2960 let dst = b.alloc_ssa(RegFile::GPR, bit_size.div_ceil(32));
2961
2962 b.push_op(OpAtom {
2963 dst: dst.into(),
2964 addr: addr,
2965 cmpr: 0.into(),
2966 data: data,
2967 atom_op: atom_op,
2968 atom_type: atom_type,
2969 addr_offset: offset,
2970 mem_space: MemSpace::Shared,
2971 mem_order: MemOrder::Strong(MemScope::CTA),
2972 mem_eviction_priority: MemEvictionPriority::Normal,
2973 });
2974 self.set_dst(&intrin.def, dst);
2975 }
2976 nir_intrinsic_shared_atomic_swap => {
2977 assert!(intrin.atomic_op() == nir_atomic_op_cmpxchg);
2978 let bit_size = intrin.def.bit_size();
2979 let (addr, offset) = self.get_io_addr_offset(&srcs[0], 24);
2980 let cmpr = self.get_src(&srcs[1]);
2981 let data = self.get_src(&srcs[2]);
2982 let atom_type = AtomType::U(bit_size);
2983
2984 assert!(intrin.def.num_components() == 1);
2985 let dst = b.alloc_ssa(RegFile::GPR, bit_size.div_ceil(32));
2986
2987 b.push_op(OpAtom {
2988 dst: dst.into(),
2989 addr: addr,
2990 cmpr: cmpr,
2991 data: data,
2992 atom_op: AtomOp::CmpExch(AtomCmpSrc::Separate),
2993 atom_type: atom_type,
2994 addr_offset: offset,
2995 mem_space: MemSpace::Shared,
2996 mem_order: MemOrder::Strong(MemScope::CTA),
2997 mem_eviction_priority: MemEvictionPriority::Normal,
2998 });
2999 self.set_dst(&intrin.def, dst);
3000 }
3001 nir_intrinsic_ssa_bar_nv => {
3002 let src = self.get_src(&srcs[0]);
3003 b.push_op(OpSrcBar { src });
3004 }
3005 nir_intrinsic_store_global => {
3006 let data = self.get_src(&srcs[0]);
3007 let size_B =
3008 (srcs[0].bit_size() / 8) * srcs[0].num_components();
3009 assert!(u32::from(size_B) <= intrin.align());
3010 let access = MemAccess {
3011 mem_type: MemType::from_size(size_B, false),
3012 space: MemSpace::Global(MemAddrType::A64),
3013 order: MemOrder::Strong(MemScope::System),
3014 eviction_priority: self
3015 .get_eviction_priority(intrin.access()),
3016 };
3017 let (addr, offset) = self.get_io_addr_offset(&srcs[1], 24);
3018
3019 b.push_op(OpSt {
3020 addr: addr,
3021 data: data,
3022 offset: offset,
3023 access: access,
3024 });
3025 }
3026 nir_intrinsic_fs_out_nv => {
3027 let data = self.get_ssa(srcs[0].as_def());
3028 assert!(data.len() == 1);
3029 let data = data[0];
3030
3031 let addr = u16::try_from(intrin.base()).unwrap();
3032 assert!(addr % 4 == 0);
3033
3034 self.fs_out_regs[usize::from(addr / 4)] = data;
3035 }
3036 nir_intrinsic_store_scratch => {
3037 let data = self.get_src(&srcs[0]);
3038 let size_B =
3039 (srcs[0].bit_size() / 8) * srcs[0].num_components();
3040 assert!(u32::from(size_B) <= intrin.align());
3041 let access = MemAccess {
3042 mem_type: MemType::from_size(size_B, false),
3043 space: MemSpace::Local,
3044 order: MemOrder::Strong(MemScope::CTA),
3045 eviction_priority: MemEvictionPriority::Normal,
3046 };
3047 let (addr, offset) = self.get_io_addr_offset(&srcs[1], 24);
3048
3049 b.push_op(OpSt {
3050 addr: addr,
3051 data: data,
3052 offset: offset,
3053 access: access,
3054 });
3055 }
3056 nir_intrinsic_store_shared => {
3057 let data = self.get_src(&srcs[0]);
3058 let size_B =
3059 (srcs[0].bit_size() / 8) * srcs[0].num_components();
3060 assert!(u32::from(size_B) <= intrin.align());
3061 let access = MemAccess {
3062 mem_type: MemType::from_size(size_B, false),
3063 space: MemSpace::Shared,
3064 order: MemOrder::Strong(MemScope::CTA),
3065 eviction_priority: MemEvictionPriority::Normal,
3066 };
3067 let (addr, offset) = self.get_io_addr_offset(&srcs[1], 24);
3068 let offset = offset + intrin.base();
3069
3070 b.push_op(OpSt {
3071 addr: addr,
3072 data: data,
3073 offset: offset,
3074 access: access,
3075 });
3076 }
3077 nir_intrinsic_emit_vertex_nv | nir_intrinsic_end_primitive_nv => {
3078 assert!(intrin.def.bit_size() == 32);
3079 assert!(intrin.def.num_components() == 1);
3080
3081 let dst = b.alloc_ssa(RegFile::GPR, 1);
3082 let handle = self.get_src(&srcs[0]);
3083 let stream_id = intrin.stream_id();
3084
3085 b.push_op(OpOut {
3086 dst: dst.into(),
3087 handle: handle,
3088 stream: stream_id.into(),
3089 out_type: if intrin.intrinsic
3090 == nir_intrinsic_emit_vertex_nv
3091 {
3092 OutType::Emit
3093 } else {
3094 OutType::Cut
3095 },
3096 });
3097 self.set_dst(&intrin.def, dst);
3098 }
3099
3100 nir_intrinsic_final_primitive_nv => {
3101 let handle = self.get_src(&srcs[0]);
3102
3103 if self.sm.sm() >= 70 {
3104 b.push_op(OpOutFinal { handle: handle });
3105 } else {
3106 b.push_op(OpRegOut { srcs: vec![handle] });
3107 }
3108 }
3109 nir_intrinsic_vote_all
3110 | nir_intrinsic_vote_any
3111 | nir_intrinsic_vote_ieq => {
3112 assert!(srcs[0].bit_size() == 1);
3113 let src = self.get_src(&srcs[0]);
3114
3115 assert!(intrin.def.bit_size() == 1);
3116 let dst = b.alloc_ssa(RegFile::Pred, 1);
3117
3118 b.push_op(OpVote {
3119 op: match intrin.intrinsic {
3120 nir_intrinsic_vote_all => VoteOp::All,
3121 nir_intrinsic_vote_any => VoteOp::Any,
3122 nir_intrinsic_vote_ieq => VoteOp::Eq,
3123 _ => panic!("Unknown vote intrinsic"),
3124 },
3125 ballot: Dst::None,
3126 vote: dst.into(),
3127 pred: src,
3128 });
3129 self.set_dst(&intrin.def, dst);
3130 }
3131 nir_intrinsic_is_sparse_texels_resident => {
3132 let src = self.get_src(&srcs[0]);
3133 let dst = b.isetp(IntCmpType::I32, IntCmpOp::Ne, src, 0.into());
3134 self.set_dst(&intrin.def, dst);
3135 }
3136 _ => panic!(
3137 "Unsupported intrinsic instruction: {}",
3138 intrin.info().name()
3139 ),
3140 }
3141 }
3142
parse_load_const( &mut self, b: &mut impl SSABuilder, load_const: &nir_load_const_instr, )3143 fn parse_load_const(
3144 &mut self,
3145 b: &mut impl SSABuilder,
3146 load_const: &nir_load_const_instr,
3147 ) {
3148 let values = &load_const.values();
3149
3150 let mut dst = Vec::new();
3151 match load_const.def.bit_size {
3152 1 => {
3153 for c in 0..load_const.def.num_components {
3154 let imm_b1 = unsafe { values[usize::from(c)].b };
3155 dst.push(b.copy(imm_b1.into())[0]);
3156 }
3157 }
3158 8 => {
3159 for dw in 0..load_const.def.num_components.div_ceil(4) {
3160 let mut imm_u32 = 0;
3161 for b in 0..4 {
3162 let c = dw * 4 + b;
3163 if c < load_const.def.num_components {
3164 let imm_u8 = unsafe { values[usize::from(c)].u8_ };
3165 imm_u32 |= u32::from(imm_u8) << b * 8;
3166 }
3167 }
3168 dst.push(b.copy(imm_u32.into())[0]);
3169 }
3170 }
3171 16 => {
3172 for dw in 0..load_const.def.num_components.div_ceil(2) {
3173 let mut imm_u32 = 0;
3174 for w in 0..2 {
3175 let c = dw * 2 + w;
3176 if c < load_const.def.num_components {
3177 let imm_u16 =
3178 unsafe { values[usize::from(c)].u16_ };
3179 imm_u32 |= u32::from(imm_u16) << w * 16;
3180 }
3181 }
3182 dst.push(b.copy(imm_u32.into())[0]);
3183 }
3184 }
3185 32 => {
3186 for c in 0..load_const.def.num_components {
3187 let imm_u32 = unsafe { values[usize::from(c)].u32_ };
3188 dst.push(b.copy(imm_u32.into())[0]);
3189 }
3190 }
3191 64 => {
3192 for c in 0..load_const.def.num_components {
3193 let imm_u64 = unsafe { values[c as usize].u64_ };
3194 dst.push(b.copy((imm_u64 as u32).into())[0]);
3195 dst.push(b.copy(((imm_u64 >> 32) as u32).into())[0]);
3196 }
3197 }
3198 _ => panic!("Unknown bit size: {}", load_const.def.bit_size),
3199 }
3200
3201 self.set_ssa(&load_const.def, dst);
3202 }
3203
parse_undef( &mut self, b: &mut impl SSABuilder, undef: &nir_undef_instr, )3204 fn parse_undef(
3205 &mut self,
3206 b: &mut impl SSABuilder,
3207 undef: &nir_undef_instr,
3208 ) {
3209 let dst = alloc_ssa_for_nir(b, &undef.def);
3210 for c in &dst {
3211 b.push_op(OpUndef { dst: (*c).into() });
3212 }
3213 self.set_ssa(&undef.def, dst);
3214 }
3215
emit_jump( &mut self, b: &mut impl SSABuilder, nb: &nir_block, target: &nir_block, )3216 fn emit_jump(
3217 &mut self,
3218 b: &mut impl SSABuilder,
3219 nb: &nir_block,
3220 target: &nir_block,
3221 ) {
3222 if target.index == self.end_block_id {
3223 b.push_op(OpExit {});
3224 } else {
3225 self.cfg.add_edge(nb.index, target.index);
3226 let target_label = self.get_block_label(target);
3227
3228 match self.peek_crs(target) {
3229 Some(SyncType::Sync) => {
3230 b.push_op(OpSync {
3231 target: target_label,
3232 });
3233 }
3234 Some(SyncType::Brk) => {
3235 b.push_op(OpBrk {
3236 target: target_label,
3237 });
3238 }
3239 Some(SyncType::Cont) => {
3240 b.push_op(OpCont {
3241 target: target_label,
3242 });
3243 }
3244 None => {
3245 b.push_op(OpBra {
3246 target: target_label,
3247 });
3248 }
3249 }
3250 }
3251 }
3252
emit_pred_jump( &mut self, b: &mut impl SSABuilder, nb: &nir_block, pred: Pred, target: &nir_block, fallthrough: &nir_block, )3253 fn emit_pred_jump(
3254 &mut self,
3255 b: &mut impl SSABuilder,
3256 nb: &nir_block,
3257 pred: Pred,
3258 target: &nir_block,
3259 fallthrough: &nir_block,
3260 ) {
3261 // The fall-through edge has to come first
3262 self.cfg.add_edge(nb.index, fallthrough.index);
3263 let op = if target.index == self.end_block_id {
3264 Op::Exit(OpExit {})
3265 } else {
3266 self.cfg.add_edge(nb.index, target.index);
3267 Op::Bra(OpBra {
3268 target: self.get_block_label(target),
3269 })
3270 };
3271 b.predicate(pred).push_op(op);
3272 }
3273
parse_block( &mut self, ssa_alloc: &mut SSAValueAllocator, phi_map: &mut PhiAllocMap, nb: &nir_block, )3274 fn parse_block(
3275 &mut self,
3276 ssa_alloc: &mut SSAValueAllocator,
3277 phi_map: &mut PhiAllocMap,
3278 nb: &nir_block,
3279 ) {
3280 let sm = self.sm;
3281 let mut b = SSAInstrBuilder::new(sm, ssa_alloc);
3282
3283 if self.sm.sm() >= 70 && nb.index == 0 && self.nir.info.shared_size > 0
3284 {
3285 // The blob seems to always do a BSYNC before accessing shared
3286 // memory. Perhaps this is to ensure that our allocation is
3287 // actually available and not in use by another thread?
3288 let label = self.label_alloc.alloc();
3289 let bar_clear = b.alloc_ssa(RegFile::Bar, 1);
3290
3291 b.push_op(OpBClear {
3292 dst: bar_clear.into(),
3293 });
3294
3295 let bar = b.alloc_ssa(RegFile::Bar, 1);
3296 b.push_op(OpBSSy {
3297 bar_out: bar.into(),
3298 bar_in: bar_clear.into(),
3299 cond: SrcRef::True.into(),
3300 target: label,
3301 });
3302
3303 b.push_op(OpBSync {
3304 bar: bar.into(),
3305 cond: SrcRef::True.into(),
3306 });
3307
3308 b.push_op(OpNop { label: Some(label) });
3309 }
3310
3311 let mut phi = OpPhiDsts::new();
3312 for ni in nb.iter_instr_list() {
3313 let Some(np) = ni.as_phi() else {
3314 break;
3315 };
3316
3317 if DEBUG.annotate() {
3318 let annotation = self
3319 .nir_instr_printer
3320 .instr_to_string(ni)
3321 .unwrap()
3322 .split_whitespace()
3323 .collect::<Vec<_>>()
3324 .join(" ");
3325 b.push_op(OpAnnotate {
3326 annotation: format!("generated by \"{}\"", annotation,),
3327 });
3328 }
3329
3330 let uniform = !nb.divergent
3331 && self.sm.sm() >= 75
3332 && !DEBUG.no_ugpr()
3333 && !np.def.divergent;
3334
3335 // This should be ensured by nak_nir_lower_cf()
3336 if uniform {
3337 for ps in np.iter_srcs() {
3338 assert!(!ps.pred().divergent);
3339 }
3340 }
3341
3342 let mut b = UniformBuilder::new(&mut b, uniform);
3343 let dst = alloc_ssa_for_nir(&mut b, np.def.as_def());
3344 for i in 0..dst.len() {
3345 let phi_id = phi_map.get_phi_id(np, i.try_into().unwrap());
3346 phi.dsts.push(phi_id, dst[i].into());
3347 }
3348 self.set_ssa(np.def.as_def(), dst);
3349 }
3350
3351 if !phi.dsts.is_empty() {
3352 b.push_op(phi);
3353 }
3354
3355 if self.sm.sm() < 75 && nb.cf_node.prev().is_none() {
3356 if let Some(_) = nb.parent().as_loop() {
3357 b.push_op(OpPCnt {
3358 target: self.get_block_label(nb),
3359 });
3360 self.push_crs(nb, SyncType::Cont);
3361 }
3362 }
3363
3364 let mut goto = None;
3365 for ni in nb.iter_instr_list() {
3366 if DEBUG.annotate() && ni.type_ != nir_instr_type_phi {
3367 let annotation = self
3368 .nir_instr_printer
3369 .instr_to_string(ni)
3370 .unwrap()
3371 .split_whitespace()
3372 .collect::<Vec<_>>()
3373 .join(" ");
3374 b.push_op(OpAnnotate {
3375 annotation: format!("generated by \"{}\"", annotation,),
3376 });
3377 }
3378
3379 let uniform = !nb.divergent
3380 && self.sm.sm() >= 75
3381 && !DEBUG.no_ugpr()
3382 && ni.def().is_some_and(|d| !d.divergent);
3383 let mut b = UniformBuilder::new(&mut b, uniform);
3384
3385 match ni.type_ {
3386 nir_instr_type_alu => {
3387 self.parse_alu(&mut b, ni.as_alu().unwrap())
3388 }
3389 nir_instr_type_jump => {
3390 let jump = ni.as_jump().unwrap();
3391 if jump.type_ == nir_jump_goto
3392 || jump.type_ == nir_jump_goto_if
3393 {
3394 goto = Some(jump);
3395 }
3396 }
3397 nir_instr_type_tex => {
3398 self.parse_tex(&mut b, ni.as_tex().unwrap())
3399 }
3400 nir_instr_type_intrinsic => {
3401 self.parse_intrinsic(&mut b, ni.as_intrinsic().unwrap())
3402 }
3403 nir_instr_type_load_const => {
3404 self.parse_load_const(&mut b, ni.as_load_const().unwrap())
3405 }
3406 nir_instr_type_undef => {
3407 self.parse_undef(&mut b, ni.as_undef().unwrap())
3408 }
3409 nir_instr_type_phi => (),
3410 _ => panic!("Unsupported instruction type"),
3411 }
3412 }
3413
3414 if self.sm.sm() < 70 {
3415 if let Some(ni) = nb.following_if() {
3416 let fb = ni.following_block();
3417 b.push_op(OpSSy {
3418 target: self.get_block_label(fb),
3419 });
3420 self.push_crs(fb, SyncType::Sync);
3421 } else if let Some(nl) = nb.following_loop() {
3422 let fb = nl.following_block();
3423 b.push_op(OpPBk {
3424 target: self.get_block_label(fb),
3425 });
3426 self.push_crs(fb, SyncType::Brk);
3427 }
3428 }
3429
3430 let succ = nb.successors();
3431 for sb in succ {
3432 let sb = match sb {
3433 Some(b) => b,
3434 None => continue,
3435 };
3436
3437 let mut phi = OpPhiSrcs::new();
3438
3439 for ni in sb.iter_instr_list() {
3440 let Some(np) = ni.as_phi() else {
3441 break;
3442 };
3443
3444 if DEBUG.annotate() {
3445 let annotation = self
3446 .nir_instr_printer
3447 .instr_to_string(ni)
3448 .unwrap()
3449 .split_whitespace()
3450 .collect::<Vec<_>>()
3451 .join(" ");
3452 b.push_op(OpAnnotate {
3453 annotation: format!("generated by \"{}\"", annotation,),
3454 });
3455 }
3456
3457 for ps in np.iter_srcs() {
3458 if ps.pred().index == nb.index {
3459 let src = *self.get_src(&ps.src).as_ssa().unwrap();
3460 for (i, src) in src.iter().enumerate() {
3461 let phi_id =
3462 phi_map.get_phi_id(np, i.try_into().unwrap());
3463 phi.srcs.push(phi_id, (*src).into());
3464 }
3465 break;
3466 }
3467 }
3468 }
3469
3470 if !phi.srcs.is_empty() {
3471 b.push_op(phi);
3472 }
3473 }
3474
3475 if let Some(goto) = goto {
3476 let target = goto.target().unwrap();
3477 if goto.type_ == nir_jump_goto {
3478 self.emit_jump(&mut b, nb, target);
3479 } else {
3480 let cond = self.get_ssa(goto.condition.as_def())[0];
3481 let else_target = goto.else_target().unwrap();
3482
3483 /* Next block in the NIR CF list */
3484 let next_block = nb.cf_node.next().unwrap().as_block().unwrap();
3485
3486 if else_target as *const _ == next_block as *const _ {
3487 self.emit_pred_jump(
3488 &mut b,
3489 nb,
3490 // This is the branch to jump to the else
3491 cond.into(),
3492 target,
3493 else_target,
3494 );
3495 } else if target as *const _ == next_block as *const _ {
3496 self.emit_pred_jump(
3497 &mut b,
3498 nb,
3499 Pred::from(cond).bnot(),
3500 else_target,
3501 target,
3502 );
3503 } else {
3504 panic!(
3505 "One of the two goto targets must be the next block in \
3506 the NIR CF list"
3507 );
3508 }
3509 }
3510 } else {
3511 if let Some(ni) = nb.following_if() {
3512 let cond = self.get_ssa(ni.condition.as_def())[0];
3513 self.emit_pred_jump(
3514 &mut b,
3515 nb,
3516 // This is the branch to jump to the else
3517 Pred::from(cond).bnot(),
3518 ni.first_else_block(),
3519 ni.first_then_block(),
3520 );
3521 } else {
3522 assert!(succ[1].is_none());
3523 let s0 = succ[0].unwrap();
3524 self.emit_jump(&mut b, nb, s0);
3525 }
3526 }
3527
3528 let bb = BasicBlock {
3529 label: self.get_block_label(nb),
3530 uniform: !nb.divergent,
3531 instrs: b.as_vec(),
3532 };
3533 self.cfg.add_node(nb.index, bb);
3534 }
3535
parse_if( &mut self, ssa_alloc: &mut SSAValueAllocator, phi_map: &mut PhiAllocMap, ni: &nir_if, )3536 fn parse_if(
3537 &mut self,
3538 ssa_alloc: &mut SSAValueAllocator,
3539 phi_map: &mut PhiAllocMap,
3540 ni: &nir_if,
3541 ) {
3542 self.parse_cf_list(ssa_alloc, phi_map, ni.iter_then_list());
3543 self.parse_cf_list(ssa_alloc, phi_map, ni.iter_else_list());
3544
3545 if self.sm.sm() < 70 {
3546 let next_block = ni.cf_node.next().unwrap().as_block().unwrap();
3547 self.pop_crs(next_block, SyncType::Sync);
3548 }
3549 }
3550
parse_loop( &mut self, ssa_alloc: &mut SSAValueAllocator, phi_map: &mut PhiAllocMap, nl: &nir_loop, )3551 fn parse_loop(
3552 &mut self,
3553 ssa_alloc: &mut SSAValueAllocator,
3554 phi_map: &mut PhiAllocMap,
3555 nl: &nir_loop,
3556 ) {
3557 self.parse_cf_list(ssa_alloc, phi_map, nl.iter_body());
3558
3559 if self.sm.sm() < 70 {
3560 let header = nl.iter_body().next().unwrap().as_block().unwrap();
3561 self.pop_crs(header, SyncType::Cont);
3562 let next_block = nl.cf_node.next().unwrap().as_block().unwrap();
3563 self.pop_crs(next_block, SyncType::Brk);
3564 }
3565 }
3566
parse_cf_list( &mut self, ssa_alloc: &mut SSAValueAllocator, phi_map: &mut PhiAllocMap, list: ExecListIter<nir_cf_node>, )3567 fn parse_cf_list(
3568 &mut self,
3569 ssa_alloc: &mut SSAValueAllocator,
3570 phi_map: &mut PhiAllocMap,
3571 list: ExecListIter<nir_cf_node>,
3572 ) {
3573 for node in list {
3574 match node.type_ {
3575 nir_cf_node_block => {
3576 let nb = node.as_block().unwrap();
3577 self.parse_block(ssa_alloc, phi_map, nb);
3578 }
3579 nir_cf_node_if => {
3580 let ni = node.as_if().unwrap();
3581 self.parse_if(ssa_alloc, phi_map, ni);
3582 }
3583 nir_cf_node_loop => {
3584 let nl = node.as_loop().unwrap();
3585 self.parse_loop(ssa_alloc, phi_map, nl);
3586 }
3587 _ => panic!("Invalid inner CF node type"),
3588 }
3589 }
3590 }
3591
parse_function_impl(&mut self, nfi: &nir_function_impl) -> Function3592 pub fn parse_function_impl(&mut self, nfi: &nir_function_impl) -> Function {
3593 let mut ssa_alloc = SSAValueAllocator::new();
3594 let end_nb = nfi.end_block();
3595 self.end_block_id = end_nb.index;
3596
3597 let mut phi_alloc = PhiAllocator::new();
3598 let mut phi_map = PhiAllocMap::new(&mut phi_alloc);
3599
3600 self.parse_cf_list(&mut ssa_alloc, &mut phi_map, nfi.iter_body());
3601
3602 let cfg = std::mem::take(&mut self.cfg).as_cfg();
3603 assert!(cfg.len() > 0);
3604 for i in 0..cfg.len() {
3605 if cfg[i].falls_through() {
3606 assert!(cfg.succ_indices(i)[0] == i + 1);
3607 }
3608 }
3609
3610 let mut f = Function {
3611 ssa_alloc: ssa_alloc,
3612 phi_alloc: phi_alloc,
3613 blocks: cfg,
3614 };
3615 f.repair_ssa();
3616 f
3617 }
3618
parse_shader(mut self) -> Shader<'a>3619 pub fn parse_shader(mut self) -> Shader<'a> {
3620 let mut functions = Vec::new();
3621 for nf in self.nir.iter_functions() {
3622 if let Some(nfi) = nf.get_impl() {
3623 let f = self.parse_function_impl(nfi);
3624 functions.push(f);
3625 }
3626 }
3627
3628 // Tessellation evaluation shaders MUST claim to read gl_TessCoord or
3629 // the hardware will throw an SPH error.
3630 if matches!(self.info.stage, ShaderStageInfo::Tessellation(_)) {
3631 match &mut self.info.io {
3632 ShaderIoInfo::Vtg(io) => {
3633 let tc = NAK_ATTR_TESS_COORD;
3634 io.mark_attrs_written(tc..(tc + 8));
3635 }
3636 _ => panic!("Tessellation must have ShaderIoInfo::Vtg"),
3637 }
3638 }
3639
3640 Shader {
3641 sm: self.sm,
3642 info: self.info,
3643 functions: functions,
3644 }
3645 }
3646 }
3647
nak_shader_from_nir<'a>( nak: &nak_compiler, ns: &'a nir_shader, sm: &'a dyn ShaderModel, ) -> Shader<'a>3648 pub fn nak_shader_from_nir<'a>(
3649 nak: &nak_compiler,
3650 ns: &'a nir_shader,
3651 sm: &'a dyn ShaderModel,
3652 ) -> Shader<'a> {
3653 ShaderFromNir::new(nak, ns, sm).parse_shader()
3654 }
3655