From 2ac86e847d4cb9b79925910585fa6fea0b07862e Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Wed, 15 Nov 2023 23:08:16 +0000 Subject: [PATCH 1/8] extern c dlopen capers --- cpp/CMakeLists.txt | 4 ++ cpp/src/utilities/stream_pool.cpp | 44 ++++++++++++++++++- cpp/tests/utilities/default_stream.cpp | 6 +++ cpp/tests/utilities/identify_stream_usage.cpp | 9 ++++ 4 files changed, 62 insertions(+), 1 deletion(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index dc12564c656..98b8e159c03 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -719,6 +719,10 @@ target_compile_definitions( "$:${CUDF_CUDA_DEFINITIONS}>>" ) +#if(CUDF_BUILD_STREAMS_TEST_UTIL) +# target_compile_definitions(cudf PRIVATE STREAM_MODE_TESTING) +#endif() + # Disable Jitify log printing. See https://github.com/NVIDIA/jitify/issues/79 target_compile_definitions(cudf PRIVATE "JITIFY_PRINT_LOG=0") diff --git a/cpp/src/utilities/stream_pool.cpp b/cpp/src/utilities/stream_pool.cpp index b3b20889ef8..eac8764cd77 100644 --- a/cpp/src/utilities/stream_pool.cpp +++ b/cpp/src/utilities/stream_pool.cpp @@ -26,6 +26,7 @@ #include #include #include +#include namespace cudf::detail { @@ -157,13 +158,52 @@ class debug_cuda_stream_pool : public cuda_stream_pool { std::size_t get_stream_pool_size() const override { return 1UL; } }; +/** + * @brief Implementation of `cuda_stream_pool` that always returns the `cudf::test::get_default_stream()` + */ +class test_cuda_stream_pool : public cuda_stream_pool { + rmm::cuda_stream_view load_test_default_stream() { + std::cout << "Inside load_test_default_stream()\n"; + void *handle; + handle = dlopen("libcudftest_default_stream.so", RTLD_LAZY); + if(!handle) { + std::cerr << "cannot load libcudftest_default_stream " << dlerror() << std::endl; + exit(EXIT_FAILURE); + } + rmm::cuda_stream_view (*tds)(); + void *utds = dlsym(handle, "cudf_test_get_default_stream"); + if (utds == NULL) { + std::cerr << "Error accessing the symbol cudf_test_get_default_stream " << dlerror() << "\n"; + exit(EXIT_FAILURE); + } + tds = reinterpret_cast(utds); + return (*tds)(); + } + public: + rmm::cuda_stream_view get_stream() override { return load_test_default_stream(); } + rmm::cuda_stream_view get_stream(stream_id_type stream_id) override + { + return load_test_default_stream(); + } + + std::vector get_streams(std::size_t count) override + { + std::cout << "getting cudf test streams\n"; + return std::vector(count, load_test_default_stream()); + } + + std::size_t get_stream_pool_size() const override { return 1UL; } +}; + /** * @brief Initialize global stream pool. */ cuda_stream_pool* create_global_cuda_stream_pool() { if (getenv("LIBCUDF_USE_DEBUG_STREAM_POOL")) return new debug_cuda_stream_pool(); - + if (getenv("LIBCUDF_USE_TEST_STREAM_POOL")) { + return new test_cuda_stream_pool(); + } return new rmm_cuda_stream_pool(); } @@ -236,6 +276,8 @@ cuda_stream_pool& global_cuda_stream_pool() std::vector fork_streams(rmm::cuda_stream_view stream, std::size_t count) { auto const streams = global_cuda_stream_pool().get_streams(count); + for(size_t i = 0; i < streams.size(); i++) + std::cout << streams[i] << std::endl; auto const event = event_for_thread(); CUDF_CUDA_TRY(cudaEventRecord(event, stream)); std::for_each(streams.begin(), streams.end(), [&](auto& strm) { diff --git a/cpp/tests/utilities/default_stream.cpp b/cpp/tests/utilities/default_stream.cpp index 52752f78bb9..767594b0bbb 100644 --- a/cpp/tests/utilities/default_stream.cpp +++ b/cpp/tests/utilities/default_stream.cpp @@ -18,6 +18,12 @@ #include +extern "C" { + rmm::cuda_stream_view const cudf_test_get_default_stream() { + return cudf::get_default_stream(); + } +} + namespace cudf { namespace test { diff --git a/cpp/tests/utilities/identify_stream_usage.cpp b/cpp/tests/utilities/identify_stream_usage.cpp index ab2a85a0842..e1352637495 100644 --- a/cpp/tests/utilities/identify_stream_usage.cpp +++ b/cpp/tests/utilities/identify_stream_usage.cpp @@ -72,6 +72,15 @@ bool stream_is_invalid(cudaStream_t stream) { #ifdef STREAM_MODE_TESTING // In this mode the _only_ valid stream is the one returned by cudf::test::get_default_stream. + /* +#if defined(CUDF_USE_PER_THREAD_DEFAULT_STREAM) + rmm::cuda_stream_view default_stream_value{rmm::cuda_stream_per_thread}; +#else + rmm::cuda_stream_view default_stream_value{}; +#endif + return (stream == default_stream_value.value() || stream == cudaStreamDefault || + stream == cudaStreamLegacy || stream == cudaStreamPerThread); + */ return (stream != cudf::test::get_default_stream().value()); #else // We explicitly list the possibilities rather than using From 1247e6dc0ecfe8abda79cde03ea2f91c260cb8a2 Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Thu, 16 Nov 2023 23:39:15 +0000 Subject: [PATCH 2/8] moved test stream pool creation to preload library to override create_global_cuda_stream_pool --- .../cudf/detail/utilities/stream_pool.hpp | 56 ++++++++++ cpp/src/utilities/stream_pool.cpp | 101 ------------------ cpp/tests/CMakeLists.txt | 1 + cpp/tests/streams/pool_test.cu | 32 ++++++ cpp/tests/utilities/default_stream.cpp | 6 -- cpp/tests/utilities/identify_stream_usage.cpp | 38 +++++-- 6 files changed, 118 insertions(+), 116 deletions(-) create mode 100644 cpp/tests/streams/pool_test.cu diff --git a/cpp/include/cudf/detail/utilities/stream_pool.hpp b/cpp/include/cudf/detail/utilities/stream_pool.hpp index 95384a9d73e..12c4c58ba87 100644 --- a/cpp/include/cudf/detail/utilities/stream_pool.hpp +++ b/cpp/include/cudf/detail/utilities/stream_pool.hpp @@ -25,6 +25,62 @@ namespace cudf::detail { +class cuda_stream_pool { + public: + // matching type used in rmm::cuda_stream_pool::get_stream(stream_id) + using stream_id_type = std::size_t; + + virtual ~cuda_stream_pool() = default; + + /** + * @brief Get a `cuda_stream_view` of a stream in the pool. + * + * This function is thread safe with respect to other calls to the same function. + * + * @return Stream view. + */ + virtual rmm::cuda_stream_view get_stream() = 0; + + /** + * @brief Get a `cuda_stream_view` of the stream associated with `stream_id`. + * + * Equivalent values of `stream_id` return a `cuda_stream_view` to the same underlying stream. + * This function is thread safe with respect to other calls to the same function. + * + * @param stream_id Unique identifier for the desired stream + * @return Requested stream view. + */ + virtual rmm::cuda_stream_view get_stream(stream_id_type stream_id) = 0; + + /** + * @brief Get a set of `cuda_stream_view` objects from the pool. + * + * An attempt is made to ensure that the returned vector does not contain duplicate + * streams, but this cannot be guaranteed if `count` is greater than the value returned by + * `get_stream_pool_size()`. + * + * This function is thread safe with respect to other calls to the same function. + * + * @param count The number of stream views to return. + * @return Vector containing `count` stream views. + */ + virtual std::vector get_streams(std::size_t count) = 0; + + /** + * @brief Get the number of stream objects in the pool. + * + * This function is thread safe with respect to other calls to the same function. + * + * @return the number of stream objects in the pool + */ + virtual std::size_t get_stream_pool_size() const = 0; +}; + +/** + * @brief Initialize global stream pool. + */ +cuda_stream_pool* create_global_cuda_stream_pool(); + /** * @brief Acquire a set of `cuda_stream_view` objects and synchronize them to an event on another * stream. diff --git a/cpp/src/utilities/stream_pool.cpp b/cpp/src/utilities/stream_pool.cpp index eac8764cd77..121873ad44b 100644 --- a/cpp/src/utilities/stream_pool.cpp +++ b/cpp/src/utilities/stream_pool.cpp @@ -26,12 +26,9 @@ #include #include #include -#include namespace cudf::detail { -namespace { - // TODO: what is a good number here. what's the penalty for making it larger? // Dave Baranec rule of thumb was max_streams_needed * num_concurrent_threads, // where num_concurrent_threads was estimated to be 4. so using 32 will allow @@ -59,57 +56,6 @@ std::size_t constexpr STREAM_POOL_SIZE = 32; } while (0) #endif -class cuda_stream_pool { - public: - // matching type used in rmm::cuda_stream_pool::get_stream(stream_id) - using stream_id_type = std::size_t; - - virtual ~cuda_stream_pool() = default; - - /** - * @brief Get a `cuda_stream_view` of a stream in the pool. - * - * This function is thread safe with respect to other calls to the same function. - * - * @return Stream view. - */ - virtual rmm::cuda_stream_view get_stream() = 0; - - /** - * @brief Get a `cuda_stream_view` of the stream associated with `stream_id`. - * - * Equivalent values of `stream_id` return a `cuda_stream_view` to the same underlying stream. - * This function is thread safe with respect to other calls to the same function. - * - * @param stream_id Unique identifier for the desired stream - * @return Requested stream view. - */ - virtual rmm::cuda_stream_view get_stream(stream_id_type stream_id) = 0; - - /** - * @brief Get a set of `cuda_stream_view` objects from the pool. - * - * An attempt is made to ensure that the returned vector does not contain duplicate - * streams, but this cannot be guaranteed if `count` is greater than the value returned by - * `get_stream_pool_size()`. - * - * This function is thread safe with respect to other calls to the same function. - * - * @param count The number of stream views to return. - * @return Vector containing `count` stream views. - */ - virtual std::vector get_streams(std::size_t count) = 0; - - /** - * @brief Get the number of stream objects in the pool. - * - * This function is thread safe with respect to other calls to the same function. - * - * @return the number of stream objects in the pool - */ - virtual std::size_t get_stream_pool_size() const = 0; -}; - /** * @brief Implementation of `cuda_stream_pool` that wraps an `rmm::cuda_stram_pool`. */ @@ -158,52 +104,9 @@ class debug_cuda_stream_pool : public cuda_stream_pool { std::size_t get_stream_pool_size() const override { return 1UL; } }; -/** - * @brief Implementation of `cuda_stream_pool` that always returns the `cudf::test::get_default_stream()` - */ -class test_cuda_stream_pool : public cuda_stream_pool { - rmm::cuda_stream_view load_test_default_stream() { - std::cout << "Inside load_test_default_stream()\n"; - void *handle; - handle = dlopen("libcudftest_default_stream.so", RTLD_LAZY); - if(!handle) { - std::cerr << "cannot load libcudftest_default_stream " << dlerror() << std::endl; - exit(EXIT_FAILURE); - } - rmm::cuda_stream_view (*tds)(); - void *utds = dlsym(handle, "cudf_test_get_default_stream"); - if (utds == NULL) { - std::cerr << "Error accessing the symbol cudf_test_get_default_stream " << dlerror() << "\n"; - exit(EXIT_FAILURE); - } - tds = reinterpret_cast(utds); - return (*tds)(); - } - public: - rmm::cuda_stream_view get_stream() override { return load_test_default_stream(); } - rmm::cuda_stream_view get_stream(stream_id_type stream_id) override - { - return load_test_default_stream(); - } - - std::vector get_streams(std::size_t count) override - { - std::cout << "getting cudf test streams\n"; - return std::vector(count, load_test_default_stream()); - } - - std::size_t get_stream_pool_size() const override { return 1UL; } -}; - -/** - * @brief Initialize global stream pool. - */ cuda_stream_pool* create_global_cuda_stream_pool() { if (getenv("LIBCUDF_USE_DEBUG_STREAM_POOL")) return new debug_cuda_stream_pool(); - if (getenv("LIBCUDF_USE_TEST_STREAM_POOL")) { - return new test_cuda_stream_pool(); - } return new rmm_cuda_stream_pool(); } @@ -271,13 +174,9 @@ cuda_stream_pool& global_cuda_stream_pool() return *pools[device_id.value()]; } -} // anonymous namespace - std::vector fork_streams(rmm::cuda_stream_view stream, std::size_t count) { auto const streams = global_cuda_stream_pool().get_streams(count); - for(size_t i = 0; i < streams.size(); i++) - std::cout << streams[i] << std::endl; auto const event = event_for_thread(); CUDF_CUDA_TRY(cudaEventRecord(event, stream)); std::for_each(streams.begin(), streams.end(), [&](auto& strm) { diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 196a4f2d038..6efbed86c34 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -637,6 +637,7 @@ ConfigureTest(STREAM_INTEROP_TEST streams/interop_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_JSONIO_TEST streams/io/json_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_LISTS_TEST streams/lists_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_NULL_MASK_TEST streams/null_mask_test.cpp STREAM_MODE testing) +ConfigureTest(STREAM_POOL_TEST streams/pool_test.cu STREAM_MODE testing) ConfigureTest(STREAM_REPLACE_TEST streams/replace_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_SEARCH_TEST streams/search_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_SORTING_TEST streams/sorting_test.cpp STREAM_MODE testing) diff --git a/cpp/tests/streams/pool_test.cu b/cpp/tests/streams/pool_test.cu new file mode 100644 index 00000000000..65c8ff91b51 --- /dev/null +++ b/cpp/tests/streams/pool_test.cu @@ -0,0 +1,32 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include + +class PoolTest : public cudf::test::BaseFixture {}; + +__global__ void do_nothing_kernel() {} + +TEST_F(PoolTest, ForkStreams) +{ + auto streams = cudf::detail::fork_streams(cudf::test::get_default_stream(), 4); + do_nothing_kernel<<<1, 1, 0, streams[0].value()>>>(); + do_nothing_kernel<<<1, 1, 0, streams[1].value()>>>(); +} diff --git a/cpp/tests/utilities/default_stream.cpp b/cpp/tests/utilities/default_stream.cpp index 767594b0bbb..52752f78bb9 100644 --- a/cpp/tests/utilities/default_stream.cpp +++ b/cpp/tests/utilities/default_stream.cpp @@ -18,12 +18,6 @@ #include -extern "C" { - rmm::cuda_stream_view const cudf_test_get_default_stream() { - return cudf::get_default_stream(); - } -} - namespace cudf { namespace test { diff --git a/cpp/tests/utilities/identify_stream_usage.cpp b/cpp/tests/utilities/identify_stream_usage.cpp index e1352637495..7f672869a6f 100644 --- a/cpp/tests/utilities/identify_stream_usage.cpp +++ b/cpp/tests/utilities/identify_stream_usage.cpp @@ -15,6 +15,7 @@ */ #include +#include #include #include @@ -66,21 +67,40 @@ rmm::cuda_stream_view const get_default_stream() } // namespace test #endif +#ifdef STREAM_MODE_TESTING +namespace detail { + +/** + * @brief Implementation of `cuda_stream_pool` that always returns the + * `cudf::test::get_default_stream()` + */ +class test_cuda_stream_pool : public cuda_stream_pool { + public: + rmm::cuda_stream_view get_stream() override { return cudf::test::get_default_stream(); } + rmm::cuda_stream_view get_stream(stream_id_type stream_id) override + { + return cudf::test::get_default_stream(); + } + + std::vector get_streams(std::size_t count) override + { + return std::vector(count, cudf::test::get_default_stream()); + } + + std::size_t get_stream_pool_size() const override { return 1UL; } +}; + +cuda_stream_pool* create_global_cuda_stream_pool() { return new test_cuda_stream_pool(); } + +} // namespace detail +#endif + } // namespace cudf bool stream_is_invalid(cudaStream_t stream) { #ifdef STREAM_MODE_TESTING // In this mode the _only_ valid stream is the one returned by cudf::test::get_default_stream. - /* -#if defined(CUDF_USE_PER_THREAD_DEFAULT_STREAM) - rmm::cuda_stream_view default_stream_value{rmm::cuda_stream_per_thread}; -#else - rmm::cuda_stream_view default_stream_value{}; -#endif - return (stream == default_stream_value.value() || stream == cudaStreamDefault || - stream == cudaStreamLegacy || stream == cudaStreamPerThread); - */ return (stream != cudf::test::get_default_stream().value()); #else // We explicitly list the possibilities rather than using From 3fa0140a366dc47b2f4d7c95c256d2bd0158d5c9 Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Fri, 17 Nov 2023 00:24:17 +0000 Subject: [PATCH 3/8] Cmakelists cleanup --- cpp/CMakeLists.txt | 4 ---- 1 file changed, 4 deletions(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 98b8e159c03..dc12564c656 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -719,10 +719,6 @@ target_compile_definitions( "$:${CUDF_CUDA_DEFINITIONS}>>" ) -#if(CUDF_BUILD_STREAMS_TEST_UTIL) -# target_compile_definitions(cudf PRIVATE STREAM_MODE_TESTING) -#endif() - # Disable Jitify log printing. See https://github.com/NVIDIA/jitify/issues/79 target_compile_definitions(cudf PRIVATE "JITIFY_PRINT_LOG=0") From c612fd64ee5b59691925d360d3c8e64bb1719678 Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Tue, 12 Dec 2023 00:57:30 +0000 Subject: [PATCH 4/8] adding comments --- cpp/tests/streams/pool_test.cu | 10 +++++----- cpp/tests/utilities/identify_stream_usage.cpp | 14 +++++++++----- 2 files changed, 14 insertions(+), 10 deletions(-) diff --git a/cpp/tests/streams/pool_test.cu b/cpp/tests/streams/pool_test.cu index 65c8ff91b51..1d3f20b9b51 100644 --- a/cpp/tests/streams/pool_test.cu +++ b/cpp/tests/streams/pool_test.cu @@ -20,13 +20,13 @@ #include #include -class PoolTest : public cudf::test::BaseFixture {}; +class StreamPoolTest : public cudf::test::BaseFixture {}; __global__ void do_nothing_kernel() {} -TEST_F(PoolTest, ForkStreams) +TEST_F(StreamPoolTest, ForkStreams) { - auto streams = cudf::detail::fork_streams(cudf::test::get_default_stream(), 4); - do_nothing_kernel<<<1, 1, 0, streams[0].value()>>>(); - do_nothing_kernel<<<1, 1, 0, streams[1].value()>>>(); + auto streams = cudf::detail::fork_streams(cudf::test::get_default_stream(), 2); + do_nothing_kernel<<<1, 32, 0, streams[0].value()>>>(); + do_nothing_kernel<<<1, 32, 0, streams[1].value()>>>(); } diff --git a/cpp/tests/utilities/identify_stream_usage.cpp b/cpp/tests/utilities/identify_stream_usage.cpp index 7f672869a6f..52c454e906b 100644 --- a/cpp/tests/utilities/identify_stream_usage.cpp +++ b/cpp/tests/utilities/identify_stream_usage.cpp @@ -32,10 +32,14 @@ #include #include +// This file is compiled into a separate library that is dynamically loaded with LD_PRELOAD at runtime to libcudf +// to override some stream-related symbols in libcudf. The goal of such a library is to verify if the stream/stream pool +// is being correctly forwarded between API calls. +// // We control whether to override cudf::test::get_default_stream or -// cudf::get_default_stream with a compile-time flag. Thesee are the two valid -// options: -// 1. STREAM_MODE_TESTING=OFF: In this mode, cudf::get_default_stream will +// cudf::get_default_stream with a compile-time flag. The behaviour of tests +// depend on whether STREAM_MODE_TESTING is defined: +// 1. If STREAM_MODE_TESTING is not defined, cudf::get_default_stream will // return a custom stream and stream_is_invalid will return true if any CUDA // API is called using any of CUDA's default stream constants // (cudaStreamLegacy, cudaStreamDefault, or cudaStreamPerThread). This check @@ -45,7 +49,7 @@ // is not sufficient to guarantee a stream-ordered API because it will not // identify places in the code that use cudf::get_default_stream instead of // properly forwarding along a user-provided stream. -// 2. STREAM_MODE_TESTING=ON: In this mode, cudf::test::get_default_stream +// 2. If STREAM_MODE_TESTING compiler option is defined, cudf::test::get_default_stream // returns a custom stream and stream_is_invalid returns true if any CUDA // API is called using any stream other than cudf::test::get_default_stream. // This is a necessary and sufficient condition to ensure that libcudf is @@ -77,7 +81,7 @@ namespace detail { class test_cuda_stream_pool : public cuda_stream_pool { public: rmm::cuda_stream_view get_stream() override { return cudf::test::get_default_stream(); } - rmm::cuda_stream_view get_stream(stream_id_type stream_id) override + [[maybe_unused]] rmm::cuda_stream_view get_stream(stream_id_type stream_id) override { return cudf::test::get_default_stream(); } From 4e04b8ee10870428f6869872e231c8a4cb2cc94b Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Tue, 12 Dec 2023 01:12:57 +0000 Subject: [PATCH 5/8] forgot to run precommit checks --- cpp/tests/utilities/identify_stream_usage.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/tests/utilities/identify_stream_usage.cpp b/cpp/tests/utilities/identify_stream_usage.cpp index 52c454e906b..bdc338d2c92 100644 --- a/cpp/tests/utilities/identify_stream_usage.cpp +++ b/cpp/tests/utilities/identify_stream_usage.cpp @@ -32,12 +32,12 @@ #include #include -// This file is compiled into a separate library that is dynamically loaded with LD_PRELOAD at runtime to libcudf -// to override some stream-related symbols in libcudf. The goal of such a library is to verify if the stream/stream pool -// is being correctly forwarded between API calls. +// This file is compiled into a separate library that is dynamically loaded with LD_PRELOAD at +// runtime to libcudf to override some stream-related symbols in libcudf. The goal of such a library +// is to verify if the stream/stream pool is being correctly forwarded between API calls. // // We control whether to override cudf::test::get_default_stream or -// cudf::get_default_stream with a compile-time flag. The behaviour of tests +// cudf::get_default_stream with a compile-time flag. The behaviour of tests // depend on whether STREAM_MODE_TESTING is defined: // 1. If STREAM_MODE_TESTING is not defined, cudf::get_default_stream will // return a custom stream and stream_is_invalid will return true if any CUDA From a7204444a861f51e5578f72a5bb976edd3b44cf3 Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Tue, 12 Dec 2023 14:37:58 -0800 Subject: [PATCH 6/8] Update cpp/include/cudf/detail/utilities/stream_pool.hpp - clarifying comment Co-authored-by: Vyas Ramasubramani --- cpp/include/cudf/detail/utilities/stream_pool.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/utilities/stream_pool.hpp b/cpp/include/cudf/detail/utilities/stream_pool.hpp index 12c4c58ba87..19ef26a10cb 100644 --- a/cpp/include/cudf/detail/utilities/stream_pool.hpp +++ b/cpp/include/cudf/detail/utilities/stream_pool.hpp @@ -67,7 +67,7 @@ class cuda_stream_pool { virtual std::vector get_streams(std::size_t count) = 0; /** - * @brief Get the number of stream objects in the pool. + * @brief Get the number of unique stream objects in the pool. * * This function is thread safe with respect to other calls to the same function. * From bb54197fafaf56599a842915db79fd8b0622973a Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Wed, 13 Dec 2023 14:51:10 -0800 Subject: [PATCH 7/8] Update cpp/tests/streams/pool_test.cu - looping through streams in test Co-authored-by: Vukasin Milovanovic --- cpp/tests/streams/pool_test.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/tests/streams/pool_test.cu b/cpp/tests/streams/pool_test.cu index 1d3f20b9b51..46b669133bd 100644 --- a/cpp/tests/streams/pool_test.cu +++ b/cpp/tests/streams/pool_test.cu @@ -27,6 +27,7 @@ __global__ void do_nothing_kernel() {} TEST_F(StreamPoolTest, ForkStreams) { auto streams = cudf::detail::fork_streams(cudf::test::get_default_stream(), 2); - do_nothing_kernel<<<1, 32, 0, streams[0].value()>>>(); - do_nothing_kernel<<<1, 32, 0, streams[1].value()>>>(); + for (auto& stream : streams){ + do_nothing_kernel<<<1, 32, 0, stream.value()>>>(); + } } From c4e72d23d8a827b1bca979996fccf15887467bc2 Mon Sep 17 00:00:00 2001 From: Shruti Shivakumar Date: Thu, 14 Dec 2023 06:11:34 +0000 Subject: [PATCH 8/8] addressing PR reviews --- cpp/tests/streams/pool_test.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/tests/streams/pool_test.cu b/cpp/tests/streams/pool_test.cu index 46b669133bd..0f92e1c0c2b 100644 --- a/cpp/tests/streams/pool_test.cu +++ b/cpp/tests/streams/pool_test.cu @@ -14,12 +14,12 @@ * limitations under the License. */ -#include -#include - #include #include +#include +#include + class StreamPoolTest : public cudf::test::BaseFixture {}; __global__ void do_nothing_kernel() {} @@ -27,7 +27,7 @@ __global__ void do_nothing_kernel() {} TEST_F(StreamPoolTest, ForkStreams) { auto streams = cudf::detail::fork_streams(cudf::test::get_default_stream(), 2); - for (auto& stream : streams){ + for (auto& stream : streams) { do_nothing_kernel<<<1, 32, 0, stream.value()>>>(); } }