Shared memory bank conflict problem

[FindDefinition]

I'm debugging a turing gemm kernel:
TileShape: 128x256x32
WarpTileShape: 64x64x32
layout: ttt
dtypes: all cutlass::half_t
tensor op: 1688
CUDA: 11.3

When I profile this kernel in nsight compute, lots of bank conflicts detected:

After debug, two issues found:

1. the 128bit access stmts are compiled to st.shared.u32 instead of st.shared.v4.u32 in ptx
2. the padding of output smem is wrong, the padding should be [0, 8] in this kernel instead of [0, 16].

Here is the problem:
How can I force the compiler generate st.shared.v4.u32? use inline asm? does the st.shared.v4.u32 require registers address to be aligned?

gemm code:

  using ElementAccumulator = cutlass::half_t;        
  using ElementComputeEpilogue = cutlass::half_t; 
  using ElementInputA = cutlass::half_t;
  using ElementInputB = cutlass::half_t;
  using ElementOutput = cutlass::half_t;
  using LayoutInputA = cutlass::layout::RowMajor;
  using LayoutInputB = cutlass::layout::RowMajor;
  using LayoutOutput = cutlass::layout::RowMajor;
  using MMAOp = cutlass::arch::OpClassTensorOp;
  // This code section describes CUDA SM architecture number
  using SmArch = cutlass::arch::Sm75;
  // This code section describes the tile size a thread block will compute
  using ShapeMMAThreadBlock =
      cutlass::gemm::GemmShape<128, 256, 32>;
  // This code section describes tile size a warp will compute
  using ShapeMMAWarp = cutlass::gemm::GemmShape<64, 64, 32>; 
  // This code section describes the size of MMA op
  using ShapeMMAOp = cutlass::gemm::GemmShape<16, 8, 8>;
  // This code section describes how threadblocks are scheduled on GPU
  using SwizzleThreadBlock = cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>;
  // This code section describes the epilogue part of the kernel
  using EpilogueOp = cutlass::epilogue::thread::LinearCombination<
      ElementOutput,                                     // <- data type of output matrix
      128 / cutlass::sizeof_bits<ElementOutput>::value,  // <- the number of elements per vectorized
                                                      // memory access. For a byte, it's 16
                                                      // elements. This becomes the vector width of
                                                      // math instructions in the epilogue too
      ElementAccumulator,                                // <- data type of accumulator
      ElementComputeEpilogue>;  // <- data type for alpha/beta in linear combination function
  constexpr int NumStages = 2;
  using Gemm = cutlass::gemm::device::Gemm<ElementInputA,
                                          LayoutInputA,
                                          ElementInputB,
                                          LayoutInputB,
                                          ElementOutput,
                                          LayoutOutput,
                                          ElementAccumulator,
                                          MMAOp,
                                          SmArch,
                                          ShapeMMAThreadBlock,
                                          ShapeMMAWarp,
                                          ShapeMMAOp,
                                          EpilogueOp,
                                          SwizzleThreadBlock,
                                          NumStages>;
  ElementComputeEpilogue alpha = ElementComputeEpilogue(1);
  ElementComputeEpilogue beta = ElementComputeEpilogue(0);

[hwu36]

Hmm. Does cutlass have this problem or your customized kernel have this problem?

You can use inline ptx to enforce 128bit shared memory. Just implement shared_store in https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/arch/memory.h . You don't need predicate for the shared store. Does shared store bank conflict happen in the mainloop or epilogue?

[FindDefinition]

I modify the volta tensorop example to fit turing fp16 gemm, then compile it by manual command instead of cutlass cmake.
both main loop and epilogue have bank conflicts. the main loop only has unexpected st.shared.u32/ld.shared.u32, in epilgue, the output smem padding is wrong. when I fix this problem, number of bank conflicts is decreased.
Update:
I try the example with cutlass cmake, unexpected st.shared.u32/ld.shared.u32 vanishes, but bank conflicts still exist until use correct output smem padding. I will figure out why manual command produce unexpected st.shared ptx

[hwu36]

Glad to see the bank conflict in the mainloop is gone.

Do you have interest in fixing the epilogue padding issues and file a PR?

[FindDefinition]

Yes, I will fix it and add some comments to explain why use that padding.

[hwu36]

you can use make VERBOSE=1 to see the command line used by cmake.

[butterluo]

Any update? Has the issue been fixed?

[hwu36]

i think our sm75 hmma kernel has no bank conflict.

[Maximilianxu]

FIXME: I think this issue still exists due to the epilogue padding size when the AccumulatorType is half and the warp stores accumulators to smem before the epilogue.