NVIDIA / NVIDIA/cutlass

uint1b_t ReduceArray specializations are broken (logical_or never compiles / always true; logical_and tests bytes not bits)

Open
#3,510 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

Description

The uint1b_t specializations of ReduceArrayOperation in include/cutlass/reduction/thread/reduction_operators.h (lines ~157-204) are broken in three independent ways.

logical_or: wrong parameter type AND constant-true result
template <int N>
struct ReduceArrayOperation<logical_or<uint1b_t>, uint1b_t, N> {
  ...
  uint1b_t operator()(
    logical_and<uint1b_t> const &reduction_op,   // <-- copy-paste: accepts only logical_and
    ArrayType const &array) const {
    ...
    bool item = true;                            // <-- identity of OR is false
    for (...) { bits = ptr[byte]; item = (item || bits); }
    return uint1b_t{item};                       // <-- always true

Consequences, both verified against the current tree:

  1. The public helper does not compile:
    cutlass::reduction::thread::detail::ReduceArray(logical_or<uint1b_t>{}, array) fails with "operator() cannot be called with the given argument list" because the parameter type names logical_and.
  2. Even calling it directly with a logical_and argument compiles but ignores its input entirely: item starts true and is only ever OR-ed into, so the result is constant true.
logical_and: tests whole bytes, not bits
bool item = false;
for (int byte = 0; byte < (N + 7) / 8; ++byte) {
  uint8_t bits = ptr[byte];
  item = (item || !bits);      // "some byte is entirely zero"
}
return uint1b_t{!item};

This returns "no byte is entirely zero", not "all elements are set". With N = 8 and element 7 cleared, byte 0 is 0x7F, no byte is zero, and the reduction reports true:

AND [1]*7+[0]: got=1 expected=0

There is also no masking of padding bits when N is not a multiple of 8.

Neither specialization has in-tree callers today, so both are latent public-API defects.

Suggested fix

For logical_or: change the parameter type to logical_or<uint1b_t> and initialize bool item = false;.
For logical_and: loop over the ceil(N/8) bytes, mask the final byte's padding bits, and test per-byte zero as intended - or simply iterate elements via the array API.

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 in include/cutlass/reduction/thread/reduction_operators.h around lines 157-204 and inspect the uint1b_t logical_or and logical_and ReduceArrayOperation specializations. Exercise the public ReduceArray helper with the reported OR and AND cases, including a cleared bit and an N value that is not a multiple of 8. Done means both reductions compile and return bitwise logical results without treating padding bits as elements.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
backend
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
78/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.