@@ -114,6 +114,10 @@ Version | Description |
 | 13.2.78 | [m035-xor-not-predicate-fold](known-miscompiles/m035-xor-not-predicate-fold/NOTES.md): Likely same root cause as m033; `xor.b32` by `0xffffffff` feeding a predicate selects the wrong arm. |
 | 13.2.78 | [m034-bfind-zero-branch-fold](known-miscompiles/m034-bfind-zero-branch-fold/NOTES.md): Branch fold treats `bfind.u32 0` as `0` instead of `0xffffffff`. |
 | 13.2.78 | [m036-mulhi-control-fold](known-miscompiles/m036-mulhi-control-fold/NOTES.md): Control-flow fold around `mul.hi.s32` uses an incorrect folded constant. |
+| 13.2.78 | [m037-bmsk-clz-bfi-fold](known-miscompiles/m037-bmsk-clz-bfi-fold/NOTES.md): `bmsk` / `clz` / `bfi` / `mad.lo` value-chain fold sets an extra output bit. |
+| 13.2.78 | [m038-structured-empty-else-fold](known-miscompiles/m038-structured-empty-else-fold/NOTES.md): Always-false structured branch with an empty else arm folds as if the untaken then arm executed. |
+| 13.2.78 | [m039-else-redefinition-fold](known-miscompiles/m039-else-redefinition-fold/NOTES.md): Branch fold drops the executed else-path redefinition of a value initialized before the branch. |
+| 13.2.78 | [m040-mulwide-neg-shr-fold](known-miscompiles/m040-mulwide-neg-shr-fold/NOTES.md): `mul.wide` low word feeding wrapped negation and logical shift loses the shifted high-bit contribution. |
 
 ## Running
 
@@ -205,6 +209,7 @@ that feature.
 | `DIV_DISABLE_MULHI` | `mul.hi.u32` and `mul.hi.s32`. |
 | `DIV_DISABLE_SIGNED_MULHI` | `mul.hi.s32` only. |
 | `DIV_DISABLE_BITWISE_BINOPS` | `and.b32`, `or.b32`, `xor.b32`. |
+| `DIV_DISABLE_XOR` | `xor.b32` while retaining `and.b32` and `or.b32`. |
 | `DIV_DISABLE_PRMT` | `prmt.b32`. |
 | `DIV_DISABLE_NOT` | `not.b32` and xor-by-`0xffffffff` forms. |
 | `DIV_DISABLE_CLZ` | `clz.b32`. |
 
@@ -260,17 +260,30 @@ fn find_greedy_removal(
 }
 
 /// Indices into `ptx.lines()` that the reducer is allowed to try removing.
-/// Sacrosanct: prologue (everything up to and including `bra block_0;`),
-/// epilogue address arithmetic, labels, structural lines, `ret;`, braces.
+/// Sacrosanct: prologue (everything through the generated ABI/register
+/// initialization), epilogue address arithmetic, labels, structural lines,
+/// `ret;`, braces.
 fn removable_indices(ptx: &str) -> Result<Vec<usize>> {
     let lines: Vec<&str> = ptx.lines().collect();
-    let prologue_end = lines
+    let prologue_end = if let Some(i) = lines
         .iter()
         .position(|l| {
             let t = l.trim();
             t.starts_with("bra") && t.contains("block_0;")
-        })
-        .ok_or_else(|| anyhow!("could not locate `bra block_0;` to mark end of prologue"))?;
+        }) {
+        i
+    } else {
+        lines
+            .iter()
+            .enumerate()
+            .skip_while(|(_, l)| !l.trim().starts_with("ld.global.u32"))
+            .find_map(|(i, l)| (i > 0 && l.trim().is_empty()).then_some(i))
+            .ok_or_else(|| {
+                anyhow!(
+                    "could not locate `bra block_0;` or structured register-init end to mark end of prologue"
+                )
+            })?
+    };
     let exit_start = lines
         .iter()
         .position(|l| l.trim() == "exit:")
@@ -283,6 +296,11 @@ fn removable_indices(ptx: &str) -> Result<Vec<usize>> {
         if t.is_empty() || t.ends_with(':') {
             continue;
         }
+        // Removing a predicate definition can leave an undefined branch or
+        // selp predicate while still producing deterministic-looking output.
+        if t.starts_with("setp.") {
+            continue;
+        }
         out.push(i);
     }
     // Epilogue: store lines only; address arithmetic and `ret;` stay put.
 
