[BUG] EVT-based GEMM in example 47 produces wrong results when m!=n in the problem size

[alexsamardzic]

Describe the bug

Results produced by EVT-based DeviceGemmStreamK GEMM in example 47_ampere_gemm_universal_streamk/ampere_gemm_universal_streamk_broadcast.cu seem to be wrong for the case when m and n in the problem size differ.

Steps/Code to reproduce bug

Run example say as follows:

./47_ampere_gemm_universal_streamk_broadcast --m=128 --n=256 --k=64

Large relative error will be reported.

To further verify that results calculated are wrong, I've set all inputs to 1 by adding following block after inputs initialization:

  cutlass::reference::host::TensorFill(options.tensor_a.host_view(), ElementA(1));
  cutlass::reference::host::TensorFill(options.tensor_b.host_view(), ElementB(1));
  cutlass::reference::host::TensorFill(options.tensor_c1.host_view(), ElementC(1));
  cutlass::reference::host::TensorFill(options.tensor_c2.host_view(), ElementC(1));
  cutlass::reference::host::TensorFill(options.tensor_Vector.host_view(), ElementC(1));

and then printed reference results and results produced by DeviceGemmStreamK GEMM, just before the end of the main() function:

  std::cout << options.tensor_ref_d.host_view() << "\n\n\n";
  std::cout << options.tensor_d.host_view() << "\n\n\n";

All values in reference results are 67, as expected, while right half of result matrix produced by DeviceGemmStreamK GEMM is all zero.

Expected behavior

Zero or near-zero relative error is expected to be produced for given problem size.

Environment details (please complete the following information):

Tested with latest CUTLASS main, on a machine with A100.

Additional context

Note that the problem is actually quite probably with StreamK, as all the other StreamK GEMMs in the example will also report the same relative error - but I found the issue while actually trying to mimic this example in order to use EVT in my code, so I was focused on DeviceGemmStreamK in this example.

[hwu36]

have you tried #1120?

[alexsamardzic]

Thanks, that change fixes it.