NVIDIA / NVIDIA/cutlass

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

Open
#1,981 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug inactive-90d
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

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.

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Reproduce the failure in examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu using the configuration from test/unit/conv/device/conv2d_fprop_implicit_gemm_s8nhwc_s8nhwc_s32nhwc_tensor_op_s32_sm80.cu. Read include/cutlass/conv/device/implicit_gemm_convolution.h around can_implement and compare it with media/docs/functionality.md. Done means the supported SM80 and SM70 data types and the relevant C++ and Python convolution paths are identified and the reported mismatch is explained or corrected.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
machine-learning
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.