[BUG] warpspeed scan causes OOB reads in some Thrust tests #8136
Description
QA reports:
88/350 Test #88: thrust.cpp.cuda.test.scan ..........................................Subprocess aborted***Exception: 0.36 sec
Testing Device 0: "NVIDIA GeForce RTX 5090"
Running 29 unit tests.
............../host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [168,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [170,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [171,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [172,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [173,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [174,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [177,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [178,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [179,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [180,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [182,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [183,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [185,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [186,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [187,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [188,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [224,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [225,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [227,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [229,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [231,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [233,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [235,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [237,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [239,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [241,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [243,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [245,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [247,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [249,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [251,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [253,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [255,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [192,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [193,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [194,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [195,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [198,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [199,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [200,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [201,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [202,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [203,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [205,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [207,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [209,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [211,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [213,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [215,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [217,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [219,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [221,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
/host_compiler_testing/COPTer/projects/files/thrust_cpp17/sources/lib/cmake/libcudacxx/../../../libcudacxx/include/cuda/std/array:172: operator[]: block: [0,0,0], thread: [223,0,0] Assertion `out-of-bounds access in std::array<T, N>` failed.
terminate called after throwing an instance of 'thrust::_V_300400_SM_1200::system::system_error'
what(): inclusive_scan failed to synchronize: cudaErrorAssert: device-side assert triggered
See NVBug5999870. They use an unreleased nvcc newer than 13.2, but I see the same issue with nvcc 13.2 on RTX 5090 in this branch #7565 (which has the same SASS as main, but different host code). I can't reproduce it on main. Maybe the CI did not catch it cause it uses an RTX PRO 6000.
Metadata
Metadata
Assignees
Labels
No labels
Type
Projects
Status