[BUG] Modifying the block/warptile shapes and the output datatype in the unit test causes the tests to fail.
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
I modified the block/warptile shapes and the output datatype in https://github.com/NVIDIA/cutlass/blob/main/test/unit/gemm/device/gemm_s8t_s8n_s32t_tensor_op_s32_sm80.cu, and found some shapes cause the tests to fail.
I modified the ElementOutput to cutlass::half_t and tested various block/warptile shapes. While some shapes passed, others failed.
Passed block/warptile shapes, for example, are
<32,128,64>/<16,64,64>,
<32,128,64>/<32,32,64>,
<32,128,64>/<32,64,64>,
<32,256,64>/<32,64,64>,
<64,32,64>/<32,32,64>,
<64,32,64>/<64,32,64>, etc.
Failed block/warptile shapes, I found, are
<16,16,64>/<16,16,64>,
<16,128,64>/<16,128,64>,
<16,256,64>/<16,128,64>,
<32,16,64>/<32,16,64>,
<32,128,64>/<16,128,64>,
<32,128,64>/<32,128,64>,
<32,256,64>/<16,128,64>,
<32,256,64>/<32,128,64>,
<64,16,64>/<64,16,64>,
<64,128,64>/<16,128,64>,
<64,128,64>/<32,128,64>,
<64,128,64>/<64,128,64>,
<64,256,64>/<16,128,64>,
<64,256,64>/<32,128,64>,
<64,256,64>/<64,128,64>,
<128,16,64>/<128,16,64>,
<128,32,64>/<128,32,64>,
<128,64,64>/<128,64,64>,
<128,128,64>/<16,128,64>,
<128,128,64>/<32,128,64>,
<128,128,64>/<64,128,64>,
<128,256,64>/<128,128,64>,
<256,32,64>/<128,32,64>,
<256,64,64>/<128,64,64>,
<256,128,64>/<32,128,64>,
<256,128,64>/<64,128,64>,
<256,128,64>/<128,128,64>.
Steps/Code to reproduce bug
CUTLASS_TEST_L1(SM80_Device_Gemm_s8t_s8n_s32t_tensor_op_s32, 64x128x64_32x128x64, {
using ElementOutput = cutlass::half_t;
using ElementAccumulator = int32_t;
using ElementCompute = int32_t;
using Gemm = cutlass::gemm::device::Gemm<
int8_t, cutlass::layout::RowMajor, int8_t,
cutlass::layout::ColumnMajor, ElementOutput, cutlass::layout::RowMajor,
ElementAccumulator, cutlass::arch::OpClassTensorOp, cutlass::arch::Sm80,
cutlass::gemm::GemmShape<64, 128, 64>,
cutlass::gemm::GemmShape<32, 128, 64>, cutlass::gemm::GemmShape<16, 8, 32>,
cutlass::epilogue::thread::LinearCombination<
ElementOutput, 128 / cutlass::sizeof_bits<ElementOutput>::value,
ElementAccumulator, ElementCompute>,
cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>, 3>;
EXPECT_TRUE(test::gemm::device::TestAllGemm<Gemm>());
} )
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 test/unit/gemm/device/gemm_s8t_s8n_s32t_tensor_op_s32_sm80.cu and the shown SM80_Device_Gemm_s8t_s8n_s32t_tensor_op_s32 test case. Reproduce the listed passing and failing block/warptile shapes after changing ElementOutput to cutlass::half_t, then compare their results and failures. Done means the reported failing configurations pass without regressing the listed passing cases.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance, testing
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 25/100