[QST] Why is there bank conflict in this simple layout?

I modified the tiled_copy.cu example in cute/tutorial to use the following layout

  auto tensor_shape = cute::Shape<cute::_128, cute::_32>{};
  auto block_shape = cute::Shape<cute::_16, cute::_32>{};
...
  Tensor tensor_S = make_tensor(make_gmem_ptr(thrust::raw_pointer_cast(d_S.data())), make_layout(tensor_shape, GenRowMajor{}));
  Tensor tensor_D = make_tensor(make_gmem_ptr(thrust::raw_pointer_cast(d_D.data())), make_layout(tensor_shape, GenRowMajor{}));
...
  // Thread arrangement
  Layout thr_layout = make_layout(make_shape(Int<16>{}, Int<16>{}));

  // Vector dimensions
  Layout vec_layout = make_layout(make_shape(Int<1>{}, Int<2>{}));
...
  using Atom = Copy_Atom<DefaultCopy, Element>;
  auto tiled_copy =
    make_tiled_copy(
      Atom{},                       // access size
      ThreadLayout{},               // thread layout
      VecLayout{});                 // vector layout (e.g. 1x2)
...

I would expect that there's no bank conflict when copying from global memory to shared memory, but running it under ncu suggests otherwise

ncu --section MemoryWorkloadAnalysis --metric l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum --launch-count 1 ./examples/cute/tutorial/tiled_copy
==PROF== Connected to process 3796174 (/home/sean/manifest2/packages/state_kernel/csrc/cutlass/build/examples/cute/tutorial/tiled_copy)
==PROF== Profiling "copy_kernel_vectorized" - 0 (1/1): 0%....50%....100% - 7 passes
Success.
==PROF== Disconnected from process 3796174
[3796174] tiled_copy@127.0.0.1
  void copy_kernel_vectorized<Tensor<ViewEngine<gmem_ptr<half_t *>>, Layout<tuple<tuple<C<16>, C<32>>, C<8>, C<1>>, tuple<tuple<C<32>, C<1>>, C<512>, C<0>>>>, Tensor<ViewEngine<gmem_ptr<half_t *>>, Layout<tuple<tuple<C<16>, C<32>>, C<8>, C<1>>, tuple<tuple<C<32>, C<1>>, C<512>, C<0>>>>, Layout<tuple<C<16>, C<16>>, tuple<C<1>, C<16>>>, Layout<tuple<C<1>, C<2>>, tuple<C<0>, C<1>>>>(T1, T2, T3, T4) (8, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.6
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                      128
    -------------------------------------------------------- ----------- ------------

    Section: Memory Workload Analysis
    --------------------------- ----------- ------------
    Metric Name                 Metric Unit Metric Value
    --------------------------- ----------- ------------
    Memory Throughput               Gbyte/s         3.78
    Mem Busy                              %         1.35
    Max Bandwidth                         %         0.74
    L1/TEX Hit Rate                       %        90.62
    L2 Compression Success Rate           %            0
    L2 Compression Ratio                               0
    L2 Hit Rate                           %        85.34
    Mem Pipes Busy                        %         0.36
    --------------------------- ----------- ------------

The full code is as follows:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cute/layout.hpp>
#include <cute/tensor.hpp>

#include "cutlass/util/print_error.hpp"
#include "cutlass/util/GPU_Clock.hpp"
#include "cutlass/util/helper_cuda.hpp"


/// Vectorized copy kernel.
///
/// Uses `make_tiled_copy()` to perform a copy using vector instructions. This operation
/// has the precondition that pointers are aligned to the vector size.
///
template <class TensorS, class TensorD, class ThreadLayout, class VecLayout>
__global__ void copy_kernel_vectorized(TensorS S, TensorD D, ThreadLayout, VecLayout)
{
  using namespace cute;
  using Element = typename TensorS::value_type;
  constexpr auto tensor_shape = cute::Shape<cute::_128, cute::_32>{};
  constexpr auto block_shape = cute::Shape<cute::_16, cute::_32>{};
  __shared__ char smem_[size(tensor_shape) * sizeof(Element)];

  Tensor smem_D = make_tensor(make_smem_ptr(const_cast<Element *>(reinterpret_cast<const Element *>(smem_))), tensor_shape);


  // Slice the tensors to obtain a view into each tile.
  Tensor tile_S = S(make_coord(_, _), blockIdx.x, blockIdx.y);  // (BlockShape_M, BlockShape_N)
  Tensor tile_D = D(make_coord(_, _), blockIdx.x, blockIdx.y);  // (BlockShape_M, BlockShape_N)
  Tensor tile_smem_D = tiled_divide(smem_D, block_shape)(make_coord(_, _), blockIdx.x, blockIdx.y);

  static_assert(rank(tile_smem_D) == 2, "tile_smem_D should have 2 dims");
  static_assert(size<0>(tile_smem_D) == size<0>(tile_D), "tile_smem_D and tile_D should have the same number of tiles in the first dimension");
  static_assert(size<1>(tile_smem_D) == size<1>(tile_D), "tile_smem_D and tile_D should have the same number of tiles in the second dimension");

  // Define `AccessType` which controls the size of the actual memory access.
  using AccessType = cutlass::AlignedArray<Element, size(VecLayout{})>;

  // A copy atom corresponds to one hardware memory access.
  using Atom = Copy_Atom<DefaultCopy, Element>;

  // Construct tiled copy, a tiling of copy atoms.
  //
  // Note, this assumes the vector and thread layouts are aligned with contigous data
  // in GMEM. Alternative thread layouts are possible but may result in uncoalesced
  // reads. Alternative vector layouts are also possible, though incompatible layouts
  // will result in compile time errors.
  auto tiled_copy =
    make_tiled_copy(
      Atom{},                       // access size
      ThreadLayout{},               // thread layout
      VecLayout{});                 // vector layout (e.g. 4x1)

  // Construct a Tensor corresponding to each thread's slice.
  auto thr_copy = tiled_copy.get_thread_slice(threadIdx.x);

  Tensor thr_tile_S = thr_copy.partition_S(tile_S);             // (CopyOp, CopyM, CopyN)
  Tensor thr_tile_D = thr_copy.partition_D(tile_D);             // (CopyOp, CopyM, CopyN)
  Tensor thr_tile_smem_D = thr_copy.partition_D(tile_smem_D); // (CopyOp, CopyM, CopyN)
  Tensor thr_tile_smem_S = thr_copy.partition_S(tile_smem_D); // (CopyOp, CopyM, CopyN)

  // Construct a register-backed Tensor with the same shape as each thread's partition
  // Use make_fragment because the first mode is the instruction-local mode
  Tensor fragment = make_fragment_like(thr_tile_D);             // (CopyOp, CopyM, CopyN)

  // Copy from GMEM to SMEM and from SMEM to GMEM
  copy(tiled_copy, thr_tile_S, thr_tile_smem_D);
  copy(tiled_copy, thr_tile_smem_S, thr_tile_D);
}

/// Main function
int main(int argc, char** argv)
{
  //
  // Given a 2D shape, perform an efficient copy
  //

  using namespace cute;
  using Element = cutlass::half_t;

  auto tensor_shape = cute::Shape<cute::_128, cute::_32>{};
  auto block_shape = cute::Shape<cute::_16, cute::_32>{};


  //
  // Allocate and initialize
  //

  thrust::host_vector<Element> h_S(size(tensor_shape));
  thrust::host_vector<Element> h_D(size(tensor_shape));

  for (size_t i = 0; i < h_S.size(); ++i) {
    h_S[i] = static_cast<Element>(static_cast<float>(i));
    h_D[i] = Element{};
  }

  thrust::device_vector<Element> d_S = h_S;
  thrust::device_vector<Element> d_D = h_D;

  //
  // Make tensors
  //

  Tensor tensor_S = make_tensor(make_gmem_ptr(thrust::raw_pointer_cast(d_S.data())), make_layout(tensor_shape, GenRowMajor{}));
  Tensor tensor_D = make_tensor(make_gmem_ptr(thrust::raw_pointer_cast(d_D.data())), make_layout(tensor_shape, GenRowMajor{}));

  //
  // Tile tensors
  //

  // Define a statically sized block (M, N).
  // Note, by convention, capital letters are used to represent static modes.


  if ((size<0>(tensor_shape) % size<0>(block_shape)) || (size<1>(tensor_shape) % size<1>(block_shape))) {
    std::cerr << "The tensor shape must be divisible by the block shape." << std::endl;
    return -1;
  }
  // Equivalent check to the above
  if (not weakly_compatible(block_shape, tensor_shape)) {
    std::cerr << "Expected the tensors to be weakly compatible with the block_shape." << std::endl;
    return -1;
  }

  // Tile the tensor (m, n) ==> ((M, N), m', n') where (M, N) is the static tile
  // shape, and modes (m', n') correspond to the number of tiles.
  //
  // These will be used to determine the CUDA kernel grid dimensions.
  Tensor tiled_tensor_S = tiled_divide(tensor_S, block_shape);      // ((M, N), m', n')
  Tensor tiled_tensor_D = tiled_divide(tensor_D, block_shape);      // ((M, N), m', n')

  // Thread arrangement
  Layout thr_layout = make_layout(make_shape(Int<16>{}, Int<16>{}));

  // Vector dimensions
  Layout vec_layout = make_layout(make_shape(Int<1>{}, Int<2>{}));

  //
  // Determine grid and block dimensions
  //

  dim3 gridDim (size<1>(tiled_tensor_D), size<2>(tiled_tensor_D));   // Grid shape corresponds to modes m' and n'
  dim3 blockDim(size(thr_layout));

  //
  // Launch the kernel
  //
  copy_kernel_vectorized<<< gridDim, blockDim >>>(
    tiled_tensor_S,
    tiled_tensor_D,
    thr_layout,
    vec_layout);

  cudaError result = cudaDeviceSynchronize();
  if (result != cudaSuccess) {
    std::cerr << "CUDA Runtime error: " << cudaGetErrorString(result) << std::endl;
    return -1;
  }

  //
  // Verify
  //

  h_D = d_D;

  int32_t errors = 0;
  int32_t const kErrorLimit = 10;

  for (size_t i = 0; i < h_D.size(); ++i) {
    if (h_S[i] != h_D[i]) {
      std::cerr << "Error. S[" << i << "]: " << h_S[i] << ",   D[" << i << "]: " << h_D[i] << std::endl;

      if (++errors >= kErrorLimit) {
        std::cerr << "Aborting on " << kErrorLimit << "nth error." << std::endl;
        return -1;
      }
    }
  }

  std::cout << "Success." << std::endl;

  return 0;
}

I'm running on an A6000 so I suspect that either I misunderstood how tiled_copy works or the scheduling optimizer (from Volta) kicks in and groups the threads in a way that messed up the bank access?

Start with cute/tutorial/tiled_copy.cu and reproduce the reported copy_kernel_vectorized launch using the provided ncu MemoryWorkloadAnalysis command on the A6000 configuration. Trace the tiled_copy thread and vector layouts, then compare the reported shared-memory bank-conflict metric with the copy results. Done means documenting whether the metric reflects the layout or scheduling behavior and identifying the relevant correction or explanation.