[BUG] [libcudf] Local Tests Fail Only When Built with PTDS

[JigaoLuo]

Describe the bug

I've encountered an issue where local tests execute successfully when built without the PTDS option. However, upon building with PTDS, two local tests fail, and I have the corresponding output. This behavior strongly indicates a bug.

Moreover, I noticed that there's no PTDS-related configuration in the CI. Hence, I believe it's essential to incorporate PTDS builds into the CI process.

Below are the details of two build scenarios based on the latest commit in the 25.06 branch, using a conda-provided virtual environment:

Steps/Code to reproduce bug: Build with PTDS

./build.sh --ptds tests libcudf
ctest --test-dir cpp/build

Then QUANTILES_TEST and DATA_CHUNK_SOURCE_TEST failed:

98% tests passed, 2 tests failed out of 108

Total Test time (real) = 751.07 sec

The following tests FAILED:
  11 - QUANTILES_TEST (Failed)
  37 - DATA_CHUNK_SOURCE_TEST (Failed)
Errors while running CTest
Output from these tests are in: /home/jluo/cudf/cpp/build/Testing/Temporary/LastTest.log
Use "--rerun-failed --output-on-failure" to re-run the failed cases verbosely.

rerun-failed

Then I did ctest --test-dir cpp/build --rerun-failed --output-on-failure and attach the relevant output file here:

PTDS_TEST_OUTPUT.txt

Expected behavior: Build without PTDS

I built the project using the following commands:

./build.sh tests libcudf
ctest --test-dir cpp/build

In this scenario, all local tests passed as expected:

...
107/108 Test #107: STREAM_TRANSPOSE_TEST ............   Passed    6.85 sec
        Start 108: STREAM_UNARY_TEST
108/108 Test #108: STREAM_UNARY_TEST ................   Passed    1.83 sec

100% tests passed, 0 tests failed out of 108

Total Test time (real) = 825.47 sec

---

Environment overview (please complete the following information)

• Environment location: Bare-metal
• Method of cuDF install: from source

(If you're unable to reproduce this bug on your end, please let me know. I'll then show the output of print_env.sh and conduct a more in-depth investigation to determine the root cause of the test failures.)

[GregoryKimball]

Thank you @JigaoLuo for studying this. Many of our partners rely on PTDS builds of libcudf, so seeing these failures is concerning.

[davidwendt]

The problem appears to be the ctest execution setup which imposes some stream detection utilities. This appears to be what is causing the failures. If you run the tests directly without ctest then these 2 tests should work ok (as well as all the other tests).

gtests/QUANTILES_TEST
gtests/DATA_CHUNK_SOURCE_TEST

We'll need to figure out why the ctest usage is not working.

[davidwendt]

I spent sometime on this today but did not make much progress. It seems to be mainly a test issue and not a problem in the libcudf code. I was able recreate the error outside of ctest by setting the following environment variables

  GTEST_CUDF_STREAM_MODE=new_cudf_default
  LD_PRELOAD=/cudf/cpp/build/libcudf_identify_stream_usage_mode_cudf.so

And then running gtest on the QUANTILES_TEST with a filter that executes only the first 2 tests:

$ ./ctest_gtest.sh gtests/QUANTILES_TEST --gtest_filter=PercentileApproxInputTypesTest/0*
Note: Google Test filter = PercentileApproxInputTypesTest/0*
[==========] Running 4 tests from 1 test suite.
[----------] Global test environment set-up.
[----------] 4 tests from PercentileApproxInputTypesTest/0, where TypeParam = signed char
[ RUN      ] PercentileApproxInputTypesTest/0.Simple
[       OK ] PercentileApproxInputTypesTest/0.Simple (1022 ms)
[ RUN      ] PercentileApproxInputTypesTest/0.Grouped
/cudf/cpp/tests/utilities/column_utilities.cu:260: Failure
Expected equality of these values:
  lhs_size
    Which is: 6
  rhs_size
    Which is: 5
Google Test trace:
/cudf/cpp/tests/quantiles/percentile_approx_test.cpp:221:  <--  line of failure

