uint1b_t ReduceArray specializations are broken (logical_or never compiles / always true; logical_and tests bytes not bits)
Nobody has claimed this yet.
- 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:
- 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 nameslogical_and. - Even calling it directly with a
logical_andargument compiles but ignores its input entirely:itemstartstrueand is only ever OR-ed into, so the result is constanttrue.
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
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- 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