NVIDIA / NVIDIA/cutlass

[BUG] Accuracy Error in CUTLASS GEMM operations.

Open
#2,334 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

Describe the bug
The result of accuracy of GEMM operation in CUTLASS (TensorOp, Simt) does not fully match accuracy of cuBLAS GEMM result.

Steps/Code to reproduce bug

using GEMM = cutlass::gemm::device::Gemm<
                            /*ElementA_ = */ float,
                            /*LayoutA_ = */ cutlass::layout::RowMajor,
                            /*ElementB_ = */ float,
                            /*LayoutB_ = */ cutlass::layout::ColumnMajor,
                            /*ElementC_ = */ float,
                            /*LayoutC_ = */ cutlass::layout::RowMajor,
                            /*ElementAccumulator_ = */ float,
                            /*OperatorClass_ = */ cutlass::arch::OpClassSimt,
                            /*ArchTag_ = */ cutlass::arch::Sm50,
                            /*ThreadblockShape_ = */ cutlass::gemm::GemmShape<128, 128, 8>,
                            /*WarpShape_ = */ cutlass::gemm::GemmShape<32, 64, 8>,
                            /*InstructionShape_ = */ cutlass::gemm::GemmShape<1, 1, 1>,
                            /*EpilogueOutputOp_ = */ cutlass::epilogue::thread::LinearCombination<float, 1, float, float, cutlass::epilogue::thread::ScaleType::Nothing, cutlass::FloatRoundStyle::round_to_nearest>,
                            /*ThreadblockSwizzle_ = */ cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<4>,
                            /*Stages = */ 2,
                            /*AlignmentA = */ 1,
                            /*AlignmentB = */ 1,
                            /*SplitKSerial = */ false,
                            /* operator = */ cutlass::arch::OpMultiplyAdd
                            >;
    GEMM gemm_operator;
    GEMM::Arguments args({m, n, k},
                        {reinterpret_cast<float *>(A.data_ptr()), k},
                        {reinterpret_cast<float *>(B.data_ptr()), k},
                        {reinterpret_cast<float *>(C.data_ptr()), n},
                        {reinterpret_cast<float *>(C.data_ptr()), n},
                        {float(1.0f), float(0.0f)});
    gemm_operator(args);

Expected behavior

I compare cuBLAS, and CUTLASS GEMM result using both bfloat16 and float32 precision to check CUTLASS get accurate result.
The code shown above is a portion of CUTLASS implementation for float32 data type.
I integrated the implementation into PyTorch via Pybind11 and verified its correctness by comparing the outputs with nn.Linear using absolute tolerance (atol) of 1e-6 and relative tolerance (rtol) of 1e-5.
Typically, the absolute difference is around 1e-3.
How should I improve this?

Environment details (please complete the following information):

  • CUDA 12.4
  • HW: RTX4070TI

Thank you.

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

Start with the provided cutlass::gemm::device::Gemm configuration and reproduce its float32 result against cuBLAS on the stated CUDA 12.4 and RTX4070TI environment. Compare the numerical behavior across the shown configuration and inputs; the issue is complete when the accuracy discrepancy is explained and a concrete correction or expected tolerance is established.

Written by the indexing model from the issue text.

Assessment

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.