Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[BUG]: Race reported in cub::BlockReduceRaking #1902

Closed
1 task done
Tracked by #1618
alliepiper opened this issue Jun 21, 2024 · 0 comments · Fixed by #2699
Closed
1 task done
Tracked by #1618

[BUG]: Race reported in cub::BlockReduceRaking #1902

alliepiper opened this issue Jun 21, 2024 · 0 comments · Fixed by #2699
Labels
bug Something isn't working right.

Comments

@alliepiper
Copy link
Collaborator

Is this a duplicate?

Type of Bug

Something else

Component

CUB

Describe the bug

https://github.com/NVIDIA/cccl/actions/runs/9606424817/job/26497170643?pr=1879

  	compute-sanitizer --tool racecheck --check-device-heap yes --leak-check full --padding 512 --track-stream-ordered-races all --check-warpgroup-mma yes --require-cuda-init no --check-exit-code yes --error-exitcode 1 --nvtx true /home/coder/cccl/build/cuda12.4-gcc13/cub-cpp17/bin/cub.cpp17.test.block_reduce.dimx_65.dimyz_1
  ========= COMPUTE-SANITIZER
  ========= Warning: Race reported between Write access at void cub::CUB_200500_600_700_800_NS::ThreadStoreVolatilePtr<unsigned char>(T1 *, T1, cub::CUB_200500_600_700_800_NS::Int2Type<(int)1>)+0x250 in /home/coder/cccl/cub/cub/thread/thread_store.cuh:284
  =========     and Read access at unsigned char cub::CUB_200500_600_700_800_NS::BlockReduceRaking<unsigned char, (int)65, (int)1, (int)1, (int)0>::Reduce<(bool)1, cuda::std::__4::plus<void>>(unsigned char, int, T2)+0x1b0 in /home/coder/cccl/cub/cub/block/specializations/block_reduce_raking.cuh:222 [7 hazards]
  =========     and Read access at unsigned char cub::CUB_200500_600_700_800_NS::BlockReduceRaking<unsigned char, (int)65, (int)1, (int)1, (int)0>::RakingReduction<(bool)1, cuda::std::__4::plus<void>, (int)1>(T2, unsigned char *, unsigned char, int, cub::CUB_200500_600_700_800_NS::Int2Type<T3>)+0x1c0 in /home/coder/cccl/cub/cub/block/specializations/block_reduce_raking.cuh:161 [7 hazards]
  =========     and Read access at unsigned char cub::CUB_200500_600_700_800_NS::BlockReduceRaking<unsigned char, (int)65, (int)1, (int)1, (int)0>::RakingReduction<(bool)1, cuda::std::__4::plus<void>, (int)1>(T2, unsigned char *, unsigned char, int, cub::CUB_200500_600_700_800_NS::Int2Type<T3>)+0x1d0 in /home/coder/cccl/cub/cub/block/specializations/block_reduce_raking.cuh:161 [7 hazards]

How to Reproduce

compute-sanitizer --tool racecheck --check-device-heap yes --leak-check full --padding 512 --track-stream-ordered-races all --check-warpgroup-mma yes --require-cuda-init no --check-exit-code yes --error-exitcode 1 --nvtx true /home/coder/cccl/build/cuda12.4-gcc13/cub-cpp17/bin/cub.cpp17.test.block_reduce.dimx_65.dimyz_1

Expected behavior

No diagnostics emitted.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response
@alliepiper alliepiper added the bug Something isn't working right. label Jun 21, 2024
This was referenced Jun 21, 2024
Draft
Open
alliepiper added a commit to alliepiper/cccl that referenced this issue Nov 4, 2024
@alliepiper
Fix race condition in block_reduce_raking.
5391cce 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes NVIDIA#1902.
alliepiper added a commit to alliepiper/cccl that referenced this issue Nov 4, 2024
@alliepiper
Fix race condition in block_reduce_raking.
c87883d 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes NVIDIA#1902.
@alliepiper alliepiper mentioned this issue Nov 4, 2024
Merged
alliepiper added a commit to alliepiper/cccl that referenced this issue Nov 5, 2024
@alliepiper
Fix race condition in block_reduce_raking.
7848b3f 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes NVIDIA#1902.
alliepiper added a commit to alliepiper/cccl that referenced this issue Nov 5, 2024
@alliepiper
Fix race condition in block_reduce_raking.
c69e266 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes NVIDIA#1902.
alliepiper added a commit that referenced this issue Nov 5, 2024
@alliepiper
Fix race condition in block_reduce_raking. (#2699)
cb8707c 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes #1902.
pciolkosz pushed a commit to pciolkosz/cccl that referenced this issue Nov 6, 2024
@alliepiper @pciolkosz
Fix race condition in block_reduce_raking. (NVIDIA#2699)
0ac99cc 
`temp_storage.raking_grid` and `temp_storage.warp_storage` are aliased.
We need to sync between using shmem to compute partials via `RakingReduction` and `WarpReduce`-ing those partial.

Fixes NVIDIA#1902.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
bug Something isn't working right.
Projects
CCCL
Status: Done
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Fix race condition in block_reduce_raking. alliepiper/cccl
1 participant
@alliepiper