cutlass icon indicating copy to clipboard operation
cutlass copied to clipboard

[BUG] Funcionality TensorOp 80+ s8 * s8 + s32 => {s32, s8} not working

Open IzanCatalan opened this issue 2 months ago • 6 comments

Describe the bug A clear and concise description of what the bug is. Hi, I have checked that fprop conv2d is not working with integers for an Nvidia A100. I have read in https://github.com/NVIDIA/cutlass/blob/main/media/docs/functionality.md that with Ampere gpus (sm80) is possible to perform convolution with integers as is also test in https://github.com/NVIDIA/cutlass/blob/main/test/unit/conv/device/conv2d_fprop_implicit_gemm_s8nhwc_s8nhwc_s32nhwc_tensor_op_s32_sm80.cu

However, when I modify https://github.com/NVIDIA/cutlass/blob/main/examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu example putting there the same config does not work, and all I get is the following output:

Got cutlass error: Error Invalid Problem at: 656 This is a reference to the lines:

  ImplicitGemm implicit_gemm_op;

  size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);

  // Allocate workspace memory
  cutlass::device_memory::allocation<int8_t> workspace(workspace_size);

  result.status = implicit_gemm_op.can_implement(arguments);
  CUTLASS_CHECK(result.status);

I debugged a little bit inside the ImplicitGemm class and the error comes from lines 107-109:

   Status status = UnderlyingKernel::Mma::IteratorA::can_implement(args.problem_size);
    if (Status::kSuccess != status) {
      return status;
    }

Is this behaviour normal? Why, in theory, according to the functionality readme (and test), can I perform int8t convolution, but it seems not to be working? What data types are available for fprop conv2d for Amperes (sm80) and Volta architectures (sm70)?

Can be the same problem with data types found if, instead of configuring a convolution with a C++ example like example16, I use cutlass with Python?

And related to the last question, for Python, is the same ImplicitGemm class use it or is there any other class called when is perform a conv2d?

Thanks.

Izan.

IzanCatalan avatar Dec 11 '24 16:12 IzanCatalan