perf: use bfi.b32 for unaligned Q6K loads — 24 fewer instructions per super-block (Refs GH-131)

Replace shl+or byte assembly in ld_global_u32_unaligned with bfi.b32 bit-field
insert. Each unaligned u32 load saves 6 instructions (9 → 3 for the packing step).
With 4 unaligned loads per Q6K super-block, this reduces instruction overhead by 24
per super-block on sm_87 Jetson Orin.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

Author: noahgift

.pmat/baseline.json

@@ -1,6 +1,6 @@
 {
   "version": "3.6.1",
-  "created_at": "2026-03-05T18:28:04.491036699Z",
+  "created_at": "2026-03-05T22:38:52.308212548Z",
   "git_context": null,
   "files": {},
   "summary": {

trueno-gpu/src/kernels/quantize/q6k/dp4a.rs

@@ -376,6 +376,8 @@ mod tests {
         assert!(ptx.contains(".visible .entry dp4a_q6k_gemv"));
         assert!(ptx.contains("dp4a.u32.s32"), "Must use dp4a instructions");
         assert!(ptx.contains("bar.sync"), "Must have barrier for cross-warp safety");
+        // GH-131: bfi.b32 used for unaligned Q6K loads (replaces shl+or assembly)
+        assert!(ptx.contains("bfi.b32"), "Must use bfi.b32 for unaligned byte packing");
     }
 
     #[test]

trueno-gpu/src/ptx/builder/global_mem.rs

@@ -7,7 +7,7 @@ use crate::ptx::instructions::{Operand, PtxInstruction, PtxOp};
 use crate::ptx::registers::VirtualReg;
 use crate::ptx::types::{PtxStateSpace, PtxType};
 
-use super::{KernelBuilder, PtxArithmetic, PtxControl};
+use super::{KernelBuilder, PtxArithmetic};
 
 impl<'a> KernelBuilder<'a> {
     // ===== Memory Operations (vectorized - not in traits) =====
@@ -154,12 +154,16 @@ impl<'a> KernelBuilder<'a> {
 
     /// Load u32 from potentially unaligned global memory address.
     ///
-    /// Uses 4 byte loads + shifts to assemble a u32, avoiding
+    /// Uses 4 byte loads + `bfi.b32` to assemble a u32, avoiding
     /// `ld.global.u32` alignment requirements (4-byte aligned).
     /// Required for Q6K super-blocks (210 bytes each, not 4-byte aligned).
     ///
     /// sm_87 (Jetson Orin) faults on misaligned ld.global.u32 with
     /// CUDA_ERROR_MISALIGNED_ADDRESS (716).
+    ///
+    /// GH-131: Optimized from shl+or (9 instructions) to bfi.b32 (3 instructions)
+    /// for the byte assembly step. Saves 6 instructions per call × 4 calls per
+    /// Q6K super-block = 24 fewer instructions per super-block.
     pub fn ld_global_u32_unaligned(&mut self, addr: VirtualReg) -> VirtualReg {
         // Load 4 consecutive bytes
         let b0 = self.ld_global_u8(addr);
@@ -173,20 +177,17 @@ impl<'a> KernelBuilder<'a> {
         let addr3 = self.add_u64(addr, off3);
         let b3 = self.ld_global_u8(addr3);
 
-        // Convert u8 (in u16 registers) to u32 and assemble little-endian
+        // Convert u8 (in u16 registers) to u32
         let w0 = self.cvt_u32_u8(b0); // byte 0 → bits [7:0]
         let w1 = self.cvt_u32_u8(b1);
         let w2 = self.cvt_u32_u8(b2);
         let w3 = self.cvt_u32_u8(b3);
-        let eight = self.mov_u32_imm(8);
-        let sixteen = self.mov_u32_imm(16);
-        let twentyfour = self.mov_u32_imm(24);
-        let s1 = self.shl_u32(w1, eight);     // byte 1 → bits [15:8]
-        let s2 = self.shl_u32(w2, sixteen);    // byte 2 → bits [23:16]
-        let s3 = self.shl_u32(w3, twentyfour); // byte 3 → bits [31:24]
-        let t01 = self.or_u32(w0, s1);
-        let t23 = self.or_u32(s2, s3);
-        self.or_u32(t01, t23)
+
+        // Assemble little-endian u32 using bfi.b32 (3 instructions vs 9 with shl+or)
+        // bfi.b32 inserts `len` bits from `insert` into `base` at position `start`
+        let t1 = self.bfi_b32(w1, w0, 8, 8);     // insert byte 1 at bits [15:8]
+        let t2 = self.bfi_b32(w2, t1, 16, 8);     // insert byte 2 at bits [23:16]
+        self.bfi_b32(w3, t2, 24, 8)               // insert byte 3 at bits [31:24]
     }
 
     /// Load u16 from global memory (for f16 as raw bits)

trueno-gpu/src/ptx/builder/warp_vote.rs

@@ -167,6 +167,32 @@ impl<'a> KernelBuilder<'a> {
         dst
     }
 
+    /// Bit field insert: insert `len` bits from `insert` into `base` at position `start`
+    ///
+    /// PTX: `bfi.b32 dst, insert, base, start, len;`
+    /// dst = base with bits [start..start+len-1] replaced by insert[0..len-1]
+    ///
+    /// GH-131: Used to pack bytes into u32 for unaligned Q6K loads on sm_87.
+    /// Replaces 3 instructions (mov+shl+or) with 1 instruction per byte insertion.
+    pub fn bfi_b32(
+        &mut self,
+        insert: VirtualReg,
+        base: VirtualReg,
+        start: u32,
+        len: u32,
+    ) -> VirtualReg {
+        let dst = self.registers.allocate_virtual(PtxType::U32);
+        self.instructions.push(
+            PtxInstruction::new(PtxOp::Bfi, PtxType::B32)
+                .dst(Operand::Reg(dst))
+                .src(Operand::Reg(insert))
+                .src(Operand::Reg(base))
+                .src(Operand::ImmI64(start as i64))
+                .src(Operand::ImmI64(len as i64)),
+        );
+        dst
+    }
+
     /// Load f32 immediate constant
     ///
     /// PAR-062: Used for NEG_INFINITY initialization