Skip to content

Fix compute-sanitizer timeouts. #6206

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

Draft
wants to merge 5 commits into
base: main
Choose a base branch
from
+48 −5
Draft

Fix compute-sanitizer timeouts. #6206

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
72d17ac
Fix compute-sanitizer timeouts.
alliepiper Oct 13, 2025
a9f352c
Incorporate -d into bisect flags var.
alliepiper Oct 13, 2025
04356d9
Add another padding bits suppression.
alliepiper Oct 13, 2025
8891a79
Skip a set of very slow tests with initcheck.
alliepiper Oct 13, 2025
a8a5b32
Merge remote-tracking branch 'origin/main' into cs-timeouts
alliepiper Oct 13, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions .github/workflows/git-bisect.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 }}
Expand Down Expand Up @@ -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}" \
Expand Down
34 changes: 34 additions & 0 deletions ci/compute-sanitizer-suppressions.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,40 @@
</frame>
</hostStack>
</record>
<!-- Another variant of the above -->
<record>
<kind>Initcheck</kind>
<what>
<text>Uninitialized __global__ memory read of size 2 bytes</text>
<size>2</size>
</what>
<where>
<func>DeviceReduceSingleTileKernel</func>
</where>
<hostStack>
<frame>
<func>cuLaunchKernel</func>
</frame>
<frame>
<func>libcudart.*</func>
</frame>
<frame>
<func>cudaLaunchKernel</func>
</frame>
<frame>
<func>.*::detail::reduce::DeviceReduceSingleTileKernel.*</func>
</frame>
<frame>
<func>.*triple_chevron::doit_host.*</func>
</frame>
<frame>
<func>.*DeviceReduce::Reduce.*</func>
</frame>
<frame>
<func>bool thrust::.*::equal.*</func>
</frame>
</hostStack>
</record>
<!--
Similar to the above, thrust::equal copies a tuple<bool, OffsetT> from host -> device
with the result of the comparison. The padding bytes trigger host API initialization
Expand Down
2 changes: 2 additions & 0 deletions ci/matrix.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
2 changes: 1 addition & 1 deletion cub/test/catch2_test_device_merge_sort.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -273,7 +273,7 @@ C2H_TEST("DeviceMergeSort::SortKeys works", "[merge][sort][device]", wide_key_ty

C2H_TEST("DeviceMergeSort::StableSortKeysCopy works and performs a stable sort when there are a lot sort-keys that "
"compare equal",
"[merge][sort][device][skip-cs-racecheck][skip-cs-memcheck]")
"[merge][sort][device][skip-cs-initcheck][skip-cs-racecheck][skip-cs-memcheck]")
{
using key_t = c2h::custom_type_t<c2h::equal_comparable_t, c2h::less_comparable_t>;
using offset_t = std::size_t;
Expand Down
6 changes: 4 additions & 2 deletions cub/test/catch2_test_device_topk_keys.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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
{
Expand Down
5 changes: 5 additions & 0 deletions cub/test/run_test.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Loading