@@ -32,6 +32,8 @@
 //!                         PTX mul.hi.s32 generation while retaining mul.hi.u32
 //!   DIV_DISABLE_BITWISE_BINOPS default: false; set 1/true/yes/on to suppress
 //!                         PTX and.b32/or.b32/xor.b32 generation
+//!   DIV_DISABLE_XOR       default: false; set 1/true/yes/on to suppress
+//!                         PTX xor.b32 generation while retaining and.b32/or.b32
 //!   DIV_DISABLE_PRMT      default: false; set 1/true/yes/on to suppress
 //!                         PTX prmt.b32 generation
 //!   DIV_DISABLE_NOT       default: false; set 1/true/yes/on to suppress
@@ -189,6 +191,8 @@ struct Args {
     #[arg(long)]
     disable_bitwise_binops: bool,
     #[arg(long)]
+    disable_xor: bool,
+    #[arg(long)]
     disable_prmt: bool,
     #[arg(long)]
     disable_not: bool,
@@ -289,6 +293,7 @@ impl Args {
         set_bool!(self.disable_mulhi, "DIV_DISABLE_MULHI");
         set_bool!(self.disable_signed_mulhi, "DIV_DISABLE_SIGNED_MULHI");
         set_bool!(self.disable_bitwise_binops, "DIV_DISABLE_BITWISE_BINOPS");
+        set_bool!(self.disable_xor, "DIV_DISABLE_XOR");
         set_bool!(self.disable_prmt, "DIV_DISABLE_PRMT");
         set_bool!(self.disable_not, "DIV_DISABLE_NOT");
         set_bool!(self.disable_clz, "DIV_DISABLE_CLZ");
@@ -389,6 +394,7 @@ impl Config {
         let disable_mulhi = env_bool("DIV_DISABLE_MULHI")?.unwrap_or(false);
         let disable_signed_mulhi = env_bool("DIV_DISABLE_SIGNED_MULHI")?.unwrap_or(false);
         let disable_bitwise_binops = env_bool("DIV_DISABLE_BITWISE_BINOPS")?.unwrap_or(false);
+        let disable_xor = env_bool("DIV_DISABLE_XOR")?.unwrap_or(false);
         let disable_prmt = env_bool("DIV_DISABLE_PRMT")?.unwrap_or(false);
         let disable_not = env_bool("DIV_DISABLE_NOT")?.unwrap_or(false);
         let disable_clz = env_bool("DIV_DISABLE_CLZ")?.unwrap_or(false);
@@ -429,6 +435,7 @@ impl Config {
             emit_mulhi: !disable_mulhi,
             emit_signed_mulhi: !disable_signed_mulhi,
             emit_bitwise_binops: !disable_bitwise_binops,
+            emit_xor: !disable_xor,
             emit_prmt: !disable_prmt,
             emit_not: !disable_not,
             emit_clz: !disable_clz,
@@ -639,7 +646,7 @@ fn main() -> Result<()> {
 
     let total_workers = cfg.gpus.len() * cfg.workers_per_gpu;
     eprintln!(
-        "fuzzx-diff: starting_seed=0x{:016x} out={} program_bytes={} max_iters={} control_flow={:?} blocks={}..{} insts_per_block={}..{} regs={} max_loop_iters={} max_immediate={} max_structured_depth={} emit_structured_loops={} emit_arbitrary_loops={} emit_lop3={} emit_minmax={} emit_sub={} emit_mulhi={} emit_signed_mulhi={} emit_bitwise_binops={} emit_prmt={} emit_not={} emit_clz={} emit_cnot={} emit_abs={} emit_signed_cmp={} emit_signed_divrem={} emit_funnel={} emit_neg={} emit_shl={} emit_signed_shr={} emit_bfind={} emit_bfi={} emit_bmsk={} emit_mad24={} emit_mul24={} emit_mul_wide={} emit_wide_int={} emit_addc={} emit_subc={} emit_i32_boundary_immediates={} emit_dp2a={} emit_set={} emit_s32_slct={} emit_video={} emit_vsub4={} gpus={:?} workers_per_gpu={} (total={})",
+        "fuzzx-diff: starting_seed=0x{:016x} out={} program_bytes={} max_iters={} control_flow={:?} blocks={}..{} insts_per_block={}..{} regs={} max_loop_iters={} max_immediate={} max_structured_depth={} emit_structured_loops={} emit_arbitrary_loops={} emit_lop3={} emit_minmax={} emit_sub={} emit_mulhi={} emit_signed_mulhi={} emit_bitwise_binops={} emit_xor={} emit_prmt={} emit_not={} emit_clz={} emit_cnot={} emit_abs={} emit_signed_cmp={} emit_signed_divrem={} emit_funnel={} emit_neg={} emit_shl={} emit_signed_shr={} emit_bfind={} emit_bfi={} emit_bmsk={} emit_mad24={} emit_mul24={} emit_mul_wide={} emit_wide_int={} emit_addc={} emit_subc={} emit_i32_boundary_immediates={} emit_dp2a={} emit_set={} emit_s32_slct={} emit_video={} emit_vsub4={} gpus={:?} workers_per_gpu={} (total={})",
         cfg.starting_seed,
         cfg.out_dir.display(),
         cfg.program_bytes,
@@ -663,6 +670,7 @@ fn main() -> Result<()> {
         cfg.gen_config.emit_mulhi,
         cfg.gen_config.emit_signed_mulhi,
         cfg.gen_config.emit_bitwise_binops,
+        cfg.gen_config.emit_xor,
         cfg.gen_config.emit_prmt,
         cfg.gen_config.emit_not,
         cfg.gen_config.emit_clz,
 
@@ -59,6 +59,7 @@ pub struct GenConfig {
     pub emit_mulhi: bool,
     pub emit_signed_mulhi: bool,
     pub emit_bitwise_binops: bool,
+    pub emit_xor: bool,
     pub emit_prmt: bool,
     pub emit_not: bool,
     pub emit_clz: bool,
@@ -116,6 +117,7 @@ impl Default for GenConfig {
             emit_mulhi: true,
             emit_signed_mulhi: true,
             emit_bitwise_binops: true,
+            emit_xor: true,
             emit_prmt: true,
             emit_not: true,
             emit_clz: true,
@@ -1027,6 +1029,7 @@ impl<'a> Generator<'a> {
                 self.cfg.emit_mulhi,
                 self.cfg.emit_signed_mulhi,
                 self.cfg.emit_bitwise_binops,
+                self.cfg.emit_xor,
             )?;
             Ok(Inst::Bin {
                 op,
@@ -1818,13 +1821,17 @@ fn pick_binop(
     emit_mulhi: bool,
     emit_signed_mulhi: bool,
     emit_bitwise_binops: bool,
+    emit_xor: bool,
 ) -> Result<BinOp> {
     let mut ops = vec![BinOp::Add, BinOp::Mul];
     if emit_sub {
         ops.push(BinOp::Sub);
     }
     if emit_bitwise_binops {
-        ops.extend([BinOp::And, BinOp::Or, BinOp::Xor]);
+        ops.extend([BinOp::And, BinOp::Or]);
+        if emit_xor {
+            ops.push(BinOp::Xor);
+        }
     }
     if emit_mulhi {
         ops.push(BinOp::MulHi);
@@ -2385,6 +2392,29 @@ mod tests {
         }
     }
 
+    #[test]
+    fn xor_generation_can_be_disabled() {
+        let cfg = GenConfig {
+            emit_xor: false,
+            ..GenConfig::default()
+        };
+
+        let mut saw_and = false;
+        let mut saw_or = false;
+        for seed in 0..512 {
+            let bytes = bytes_from_seed(seed, 4096);
+            let ptx = generate_from_bytes_with_config(&bytes, &cfg).unwrap();
+            assert!(
+                !has_mnemonic(&ptx, "xor.b32"),
+                "seed {seed:x} emitted xor.b32"
+            );
+            saw_and |= has_mnemonic(&ptx, "and.b32");
+            saw_or |= has_mnemonic(&ptx, "or.b32");
+        }
+        assert!(saw_and, "sample did not retain and.b32 coverage");
+        assert!(saw_or, "sample did not retain or.b32 coverage");
+    }
+
     #[test]
     fn prmt_generation_can_be_disabled() {
         let cfg = GenConfig {
 
@@ -51,45 +51,40 @@ in_n     = 32
 All lanes compute the same value and store to `out[0][1]`, so the repeated
 stores are benign. The standalone reproducer launches one thread.
 
+On 2026-05-15, this was further reduced from 106 to 49 lines by deleting
+semantically dead arithmetic, folding constants on the taken path, and removing
+untaken-path scaffolding while preserving the same `0x80000092` (`-O0`) versus
+`0x80000090` (`-O2`/`-O3`) result. Several apparently dead control-flow edges
+still need to remain; deleting them causes optimized ptxas to stop taking the
+buggy fold.
+
 ## Correct Scalar Trace
 
 The taken path is:
 
 ```text
-structured_if_0_then -> structured_if_1_else -> structured_if_2_else
--> structured_if_3_then
+structured_if_1_else -> structured_if_2_done -> structured_if_3_done
 ```
 
 For `input[0] = 0x55ff25dc` and `in_n = 32`:
 
 ```text
-%r11 = 60
-%r18 = mad.lo.u32 60, 49375, 60 = 0x002d3480
-%p1  = true
-
-%r19 = 8
-%r5  = 8 - 0x002d3480 = 0xffd2cb88
+%r5  = 0xffd2cb88
 %r6  = 32 * 0xffd2cb88 + 0xffd2cb88 = 0xfa2c3c88
-%r8  = 16 >> 11 = 0
 
 %p6  = setp.ge.u32 19682, 0x55ff25dc = false
-%r3  = 0 ^ 0xffd2cb88 = 0xffd2cb88
-%r19 = clz.b32 0 = 32
-%r10 = 0x002d3480 & 2 = 0
-%r15 = ((0xffd2cb88 * 0xfa2c3c88 + 31152) >> 26) = 6
+%r1  = mad.lo.u32 0xffd2cb88, 0xfa2c3c88, 31152 = 0x1b1079f0
+%r15 = 0x1b1079f0 >> 26 = 6
 %p14 = false
 %r0  = 0x40000000
-%r8  = brev.b32 0x20000000 = 4
+%r8  = 4
 %r16 = 0x40000000 ^ 33145 = 0x40008179
-%r2  = popc.b32 24696 = 6
 %r4  = mul.hi.s32 6, 0x40008179 = 1
-%r12 = 0xffd2cb88 * 0 + 1 = 1
 %p18 = setp.eq.u32 1, 0 = false
 
-%r13 = mul.hi.s32 489, 0xffd2cb88 = 0xffffffff
 %r14 = mad.lo.u32 4, 0x20000000, 0xffffffff = 0x7fffffff
 %r19 = 1 - 0x7fffffff = 0x80000002
-%r1  = mad.lo.u32 12, 12, 0x80000002 = 0x80000092
+%r1  = 0x80000002 + 144 = 0x80000092
 %p21 = setp.le.u32 4, 0x80000002 = true
 store %r1 to out[0][1]
 ```
 
@@ -8,99 +8,42 @@
     .param .u32 in_n
 )
 {
-    .reg .pred  %p<26>;
-    .reg .b32   %r<22>;
-    .reg .b64   %rd<8>;
+    .reg .pred  %p<22>;
+    .reg .b32   %r<20>;
+    .reg .b64   %rd<5>;
 
     ld.param.u64    %rd0, [in_ptr];
     ld.param.u64    %rd1, [out_ptr];
     ld.param.u32    %r0, [in_n];
     cvta.to.global.u64 %rd2, %rd0;
     ld.global.u32   %r2, [%rd2];
-    mov.u32         %r17, 0;
-    mov.u32         %r3, 0;
-    mov.u32         %r7, 0;
-    setp.eq.u32     %p24, 0, 1;
 
-    mul.hi.s32      %r14, %r0, 16;
-    mov.u32         %r11, 60;
-    mad.lo.u32      %r18, %r11, 49375, %r11;
-    shr.u32         %r13, 69160512, 25;
-    mad.lo.u32      %r5, 4, %r14, %r0;
-    setp.eq.u32     %p1, 0, 0;
-    @%p1 bra        structured_if_0_then;
-    bra             structured_if_0_else;
-structured_if_0_then:
-    setp.eq.u32     %p3, 0, 0;
-    selp.b32        %r19, 8, %r17, %p3;
-    sub.u32         %r5, %r19, %r18;
+    mov.u32         %r5, 0xffd2cb88;
     mad.lo.u32      %r6, %r0, %r5, %r5;
-    shr.u32         %r8, 16, 11;
-    bra             structured_if_0_done;
-structured_if_0_else:
-    bra             structured_if_0_done;
-structured_if_0_done:
     setp.ge.u32     %p6, 19682, %r2;
-    @%p6 bra        structured_if_1_then;
-    bra             structured_if_1_else;
-structured_if_1_then:
-    mad.lo.u32      %r19, %r19, 4194304, 27678;
-    clz.b32         %r8, %r17;
-    setp.ge.u32     %p9, %r19, 53078;
-    selp.b32        %r17, 4294967295, %r7, %p9;
-    dp4a.s32.u32    %r10, %r3, %r7, 62161;
+    @!%p6 bra       structured_if_1_else;
+    clz.b32         %r8, 0;
     bra             structured_if_1_done;
 structured_if_1_else:
-    xor.b32         %r3, %r8, %r5;
-    mad.lo.u32      %r2, %r18, %r19, 49358;
-    mul.hi.s32      %r1, %r13, 33098;
-    clz.b32         %r19, %r8;
-    cvt.u32.u16     %r2, %r2;
-    and.b32         %r10, %r18, %r13;
-    mad.lo.u32      %r1, %r3, %r6, 31152;
-    mad.lo.u32      %r8, %r19, 19382, %r5;
-    shr.u32         %r15, %r19, 11;
+    mad.lo.u32      %r1, %r5, %r6, 31152;
     shr.u32         %r15, %r1, 26;
-    mad.lo.u32      %r13, %r6, %r5, 1469367838;
-    sub.u32         %r6, 12, %r10;
     setp.ge.u32     %p14, %r15, 8;
     selp.b32        %r0, %r1, 1073741824, %p14;
-    shr.u32         %r7, 27123, 21;
-    brev.b32        %r8, 536870912;
+    mov.u32         %r8, 4;
     xor.b32         %r16, %r0, 33145;
-    bfe.u32         %r17, %r11, 25, 1;
-    popc.b32        %r2, 24696;
-    and.b32         %r1, %r0, %r19;
-    mul.hi.s32      %r4, %r2, %r16;
-    mad.lo.u32      %r12, %r5, %r10, %r4;
-    setp.eq.u32     %p18, %r12, %r7;
-    @%p18 bra       structured_if_2_then;
-    bra             structured_if_2_else;
-structured_if_2_then:
-    bra             structured_if_2_done;
-structured_if_2_else:
-    mul.hi.s32      %r13, 489, %r3;
-    mad.lo.u32      %r14, %r8, 536870912, %r13;
+    mul.hi.s32      %r4, 6, %r16;
+    setp.eq.u32     %p18, %r4, 0;
+    @%p18 bra       structured_if_2_done;
+    mad.lo.u32      %r14, %r8, 536870912, 0xffffffff;
     sub.u32         %r19, %r4, %r14;
-    mad.lo.u32      %r1, %r6, %r6, %r19;
-    and.b32         %r4, 24045, %r0;
-    bra             structured_if_2_done;
+    add.u32         %r1, %r19, 144;
 structured_if_2_done:
-    bra             structured_if_1_done;
 structured_if_1_done:
     setp.le.u32     %p21, %r8, %r19;
-    @%p21 bra       structured_if_3_then;
-    bra             structured_if_3_else;
-structured_if_3_then:
-    bra             structured_if_3_done;
-structured_if_3_else:
-    selp.b32        %r1, 5032, %r17, %p24;
-    bra             structured_if_3_done;
+    @%p21 bra       structured_if_3_done;
+    mov.u32         %r1, 0;
 structured_if_3_done:
-    bra             exit;
 
-exit:
-    cvta.to.global.u64 %rd4, %rd1;
-    st.global.u32   [%rd4 + 4], %r1;
+    st.global.u32   [%rd1 + 4], %r1;
     ret;
 }