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:
- 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;
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.
Description
The ColumnMajor-output specializations of
device::GemmUniversalWithBroadcastanddevice::GemmUniversalStreamKWithBroadcastimplement their output layout by transposing the problem (to_underlying_argumentscallsargs.transposed_problem()), but the transpose does not touch the broadcast-vector arguments even though they are consumed in the transposed coordinate system.while the fused epilogue indexes the bias by the underlying kernel's output columns:
After the transpose, the underlying kernel's column extent is the user's M, so:
n(the documented per-column broadcast for the user'sm x noutput) is read at indices up tom - 1: out-of-bounds reads wheneverm > n, and wrong elements wheneverm != n;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-317andgemm_universal_streamk_with_broadcast.h:315-317both route through this transpose. Nothing in-tree exercises the ColumnMajor-output specialization withptr_Vectorset, 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_Vectorfor the ColumnMajor-output specializations instead of silently misindexing.