diff --git a/src/nouveau/compiler/nak/builder.rs b/src/nouveau/compiler/nak/builder.rs index 021e11a2d3c..160efd2924f 100644 --- a/src/nouveau/compiler/nak/builder.rs +++ b/src/nouveau/compiler/nak/builder.rs @@ -43,13 +43,13 @@ pub trait Builder { } if is_predicate { self.push_op(OpPLop3 { - dsts: [dst.into(), Dst::None], + dsts: [dst, Dst::None], srcs: [x, y, true.into()], ops: [op, LogicOp3::new_const(false)], }); } else { self.push_op(OpLop3 { - dst: dst.into(), + dst: dst, srcs: [x, y, 0.into()], op: op, }); @@ -68,13 +68,13 @@ pub trait Builder { } }; self.push_op(OpPSetP { - dsts: [dst.into(), Dst::None], + dsts: [dst, Dst::None], ops: [cmp_op, PredSetOp::And], srcs: [x, y, true.into()], }); } else { self.push_op(OpLop2 { - dst: dst.into(), + dst: dst, srcs: [x, y], op: op, }); @@ -566,12 +566,12 @@ pub trait SSABuilder: Builder { self.push_op(OpPSetP { dsts: [tmp.into(), Dst::None], ops: [PredSetOp::And, PredSetOp::And], - srcs: [cond.into(), x.into(), true.into()], + srcs: [cond, x, true.into()], }); self.push_op(OpPSetP { dsts: [dst.into(), Dst::None], ops: [PredSetOp::And, PredSetOp::Or], - srcs: [Src::from(cond).bnot(), y.into(), tmp.into()], + srcs: [cond.bnot(), y, tmp.into()], }); } dst diff --git a/src/nouveau/compiler/nak/encode_sm70.rs b/src/nouveau/compiler/nak/encode_sm70.rs index fc927a2fe4d..117c563900b 100644 --- a/src/nouveau/compiler/nak/encode_sm70.rs +++ b/src/nouveau/compiler/nak/encode_sm70.rs @@ -815,8 +815,8 @@ impl SM70Instr { self.encode_alu( 0x00c, None, - ALUSrc::from_src(&op.srcs[0].into()), - ALUSrc::from_src(&op.srcs[1].into()), + ALUSrc::from_src(&op.srcs[0]), + ALUSrc::from_src(&op.srcs[1]), ALUSrc::None, ); @@ -843,9 +843,9 @@ impl SM70Instr { self.encode_alu( 0x012, Some(op.dst), - ALUSrc::from_src(&op.srcs[0].into()), - ALUSrc::from_src(&op.srcs[1].into()), - ALUSrc::from_src(&op.srcs[2].into()), + ALUSrc::from_src(&op.srcs[0]), + ALUSrc::from_src(&op.srcs[1]), + ALUSrc::from_src(&op.srcs[2]), ); self.set_field(72..80, op.op.lut); @@ -898,7 +898,7 @@ impl SM70Instr { 0x104, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } else { @@ -906,7 +906,7 @@ impl SM70Instr { 0x110, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } @@ -927,7 +927,7 @@ impl SM70Instr { 0x105, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } else { @@ -935,7 +935,7 @@ impl SM70Instr { 0x111, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } @@ -954,7 +954,7 @@ impl SM70Instr { 0x106, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } else { @@ -962,7 +962,7 @@ impl SM70Instr { 0x112, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } @@ -980,7 +980,7 @@ impl SM70Instr { 0x107, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } else { @@ -988,7 +988,7 @@ impl SM70Instr { 0x113, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); } @@ -1004,7 +1004,7 @@ impl SM70Instr { 0x002, Some(op.dst), ALUSrc::None, - ALUSrc::from_src(&op.src.into()), + ALUSrc::from_src(&op.src), ALUSrc::None, ); self.set_field(72..76, op.quad_lanes); @@ -1037,8 +1037,8 @@ impl SM70Instr { self.encode_alu( 0x007, Some(op.dst), - ALUSrc::from_src(&op.srcs[0].into()), - ALUSrc::from_src(&op.srcs[1].into()), + ALUSrc::from_src(&op.srcs[0]), + ALUSrc::from_src(&op.srcs[1]), ALUSrc::None, ); diff --git a/src/nouveau/compiler/nak/from_nir.rs b/src/nouveau/compiler/nak/from_nir.rs index c1bb0f20260..af33238a9d2 100644 --- a/src/nouveau/compiler/nak/from_nir.rs +++ b/src/nouveau/compiler/nak/from_nir.rs @@ -28,9 +28,9 @@ fn init_info_from_nir(nir: &nir_shader, sm: u8) -> ShaderInfo { MESA_SHADER_COMPUTE => { ShaderStageInfo::Compute(ComputeShaderInfo { local_size: [ - nir.info.workgroup_size[0].into(), - nir.info.workgroup_size[1].into(), - nir.info.workgroup_size[2].into(), + nir.info.workgroup_size[0], + nir.info.workgroup_size[1], + nir.info.workgroup_size[2], ], smem_size: nir.info.shared_size.try_into().unwrap(), }) @@ -280,7 +280,7 @@ impl<'a> ShaderFromNir<'a> { } let bits = usize::from(def.bit_size) * usize::from(def.num_components); - assert!(vec.len() == bits.div_ceil(32).into()); + assert!(vec.len() == bits.div_ceil(32)); } self.ssa_map .entry(def.index) @@ -1632,7 +1632,7 @@ impl<'a> ShaderFromNir<'a> { if mask & (1 << i) == 0 { nir_dst.push(b.copy(0.into())[0]); } else { - nir_dst.push(dst[di].into()); + nir_dst.push(dst[di]); di += 1; } } @@ -1959,7 +1959,7 @@ impl<'a> ShaderFromNir<'a> { let coord = self.get_image_coord(intrin, dim); // let sample = self.get_src(&srcs[2]); - let comps = u8::try_from(intrin.num_components).unwrap(); + let comps = intrin.num_components; assert!(intrin.def.bit_size() == 32); assert!(comps == 1 || comps == 2 || comps == 4); @@ -1985,7 +1985,7 @@ impl<'a> ShaderFromNir<'a> { // let sample = self.get_src(&srcs[2]); let data = self.get_src(&srcs[3]); - let comps = u8::try_from(intrin.num_components).unwrap(); + let comps = intrin.num_components; assert!(srcs[3].bit_size() == 32); assert!(comps == 1 || comps == 2 || comps == 4); @@ -2555,7 +2555,7 @@ impl<'a> ShaderFromNir<'a> { assert!(addr % 4 == 0); for c in 0..usize::from(intrin.num_components) { - let idx = usize::from(addr / 4) + usize::from(c); + let idx = usize::from(addr / 4) + c; self.fs_out_regs[idx] = data.as_ssa().unwrap()[c]; } } diff --git a/src/nouveau/compiler/nak/spill_values.rs b/src/nouveau/compiler/nak/spill_values.rs index 0563b53943b..2fbc0090c1f 100644 --- a/src/nouveau/compiler/nak/spill_values.rs +++ b/src/nouveau/compiler/nak/spill_values.rs @@ -156,7 +156,7 @@ impl Spill for SpillBar { assert!(dst.file() == RegFile::GPR); Instr::new_boxed(OpBMov { dst: dst.into(), - src: src.into(), + src: src, clear: false, }) } @@ -164,7 +164,7 @@ impl Spill for SpillBar { fn fill(&self, dst: Dst, src: SSAValue) -> Box { assert!(src.file() == RegFile::GPR); Instr::new_boxed(OpBMov { - dst: dst.into(), + dst: dst, src: src.into(), clear: false, }) @@ -456,7 +456,7 @@ fn spill_values( some.push(Reverse(SSANextUse::new(*ssa, next_use))); } } - while w.count(file) < limit.into() { + while w.count(file) < limit { let Some(entry) = some.pop() else { break; }; @@ -465,7 +465,7 @@ fn spill_values( // If we still have room, consider values which aren't used // inside the loop. - if w.count(file) < limit.into() { + if w.count(file) < limit { for ssa in i_b.iter() { debug_assert!(ssa.file() == file); if !lu.contains(ssa) { @@ -474,7 +474,7 @@ fn spill_values( } } - while w.count(file) < limit.into() { + while w.count(file) < limit { let Some(entry) = some.pop() else { break; }; @@ -522,7 +522,7 @@ fn spill_values( some.push(Reverse(SSANextUse::new(ssa, info.next_use))); } } - while w.count(file) < limit.into() { + while w.count(file) < limit { let Some(entry) = some.pop() else { break; }; @@ -698,8 +698,8 @@ fn spill_values( let abs_pressure = b.w.count(file) + u32::from(rel_pressure); - if abs_pressure > limit.into() { - let count = abs_pressure - u32::from(limit); + if abs_pressure > limit { + let count = abs_pressure - limit; let count = count.try_into().unwrap(); let mut spills = SpillChooser::new(bl, ip, count);