NVIDIA / NVIDIA/cutlass

[BUG] Modifying the block/warptile shapes and the output datatype in the unit test causes the tests to fail.

Open
#2,041 3 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage 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
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

  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 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.