[QST] How to efficiently convert fp32 Tensor to fp16 Tensor in Cutlass 3.x

[tridao]

I'd like to convert a fp32 Tensor (in registers) to a fp16 Tensor (in registers), ideally using the __float22half2_rn function for efficiency.
Cutlass 2.x has NumericArrayConverter that specializes to fp32 -> fp16 conversion that uses this function.

For Cutlass 3.x, I'm currently doing:

// accum is a fp32 Tensor in register
Tensor acc_fp16 = make_tensor<cutlass::half_t>(shape(accum));
for (int i = 0; i < size(accum); ++i) { acc_fp16(i) = accum(i); }

However this might not be as efficient, as it doesn't use the function for float2 -> half2 conversion. I looked at the generated PTX and it's not using anything like cvt.rn.f16x2.f32.

Should I cast the fp32 Tensor to Array, then use NumericArrayConverter, then cast the result back to Tensor?

[ccecka]

Though it's not 100% safe, depending on the Layout of the Tensor, this can work:

Tensor acc_fp32 = ...;  // FP32 Tensor
Tensor acc_fp16 = make_tensor<cutlass::half_t>(shape(acc_fp32));
{
  Tensor acc_fp32x2 = recast< float2>(acc_fp32);
  Tensor acc_fp16x2 = recast<__half2>(acc_fp16);
  for (int i = 0; i < size(acc_fp32x2); ++i) { acc_fp16x2(i) = __float22half2_rn(acc_fp32x2(i)); }
}
// Continue using acc_fp16

Note that using make_tensor<cutlass::half_t>(layout(acc_fp32)); rather than shape(acc_fp32) would be slightly safer.

[tridao]

This works, thank you!