Fix double-scheduling of dc in HSS hopper matmuls

[jacobhinkle]

When we do not have an epilogue (not even a cast), it might be the case that the original MmaOp has output which is a Fusion output. In this case the cached output which we often call dc is actually an mma_result. Currently this causes us to schedule that tensor once in scheduleMmaResults then again in scheduleEpilogue, leading to an esoteric error (see included test). This PR simply skips scheduling those tensors directly if they are already known to be mma results.

[jacobhinkle]

!test

[jacobhinkle]

We would have hit this eventually but the HSS tests are still guarded against Hopper. I'm posting this now to unblock some internal heuristics work.

[jacobhinkle]

@protonu we need to handle all possible dtypes in the epilogue.

[protonu]

Are we still using stmatrix if the output is fp32 and there wasn't a cast?

[jacobhinkle]

Yes, if you enable use_smem_epilogue in the included test we hit an error in scheduleStMatrixForMmaOutput.

[jacobhinkle]

So we need an if/else here that checks the dtype of d_smem and schedules with vectorized stores instead if not 16bit

Fuser/csrc/scheduler/hopper_multi_matmul.cpp

Lines 584 to 588 in 230f633

	MmaInputSmemSwizzle swizzle = mma_utils::tmaSwizzleSharedMemory(d_smem);
	// Schedule shared memory cache; Output from StMatrix
	mma_utils::scheduleStMatrixForMmaOutput(
	d_smem, swizzle, stmatrix_tile_m, stmatrix_tile_n);