engine-cuda: record why gf_sqr uses the general product

The three-multiply squaring is a real algebraic saving and a measured
loss. Implemented, verified bit-exact against __int128 over 40M random
values plus limb and power-of-two boundaries, then benchmarked on beast:
1072.9 vs 1137.7 MH/s over three interleaved rounds, 5.7% slower.

It trades 6.7% of the kernel's widening multiplies for 11% more
instructions -- 5 IMAD.WIDE in 56 instructions against the general
product's 6 in 48 -- and four formulations (condition-code carry, compare
carry, two mad.wide.u32 variants) all compiled to the same 5/56, so that
is the floor rather than a tuning failure.

Taken with the LAIR_INT_UNROLL result, the picture is consistent: this
kernel is bound by instruction issue and by the cards' power envelope,
not by multiply throughput or occupancy. Re-expressing the same work
cannot win; only removing work can. Left as a comment so the next person
does not spend the afternoon rediscovering it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_017ytCkttecH9SNTUEw8VT9D

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));
 }