From 206801f775df73a79f93d068c41665c24c7f9d29 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Mon, 13 Oct 2025 14:11:28 +0000 Subject: [PATCH 1/4] Fix compute-sanitizer timeouts. --- ci/matrix.yaml | 2 ++ cub/test/catch2_test_device_topk_keys.cu | 6 ++++-- cub/test/run_test.cmake | 5 +++++ 3 files changed, 11 insertions(+), 2 deletions(-) diff --git a/ci/matrix.yaml b/ci/matrix.yaml index 78851db1c6e..e285bd8293d 100644 --- a/ci/matrix.yaml +++ b/ci/matrix.yaml @@ -21,6 +21,8 @@ workflows: # args: '--preset libcudacxx-cpp20 --lit-tests "cuda/utility/basic_any.pass.cpp"' } # override: + - { jobs: ['compute_init_nolid', 'compute_race_lid0', 'compute_init_lid0'], + project: 'cub', std: 'max', sm: 'gpu', gpu: 'rtxa6000', cmake_options: '-DCMAKE_CUDA_FLAGS=-lineinfo'} pull_request: # Old CTK: Oldest/newest supported host compilers: diff --git a/cub/test/catch2_test_device_topk_keys.cu b/cub/test/catch2_test_device_topk_keys.cu index d3288dd8556..ba1c2c7837e 100644 --- a/cub/test/catch2_test_device_topk_keys.cu +++ b/cub/test/catch2_test_device_topk_keys.cu @@ -149,7 +149,9 @@ C2H_TEST("DeviceTopK::{Min,Max}Keys work with iterators", "[keys][topk][device]" } } -C2H_TEST("DeviceTopK::{Min,Max}Keys works with a large number of items", "[keys][topk][device]", num_items_types) +C2H_TEST("DeviceTopK::{Min,Max}Keys works with a large number of items", + "[keys][topk][device][skip-cs-racecheck][skip-cs-initcheck][skip-cs-synccheck]", + num_items_types) try { using key_t = cuda::std::uint32_t; @@ -191,7 +193,7 @@ catch (std::bad_alloc& e) } C2H_TEST("DeviceTopK::{Min,Max}Keys works for different offset types for num_items and k", - "[keys][topk][device]", + "[keys][topk][device][skip-cs-racecheck][skip-cs-initcheck][skip-cs-synccheck]", k_items_types) try { diff --git a/cub/test/run_test.cmake b/cub/test/run_test.cmake index 76ce5e192a2..c5d00d0878f 100644 --- a/cub/test/run_test.cmake +++ b/cub/test/run_test.cmake @@ -100,6 +100,11 @@ elseif (MODE MATCHES "^compute-sanitizer-(.*)$") message(FATAL_ERROR "CCCL_SKIP_TEST:\n${TEST} intentionally throws CUDA errors. Skipping.") endif() + # The allocator test uses blocking-kernel tricks to explicitly control kernel lifetimes. + # This causes initcheck to deadlock. + if ("${TEST}" MATCHES "/cub.*.test.allocator$" AND tool STREQUAL "initcheck") + message(FATAL_ERROR "CCCL_SKIP_TEST:\n${TEST} uses blocking-kernels that break initcheck. Skipping.") + endif() if (TYPE STREQUAL "Catch2") list(APPEND ARGS From e1ff58f3a929b82e57ed4b95075ed46a8a441bf3 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Mon, 13 Oct 2025 17:48:04 +0000 Subject: [PATCH 2/4] Incorporate -d into bisect flags var. --- .github/workflows/git-bisect.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/workflows/git-bisect.yml b/.github/workflows/git-bisect.yml index a047221360f..af22e93dc73 100644 --- a/.github/workflows/git-bisect.yml +++ b/.github/workflows/git-bisect.yml @@ -256,7 +256,7 @@ jobs: - name: Run git bisect env: GITHUB_TOKEN: ${{ github.token }} - LAUNCH_ARGS: ${{ inputs.launch_args }} + LAUNCH_ARGS: '-d ${{ inputs.launch_args }}' GITHUB_REPOSITORY: ${{ github.repository }} # AWS credentials from the configure-aws-credentials action AWS_ACCESS_KEY_ID: ${{ env.AWS_ACCESS_KEY_ID }} @@ -297,7 +297,7 @@ jobs: mkdir -p /tmp/shared rc=0 set -x - .devcontainer/launch.sh -d ${LAUNCH_ARGS} ${GPU_ARGS} \ + .devcontainer/launch.sh ${LAUNCH_ARGS} ${GPU_ARGS} \ --env "AWS_ROLE_ARN=" \ --env "AWS_REGION=${AWS_REGION}" \ --env "SCCACHE_REGION=${AWS_REGION}" \ From 09c23ad9ec3a4577a7984211d67bc0d1a0539b69 Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Mon, 13 Oct 2025 17:48:19 +0000 Subject: [PATCH 3/4] Add another padding bits suppression. --- ci/compute-sanitizer-suppressions.xml | 34 +++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/ci/compute-sanitizer-suppressions.xml b/ci/compute-sanitizer-suppressions.xml index 37aa0fe5937..22a61d2bae0 100644 --- a/ci/compute-sanitizer-suppressions.xml +++ b/ci/compute-sanitizer-suppressions.xml @@ -44,6 +44,40 @@ + + + Initcheck + + Uninitialized __global__ memory read of size 2 bytes + 2 + + + DeviceReduceSingleTileKernel + + + + cuLaunchKernel + + + libcudart.* + + + cudaLaunchKernel + + + .*::detail::reduce::DeviceReduceSingleTileKernel.* + + + .*triple_chevron::doit_host.* + + + .*DeviceReduce::Reduce.* + + + bool thrust::.*::equal.* + + +