How register liveness reducing is achieved

[wzhcz8902]

cutlass/include/cutlass/gemm/threadblock/index_remat.h

Lines 65 to 69 in 5c447dd

	/// Helper to rematerialize block Idx. Reduces register liveness.
	CUTLASS_DEVICE
	int RematerializeBlockIdxX() {
	return blockIdx.x;
	}

Can you help me understand why it helps reduce register liveness as the comment says?

[hwu36]

instead of

mov R1, blockIdx.x
...
1000 lines of assembly
...
setp R1

we prefer to do

setp blockIdx.x

just recompute simple things every time we use it so that we can have more free registers.

[mrakgr]

Drat, I was hoping the Cuda compiler did that optimization on its own. I can think of many places in my ML library where I failed to apply it.

[mrakgr]

I am testing this optimization out on some simpler cases, and as far as I can tell, the Cuda compiler is very good at optimizing array indexing. I cannot see any difference in register liveness or the instruction count. The last operand of the IMAD instruction, which is a constant confusingly keeps changing between different versions, but I cannot see any benefit to applying this optimization by hand.