[CUB] Test and document in-place BlockExchange overloads#10227
Conversation
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesIn-place BlockExchange API
Assessment against linked issues
Comment Warning |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
cub/test/catch2_test_block_exchange.cu (1)
24-28: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Parameterize
WarpTimeSlicingand instantiatecub::BlockExchange<T, BlockThreads, ItemsPerThread, WarpTimeSlicing>. The current alias always selects the defaultfalse, leaving all eight in-place overloads untested against their time-sliced implementations.As per path instructions, focus on test coverage.
Source: Path instructions
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: fade38b4-2d9d-4b76-9163-afacb85d2e3b
📒 Files selected for processing (2)
cub/cub/block/block_exchange.cuhcub/test/catch2_test_block_exchange.cu
| _CCCL_DEVICE _CCCL_FORCEINLINE void StripedToBlocked(T (&items)[ItemsPerThread]) | ||
| { | ||
| StripedToBlocked(items, items); | ||
| } |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
python3 - <<'PY'
from pathlib import Path
path = Path("cub/cub/block/block_exchange.cuh")
lines = path.read_text().splitlines()
for start,end in [(890,920),(965,990),(1036,1060),(1108,1135),(1164,1240),(1290,1320),(1380,1410)]:
print(f"\n--- {start}-{end} ---")
for i in range(start-1, min(end, len(lines))):
print(f"{i+1:4}: {lines[i]}")
PYRepository: NVIDIA/cccl
Length of output: 8469
🏁 Script executed:
python3 - <<'PY'
from pathlib import Path
import re
path = Path("cub/cub/block/block_exchange.cuh")
text = path.read_text().splitlines()
names = [
"StripedToBlocked(",
"BlockedToStriped(",
"WarpStripedToBlocked(",
"BlockedToWarpStriped(",
"ScatterToBlocked(",
"ScatterToStriped(",
"ScatterToStripedGuarded(",
"ScatterToStripedFlagged(",
]
for name in names:
print(f"\n### {name}")
for i, line in enumerate(text, 1):
if name in line:
# print a small window around each occurrence
for j in range(max(1, i-2), min(len(text), i+4)+1):
print(f"{j:4}: {text[j-1]}")
print("---")
PYRepository: NVIDIA/cccl
Length of output: 18952
🏁 Script executed:
rg -n "noexcept" cub/cub/block/block_exchange.cuhRepository: NVIDIA/cccl
Length of output: 149
suggestion: add noexcept to these forwarding overloads
Each wrapper only forwards to another BlockExchange method, so the in-place overloads should be marked noexcept too:
cub/cub/block/block_exchange.cuh#L903-L906cub/cub/block/block_exchange.cuh#L974-L977cub/cub/block/block_exchange.cuh#L1045-L1048cub/cub/block/block_exchange.cuh#L1119-L1122cub/cub/block/block_exchange.cuh#L1173-L1176cub/cub/block/block_exchange.cuh#L1224-L1227cub/cub/block/block_exchange.cuh#L1301-L1305cub/cub/block/block_exchange.cuh#L1393-L1397
📍 Affects 1 file
cub/cub/block/block_exchange.cuh#L903-L906(this comment)cub/cub/block/block_exchange.cuh#L974-L977cub/cub/block/block_exchange.cuh#L1045-L1048cub/cub/block/block_exchange.cuh#L1119-L1122cub/cub/block/block_exchange.cuh#L1173-L1176cub/cub/block/block_exchange.cuh#L1224-L1227cub/cub/block/block_exchange.cuh#L1301-L1305cub/cub/block/block_exchange.cuh#L1393-L1397
Source: Coding guidelines
1 participant
Description
Closes #932
This PR adds direct test coverage and documentation for all eight in-place
cub::BlockExchangeoverloads.The new Catch2 tests cover blocked, striped, warp-striped, and scatter exchanges across multiple block sizes, item counts, and data types. The in-place overloads are moved next to their corresponding output-parameter overloads, documented, and exposed to Doxygen by removing the
_CCCL_DOXYGEN_INVOKEDguard.Testing
pre-commit run --files cub/cub/block/block_exchange.cuh cub/test/catch2_test_block_exchange.cuChecklist