[BUG]: Racy non-trivial runs test

[gevtushenko]

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

The cub::DeviceRunLengthEncode::NonTrivialRuns test failed in ctk 12.2, llvm15, c++14 container with the following error message:

Pointer DeviceRunLengthEncode::NonTrivialRuns cub::CUB on 1000000 items, 174932 segments (avg run length 5.717), {6uchar2 key, i offset, i length}, max_segment 16, entropy_reduction 1
Synchronizing...
Synchronizing...
INCORRECT: [173]: 0 != 968	 Offsets FAIL
INCORRECT: [172]: 0 != 2	 Lengths FAIL
Count PASS

Consecutive run was successful. It's unclear if there's a race in the algorithm or a particular random workloads leads to the issue. The only change in 2.2 that's related to non-trivial runs on V100 is #294.

How to Reproduce

Run ./bin/cub.cpp20.test.device_run_length_encode.cdp_0 and hope that the issue take place.

Expected behavior

The test is always passing.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[elstehle]

To share a few findings so far:

• I have to run all the RLE test suite on average ~1300 times to observe the failure
• So far, it was always the same test case failing
• So far, I was only able to reproduce with clang host compiler
• Running with compute-sanitizer --tool racecheck reports some hazards, that I'm currently investigating
• The failure also occurs before latest fix for non-trivial runs
• I see races being reported also for commits before the latest fix for non-trivial runs