This fails intermittently but is a bit more reliable running both the Simple and Grouped tests as shown above.
The test fails checking the result of the cudf::percentile_approx() API call with a derived expected result using various calls. The always returns the same result. The expected result is what changes and causes the intermittent failure. Building the expected values is quite involved in percentile_approx_test.cpp. One step involves using cudf::groupby::groupby(keys) and calling get_groups(values) on the input values. This returns a different result and likely contributes to failure.

The intermittentness hints at a possible timing issue and using compute-sanitizer showed no errors except when using initcheck.

======== COMPUTE-SANITIZER
Note: Google Test filter = PercentileApproxInputTypesTest/0.Grouped
[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from PercentileApproxInputTypesTest/0, where TypeParam = signed char
[ RUN      ] PercentileApproxInputTypesTest/0.Grouped
========= Host API memory access error at host access to 0x74ccc0f2a200 of size 600,000 bytes
=========     Uninitialized access at 0x74ccc0fb6cc0 on access by cudaMemcpy source
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: cuMemcpyAsync [0x3dc02f] in libcuda.so.1
=========         Host Frame:  [0x48178] in libcudart.so.12
=========         Host Frame:  [0x17057] in libcudart.so.12
=========         Host Frame: cudaMemcpyAsync [0x768a0] in libcudart.so.12
=========         Host Frame: cudaMemcpyAsync_ptsz in identify_stream_usage.cpp:274 [0x6b8c] in libcudf_identify_stream_usage_mode_cudf.so
=========         Host Frame: rmm::device_buffer::copy_async(void const*, unsigned long) in device_buffer.hpp:474 [0xcd726b] in libcudf.so
=========         Host Frame: rmm::device_buffer::device_buffer(void const*, unsigned long, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in device_buffer.hpp:143 [0xcd6cfa] in libcudf.so
=========         Host Frame: std::unique_ptr<cudf::column, std::default_delete<cudf::column> > cudf::(anonymous namespace)::create_column_from_view::operator()<int, (void*)0>() [0xce7769] in libcudf.so
=========         Host Frame: cudf::column::column(cudf::column_view, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in column.cu:256 [0xce549e] in libcudf.so
=========         Host Frame: std::__detail::_MakeUniq<cudf::column>::__single_object std::make_unique<cudf::column, cudf::column_view const&, rmm::cuda_stream_view&, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>&>(cudf::column_view const&, rmm::cuda_stream_view&, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>&) in unique_ptr.h:1070 [0xcecfdd] in libcudf.so
=========         Host Frame: std::unique_ptr<cudf::column, std::default_delete<cudf::column> > cudf::detail::column_scatterer_impl<int, void>::operator()<thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default> >(cudf::column_view const&, thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default>, thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default>, cudf::column_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) const in scatter.cuh:152 [0x10e6585] in libcudf.so
=========         Host Frame: std::unique_ptr<cudf::column, std::default_delete<cudf::column> > cudf::detail::column_scatterer::operator()<int, thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default> >(cudf::column_view const&, thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default>, thrust::transform_iterator<cudf::detail::index_converter<int>, cudf::detail::input_indexalator, thrust::use_default, thrust::use_default>, cudf::column_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) const in scatter.cuh:270 [0x10da2ac] in libcudf.so
=========         Host Frame: auto cudf::detail::scatter<cudf::detail::input_indexalator>(cudf::table_view const&, cudf::detail::input_indexalator, cudf::detail::input_indexalator, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>)::{lambda(auto:1 const&, auto:2 const&)#1}::operator()<cudf::column_view, {lambda(auto:1 const&, auto:2 const&)#1}::operator()>(cudf::column_view const&, {lambda(auto:1 const&, auto:2 const&)#1}::operator() const&) const in scatter.cuh:418 [0x10d5484] in libcudf.so
=========         Host Frame: __gnu_cxx::__normal_iterator<std::unique_ptr<cudf::column, std::default_delete<cudf::column> >*, std::vector<std::unique_ptr<cudf::column, std::default_delete<cudf::column> >, std::allocator<std::unique_ptr<cudf::column, std::default_delete<cudf::column> > > > > std::transform<__gnu_cxx::__normal_iterator<cudf::column_view const*, std::vector<cudf::column_view, std::allocator<cudf::column_view> > >, __gnu_cxx::__normal_iterator<cudf::column_view const*, std::vector<cudf::column_view, std::allocator<cudf::column_view> > >, __gnu_cxx::__normal_iterator<std::unique_ptr<cudf::column, std::default_delete<cudf::column> >*, std::vector<std::unique_ptr<cudf::column, std::default_delete<cudf::column> >, std::allocator<std::unique_ptr<cudf::column, std::default_delete<cudf::column> > > > >, cudf::detail::scatter<cudf::detail::input_indexalator>(cudf::table_view const&, cudf::detail::input_indexalator, cudf::detail::input_indexalator, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>)::{lambda(auto:1 const&, auto:2 const&)#1}>(__gnu_cxx::__normal_iterator<cudf::column_view const*, std::vector<cudf::column_view, std::allocator<cudf::column_view> > >, __gnu_cxx::__normal_iterator<cudf::column_view const*, std::vector<cudf::column_view, std::allocator<cudf::column_view> > >, __gnu_cxx::__normal_iterator<cudf::column_view const*, std::vector<cudf::column_view, std::allocator<cudf::column_view> > >, cudf::detail::scatter<cudf::detail::input_indexalator>(cudf::table_view const&, cudf::detail::input_indexalator, cudf::detail::input_indexalator, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>)::{lambda(auto:1 const&, auto:2 const&)#1}, cudf::detail::scatter<cudf::detail::input_indexalator>(cudf::table_view const&, cudf::detail::input_indexalator, cudf::detail::input_indexalator, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>)::{lambda(auto:1 const&, auto:2 const&)#1}) in stl_algo.h:4349 [0x10d6dc0] in libcudf.so
=========         Host Frame: std::unique_ptr<cudf::table, std::default_delete<cudf::table> > cudf::detail::scatter<cudf::detail::input_indexalator>(cudf::table_view const&, cudf::detail::input_indexalator, cudf::detail::input_indexalator, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in scatter.cuh:413 [0x10d1b75] in libcudf.so
=========         Host Frame: cudf::detail::scatter(cudf::table_view const&, cudf::column_view const&, cudf::table_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in scatter.cu:330 [0x10c3cca] in libcudf.so
=========         Host Frame: cudf::groupby::detail::sort::sort_groupby_helper::unsorted_keys_labels(rmm::cuda_stream_view) in sort_helper.cu:221 [0x195b2db] in libcudf.so
=========         Host Frame: cudf::groupby::detail::sort::sort_groupby_helper::sorted_values(cudf::column_view const&, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in sort_helper.cu:273 [0x195b909] in libcudf.so
=========         Host Frame: cudf::groupby::detail::store_result_functor::get_sorted_values() in functors.hpp:94 [0x156e4f7] in libcudf.so
=========         Host Frame: void cudf::groupby::detail::aggregate_result_functor::operator()<(cudf::aggregation::Kind)34>(cudf::aggregation const&) in aggregate.cpp:745 [0x15747ac] in libcudf.so
=========         Host Frame: decltype(auto) cudf::detail::aggregation_dispatcher<cudf::groupby::detail::aggregate_result_functor&, cudf::groupby_aggregation&>(cudf::aggregation::Kind, cudf::groupby::detail::aggregate_result_functor&, cudf::groupby_aggregation&) in aggregation.hpp:1618 [0x1575d56] in libcudf.so
=========         Host Frame: cudf::groupby::groupby::sort_aggregate(cudf::host_span<cudf::groupby::aggregation_request const, 18446744073709551615ul>, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in aggregate.cpp:857 [0x1575fd9] in libcudf.so
=========         Host Frame: cudf::groupby::groupby::dispatch_aggregation(cudf::host_span<cudf::groupby::aggregation_request const, 18446744073709551615ul>, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in groupby.cu:80 [0x150a466] in libcudf.so
=========         Host Frame: cudf::groupby::groupby::aggregate(cudf::host_span<cudf::groupby::aggregation_request const, 18446744073709551615ul>, rmm::cuda_stream_view, cuda::mr::__4::basic_resource_ref<(cuda::mr::__4::_AllocType)1, cuda::mr::__4::device_accessible>) in groupby.cu:228 [0x150a73c] in libcudf.so
=========         Host Frame: (anonymous namespace)::percentile_approx_test(cudf::column_view const&, cudf::column_view const&, int, std::vector<double, std::allocator<double> > const&, int)::{lambda(cudf::column_view const&, int)#1}::operator()(cudf::column_view const&, int) const in percentile_approx_test.cpp:165 [0x14efc0] in QUANTILES_TEST

The uninit'd section appears to be in the middle of a column which seems like it has not finished populating (?).
Adding some strategic cudaDeviceSynchonize() calls did not help so I'm not sure if the initcheck is accurate.

Also, there is the question of why this only occurs with PTDS and when stream-checking is enabled.
Maybe it is possible we somehow introduced a 2nd stream and so certain functions are being called on different streams causing the mismatch.

[vyasr]

@davidwendt do you still have the PTDS build handy to debug with? If so, inside stream_is_invalid here could you print out the value of the stream (cast to an integer probably) and see if everything is running on the same stream value? It would also be worth printing out the value of cudf::get_default_stream().value() to ensure that it is always the same.

Another possibility: I wonder if the set of functions that we patch is out of date. Maybe CCCL or one of our other dependencies is using a newer function somewhere under the hood and that is not being checked and it is somehow being passed different stream values.

A third possibility: maybe there are some symbols that are actually different for PTDS builds? I thought that was only true of the driver API, not the runtime API, but perhaps I was wrong.

[JigaoLuo]

Quick update: The QUANTILES_TEST shows increased runtime on branch 25.08. I tested twice today:

$ ctest
Test project /home/jluo/cudf-test/cpp/build
        Start   1: generate_resource_spec
  1/108 Test   #1: generate_resource_spec ...........   Passed    0.16 sec
        Start   2: COLUMN_TEST
  2/108 Test   #2: COLUMN_TEST ......................   Passed    8.59 sec
        Start   3: SCALAR_TEST
  3/108 Test   #3: SCALAR_TEST ......................   Passed    8.46 sec
        Start   4: TIMESTAMPS_TEST
  4/108 Test   #4: TIMESTAMPS_TEST ..................   Passed    8.45 sec
        Start   5: GROUPBY_TEST
  5/108 Test   #5: GROUPBY_TEST .....................   Passed   16.87 sec
        Start   6: JOIN_TEST
  6/108 Test   #6: JOIN_TEST ........................   Passed   12.08 sec
        Start   7: IS_SORTED_TEST
  7/108 Test   #7: IS_SORTED_TEST ...................   Passed    8.57 sec
        Start   8: DATETIME_OPS_TEST
  8/108 Test   #8: DATETIME_OPS_TEST ................   Passed    8.53 sec
        Start   9: HASHING_TEST
  9/108 Test   #9: HASHING_TEST .....................   Passed    8.61 sec
        Start  10: PARTITIONING_TEST
 10/108 Test  #10: PARTITIONING_TEST ................   Passed   11.67 sec
        Start  11: QUANTILES_TEST
 11/108 Test  #11: QUANTILES_TEST ...................***Timeout 1500.06 sec
....


$ ctest -R "QUANTILES_TEST"
Test project /home/jluo/cudf-dev/cpp/build
    Start  1: generate_resource_spec
1/2 Test  #1: generate_resource_spec ...........   Passed    0.16 sec
    Start 11: QUANTILES_TEST


2/2 Test #11: QUANTILES_TEST ...................***Failed  1373.88 sec

50% tests passed, 1 tests failed out of 2

Total Test time (real) = 1374.19 sec

The following tests FAILED:
	11 - QUANTILES_TEST (Failed)
Errors while running CTest

Standalone Google tests run with normal runtime and produce all correct results in tests.

[davidwendt]

The DATA_CHUNK_SOURCE_TEST appears to be fixed by #20492
The QUANTILES_TEST is still failing similar to #20504