Skip to content

GemmUniversalWithBroadcast ColumnMajor-output transpose leaves the bias vector indexed in the wrong coordinate system #3540

Description

@VaggelisGian

Description

The ColumnMajor-output specializations of device::GemmUniversalWithBroadcast and device::GemmUniversalStreamKWithBroadcast implement their output layout by transposing the problem (to_underlying_arguments calls args.transposed_problem()), but the transpose does not touch the broadcast-vector arguments even though they are consumed in the transposed coordinate system.

// include/cutlass/gemm/kernel/gemm_with_fused_epilogue.h:208-217
Arguments transposed_problem() const {
  Arguments args(*this);
  std::swap(args.problem_size.m(), args.problem_size.n());
  std::swap(args.ptr_A, args.ptr_B);
  std::swap(args.lda, args.ldb);
  std::swap(args.batch_stride_A, args.batch_stride_B);
  return args;      // ptr_Vector / ldr / batch_stride_Vector untouched
}

while the fused epilogue indexes the bias by the underlying kernel's output columns:

// include/cutlass/gemm/kernel/gemm_with_fused_epilogue.h:630-631 (and :762-763)
if (ptr_Vector) {
  ptr_Vector += threadblock_offset.column() + threadblock_tile_offset.m() * params.ldr;
}

After the transpose, the underlying kernel's column extent is the user's M, so:

  1. a bias vector of length n (the documented per-column broadcast for the user's m x n output) is read at indices up to m - 1: out-of-bounds reads whenever m > n, and wrong elements whenever m != n;
  2. ldr, which scales the per-M-tile offset, keeps its pre-transpose meaning relative to an axis that no longer corresponds.

gemm_universal_with_broadcast.h:315-317 and gemm_universal_streamk_with_broadcast.h:315-317 both route through this transpose. Nothing in-tree exercises the ColumnMajor-output specialization with ptr_Vector set, which is why this has gone unnoticed.

Suggested fix

Decide and document the post-transpose convention for the vector: either swap the vector pointer semantics alongside the problem (and adjust the epilogue indexing), or reject ptr_Vector for the ColumnMajor-output specializations instead of silently misindexing.

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions