[BUG]: Assertion fires when running CUB's decoupled look back tests through compute-sanitizer --tool synccheck

[alliepiper]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Runtime Error

Component

CUB

Describe the bug

Running the decoupled lookback test through synccheck trips an assertion in the implementation. This only happens under synccheck.

/home/coder/cccl/cub/test/catch2_test_device_decoupled_look_back.cu:95:
  Assertion `tile_aggregate == prefix.GetBlockAggregate()` failed.

This is happening in the following kernel:

void decoupled_look_back_kernel(
  cub::CUB_200500_600_700_800_NS::ScanTileState<
    MessageT, 
    cub::CUB_200500_600_700_800_NS::Traits<MessageT>::PRIMITIVE
  >,
  MessageT *)
  
[with MessageT = unsigned char]:

The assertion fires for multiple blocks, but not all blocks:

block: [3517,0,0], thread: [32,0,0]
block: [3517,0,0], thread: [33,0,0]
block: [3517,0,0], thread: [34,0,0]
block: [3517,0,0], thread: [35,0,0]
block: [3517,0,0], thread: [36,0,0]
block: [3517,0,0], thread: [37,0,0]
block: [3517,0,0], thread: [38,0,0]
block: [3517,0,0], thread: [39,0,0]
block: [3517,0,0], thread: [40,0,0]
block: [3517,0,0], thread: [41,0,0]
block: [3517,0,0], thread: [42,0,0]
block: [3517,0,0], thread: [43,0,0]
block: [3517,0,0], thread: [44,0,0]
block: [3517,0,0], thread: [45,0,0]
block: [3517,0,0], thread: [46,0,0]
block: [3517,0,0], thread: [47,0,0]
block: [3517,0,0], thread: [48,0,0]
block: [3517,0,0], thread: [49,0,0]
block: [3517,0,0], thread: [50,0,0]
block: [3517,0,0], thread: [51,0,0]
block: [3517,0,0], thread: [52,0,0]
block: [3517,0,0], thread: [53,0,0]
block: [3517,0,0], thread: [54,0,0]
block: [3517,0,0], thread: [55,0,0]
block: [3517,0,0], thread: [56,0,0]
block: [3517,0,0], thread: [57,0,0]
block: [3517,0,0], thread: [58,0,0]
block: [3517,0,0], thread: [59,0,0]
block: [3517,0,0], thread: [60,0,0]
block: [3517,0,0], thread: [61,0,0]
block: [3517,0,0], thread: [62,0,0]
block: [3517,0,0], thread: [63,0,0]
block: [3523,0,0], thread: [192,0,0]
block: [3523,0,0], thread: [193,0,0]
block: [3523,0,0], thread: [194,0,0]
block: [3523,0,0], thread: [195,0,0]
block: [3523,0,0], thread: [196,0,0]
block: [3523,0,0], thread: [197,0,0]
block: [3523,0,0], thread: [198,0,0]
block: [3523,0,0], thread: [199,0,0]
block: [3523,0,0], thread: [200,0,0]
block: [3523,0,0], thread: [201,0,0]
block: [3523,0,0], thread: [202,0,0]
block: [3523,0,0], thread: [203,0,0]
block: [3523,0,0], thread: [204,0,0]
block: [3523,0,0], thread: [205,0,0]
block: [3523,0,0], thread: [206,0,0]
block: [3523,0,0], thread: [207,0,0]
block: [3523,0,0], thread: [208,0,0]
block: [3523,0,0], thread: [209,0,0]
block: [3523,0,0], thread: [210,0,0]
block: [3523,0,0], thread: [211,0,0]
block: [3523,0,0], thread: [212,0,0]
block: [3523,0,0], thread: [213,0,0]
block: [3523,0,0], thread: [214,0,0]
block: [3523,0,0], thread: [215,0,0]
block: [3523,0,0], thread: [216,0,0]
block: [3523,0,0], thread: [217,0,0]
block: [3523,0,0], thread: [218,0,0]
block: [3523,0,0], thread: [219,0,0]
block: [3523,0,0], thread: [220,0,0]
block: [3523,0,0], thread: [221,0,0]
block: [3523,0,0], thread: [222,0,0]
block: [3523,0,0], thread: [223,0,0]
block: [3913,0,0], thread: [96,0,0]
block: [3913,0,0], thread: [97,0,0]
block: [3913,0,0], thread: [98,0,0]
block: [3913,0,0], thread: [99,0,0]
block: [3913,0,0], thread: [100,0,0]
block: [3913,0,0], thread: [101,0,0]
block: [3913,0,0], thread: [102,0,0]
block: [3913,0,0], thread: [103,0,0]
block: [3913,0,0], thread: [104,0,0]
block: [3913,0,0], thread: [105,0,0]
block: [3913,0,0], thread: [106,0,0]
block: [3913,0,0], thread: [107,0,0]
block: [3913,0,0], thread: [108,0,0]
block: [3913,0,0], thread: [109,0,0]
block: [3913,0,0], thread: [110,0,0]
block: [3913,0,0], thread: [111,0,0]
block: [3913,0,0], thread: [112,0,0]
block: [3913,0,0], thread: [113,0,0]
block: [3913,0,0], thread: [114,0,0]
block: [3913,0,0], thread: [115,0,0]
block: [3913,0,0], thread: [116,0,0]
block: [3913,0,0], thread: [117,0,0]
block: [3913,0,0], thread: [118,0,0]
block: [3913,0,0], thread: [119,0,0]
block: [3913,0,0], thread: [120,0,0]
block: [3913,0,0], thread: [121,0,0]
block: [3913,0,0], thread: [122,0,0]
block: [3913,0,0], thread: [123,0,0]
block: [3913,0,0], thread: [124,0,0]
block: [3913,0,0], thread: [125,0,0]
block: [3913,0,0], thread: [126,0,0]
block: [3913,0,0], thread: [127,0,0]

How to Reproduce

Run the decoupled look back test through synccheck. This executable was compiled with -lineinfo, unsure if that's relevant.

~/cccl$ compute-sanitizer --tool synccheck build/cuda12.4-gcc13/cub-cpp17/bin/cub.cpp17.test.device_decoupled_look_back 
========= COMPUTE-SANITIZER
/home/coder/cccl/cub/test/catch2_test_device_decoupled_look_back.cu:95: void decoupled_look_back_kernel(cub::CUB_200500_600_700_800_NS::ScanTileState<MessageT, cub::CUB_200500_600_700_800_NS::Traits<MessageT>::PRIMITIVE>, MessageT *) [with MessageT = unsigned char]: block: [3517,0,0], thread: [32,0,0] Assertion `tile_aggregate == prefix.GetBlockAggregate()` failed.
(same error repeats)

Expected behavior

Assertions should never fire.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response