[BUG] Funcionality TensorOp 80+ s8 * s8 + s32 => {s32, s8} not working
Nobody has claimed this yet.
- 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
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- 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