NVIDIA / NVIDIA/cutlass

copy_if with a Copy_Atom dispatches on the src slice only and can store out of bounds when dst extents differ

Open
#3,488 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Which component has the problem?

CuTe C++ - include/cute/algorithm/copy.hpp, the copy_if overloads

Describe the bug

copy_if(Copy_Atom<...>, pred, src, dst) only static-asserts that src and dst have equal rank; it never compares extents. Per rest-slice it calls Copy_Atom::call(pred_v, src_v, dst_v), and the atom's instruction dispatch tests the vector width against the src slice only (copy_atom.hpp, is_constant<NumValSrc, size(src)>). When dst has a smaller extent, one wide instruction is emitted into the smaller dst slice and stores out of bounds. There is no diagnostic of any kind before the corruption.

The sibling entry point copy(src, dst) handles the same mismatch safely: its vectorization machinery recomputes max_common_vector/alignment from both layouts and degrades to elementwise copies.

Minimal repro (host-only, ASan):

#include <cute/tensor.hpp>
#include <cute/algorithm/copy.hpp>
#include <cstdio>
using namespace cute;

int main() {
  float sa[12], da[6];
  Tensor src = make_tensor(&sa[0], Shape<_4,_3>{});   // 12 elems
  Tensor dst = make_tensor(&da[0], Shape<_2,_3>{});   // 6 elems (typo)
  bool m[12]; for (int i=0;i<12;++i) m[i]=true;
  Tensor prd = make_tensor(&m[0], Shape<_4,_3>{});

  copy_if(Copy_Atom<UniversalCopy<uint128_t>, float>{}, prd, src, dst);
}
$ g++ -std=c++17 -g -fsanitize=address ... && ./a.out
==30==ERROR: AddressSanitizer: stack-buffer-overflow ...
WRITE of size 16 at 0x... thread T0
    #2 cute::copy_unpack<...> (include/cute/atom/copy_atom.hpp)
    #3 cute::Copy_Atom<...>::call(...)
    #4 cute::copy_if<...> (include/cute/algorithm/copy.hpp:69)

With plain copy(src, dst) on identical shapes the program runs ASan-clean and produces correct results, so the two entry points disagree on error behavior for the same misuse.

The generic elementwise copy_if at copy.hpp:44-62 has a milder version of the same pattern: it bounds the loop by size(dst) and indexes pred(i)/src(i) without any cross-check.

Expected behavior

Either a compile-time/static diagnostic when size(src) != size(dst) per rest-slice (mirroring what copy() effectively enforces), or dispatch keyed off both slices. Raw atom calls are normally fed by TiledCopy partitioning which guarantees matching widths, so this may be an accepted contract - but today the failure mode is silent memory corruption on a plausible typo, where the sibling API reports nothing yet also does the right thing.

Environment details
  • Current main, header-only, host-only repro (no GPU needed), g++ + AddressSanitizer

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start by reproducing the host-only example with AddressSanitizer, then inspect the copy_if overloads in include/cute/algorithm/copy.hpp and the dispatch in include/cute/atom/copy_atom.hpp. Compare this path with copy()'s vectorization checks and the generic elementwise overload. Done means mismatched source and destination extents produce a defined diagnostic or safe behavior, with a regression covering the repro.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
backend, performance
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Active
Clarity
Mostly clear
Newbie friendliness
68/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.