[BUG] Accuracy Error in CUTLASS GEMM operations.
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
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
- 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
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