Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
xiaonans opened this issue Jan 16, 2025 · 0 comments
Labels
? - Needs Triage bug Something isn't working

Comments

@xiaonans
Copy link

xiaonans commented Jan 16, 2025

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>());
} )
@xiaonans xiaonans added ? - Needs Triage bug Something isn't working labels Jan 16, 2025
@xiaonans xiaonans changed the title [BUG] Modifying the block or warptile shapes and the output datatype in the unit test causes the tests to fail. [BUG] Modifying the block/warptile shapes and the output datatype in the unit test causes the tests to fail. Jan 16, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
? - Needs Triage bug Something isn't working
Projects
None yet
Development

No branches or pull requests

1 participant