diff --git a/crates/engine-cuda/src/kernels/mining.cu b/crates/engine-cuda/src/kernels/mining.cu
index c1188d0..b55cbbf 100644
--- a/crates/engine-cuda/src/kernels/mining.cu
+++ b/crates/engine-cuda/src/kernels/mining.cu
@@ -90,6 +90,20 @@ __device__ __forceinline__ u64 gf_mul(u64 a, u64 b) {
     return gf_reduce(a * b, __umul64hi(a, b));
 }
 
+// Squaring deliberately uses the general product. nvcc does not specialise
+// `a * a`: it emits the same six IMAD.WIDE in 48 instructions as gf_mul
+// (measured, sm_120). The textbook saving is real but does not pay here --
+// a = a1*2^32 + a0 gives a^2 = a0^2 + a0*a1*2^33 + a1^2*2^64, three 32x32
+// multiplies instead of four, and it was implemented, verified bit-exact
+// against __int128 over 40M values, and measured 5.7% SLOWER on beast's 5090s
+// (1072.9 vs 1137.7 MH/s, three interleaved rounds). It saves 6.7% of the
+// kernel's widening multiplies and costs 11% more instructions: 5 multiplies
+// in 56 instructions against 6 in 48. Four formulations (condition-code
+// carry, compare carry, and two mad.wide.u32 variants) all compiled to the
+// same 5/56, so 56 is the floor. The scarce resource in this kernel is
+// instruction issue, not multiply throughput -- which is also why raising
+// occupancy lost (see LAIR_INT_UNROLL). Anything that wins here has to remove
+// work, not re-express it.
 __device__ __forceinline__ u64 gf_sqr(u64 a) {
     return gf_reduce(a * a, __umul64hi(a, a));
 }