Skip to content

[libcu++] Use cccl runtime in thrust tests pt 4 #9778

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

Open
pciolkosz wants to merge 3 commits into
base: main
Choose a base branch
from
+263 −212
Open

[libcu++] Use cccl runtime in thrust tests pt 4 #9778

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
0bbbfd8
Use CCCL Runtime in Thrust count and equal tests
pciolkosz Jun 30, 2026
12f0b11
Use CCCL Runtime in Thrust generate tests
pciolkosz Jul 9, 2026
db2427c
Hoist Runtime random sample generator
pciolkosz Jul 9, 2026
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
39 changes: 39 additions & 0 deletions thrust/testing/cuda/cccl_runtime_test_helper.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,8 +5,10 @@

#include <cuda/__cccl_config>
#include <cuda/algorithm>
#include <cuda/buffer>
#include <cuda/devices>
#include <cuda/launch>
#include <cuda/memory_resource>
#include <cuda/std/cstddef>
#include <cuda/std/initializer_list>
#include <cuda/std/span>
Expand Down Expand Up @@ -42,6 +44,30 @@ _CCCL_DEVICE_API inline void assert_device(bool condition, const char* expressio
}
}

template <typename T>
[[nodiscard]] _CCCL_HOST_API inline auto make_host_buffer(cuda::stream_ref stream, cuda::std::size_t size)
{
using resource_t = cuda::mr::synchronous_resource_adapter<cuda::mr::legacy_pinned_memory_resource>;
resource_t resource{cuda::mr::legacy_pinned_memory_resource{cuda::device_ref{0}}};
return cuda::make_buffer<T>(stream, resource, size, cuda::no_init);
Comment on lines +47 to +52

@coderabbitai coderabbitai Bot Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Inspect the helper and related test call sites.
git ls-files thrust/testing/cuda/cccl_runtime_test_helper.cuh thrust/testing/cuda | sed 's#^`#FILE` #'
echo
wc -l thrust/testing/cuda/cccl_runtime_test_helper.cuh
echo
cat -n thrust/testing/cuda/cccl_runtime_test_helper.cuh | sed -n '1,220p'
echo
echo "---- call sites ----"
rg -n "make_host_buffer|random_samples_buffer|current_test_device|legacy_pinned_memory_resource|stream_ref" thrust/testing/cuda -S
echo
echo "---- docs / declarations ----"
rg -n "legacy_pinned_memory_resource|synchronous_resource_adapter|make_buffer<|no_init" libcudacxx thrust -S

Repository: NVIDIA/cccl

Length of output: 50367

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Focus on the helper docs and how these tests bind streams/devices.
echo "---- legacy_pinned_memory_resource docs ----"
sed -n '1,180p' libcudacxx/include/cuda/__memory_resource/legacy_pinned_memory_resource.h
echo
echo "---- synchronous_resource_adapter docs ----"
sed -n '1,220p' libcudacxx/include/cuda/__memory_resource/synchronous_resource_adapter.h
echo
echo "---- stream creation / device binding in thrust testing cuda ----"
rg -n "make_stream|stream_ref|current_test_device\(\)|device_ref\{|set_device|device\)" thrust/testing/cuda/*.cu thrust/testing/cuda/*.cuh -S

Repository: NVIDIA/cccl

Length of output: 29273

🏁 Script executed:

#!/bin/bash
set -euo pipefail

echo "---- specific callers of random_samples_buffer / make_host_buffer ----"
sed -n '1,140p' thrust/testing/cuda/count.cu
echo
sed -n '1,180p' thrust/testing/cuda/equal.cu
echo
sed -n '1,180p' thrust/testing/cuda/generate.cu
echo
sed -n '1,140p' thrust/testing/cuda/tabulate.cu
echo
sed -n '1,160p' thrust/testing/cuda/set_union.cu

Repository: NVIDIA/cccl

Length of output: 16289

🏁 Script executed:

#!/bin/bash
set -euo pipefail

sed -n '1,120p' libcudacxx/test/libcudacxx/cuda/ccclrt/algorithm/common.cuh
echo
rg -n "make_host_buffer\\(|random_samples_buffer\\(" thrust/testing/cuda libcudacxx/test -S

Repository: NVIDIA/cccl

Length of output: 3535

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Inspect the stream constructor that takes a device.
sed -n '1,120p' libcudacxx/include/cuda/__stream/stream.h

Repository: NVIDIA/cccl

Length of output: 4192

🏁 Script executed:

#!/bin/bash
set -euo pipefail

sed -n '1,120p' libcudacxx/test/libcudacxx/cuda/ccclrt/algorithm/common.cuh

Repository: NVIDIA/cccl

Length of output: 3223

🏁 Script executed:

#!/bin/bash
set -euo pipefail

cat -n libcudacxx/test/libcudacxx/cuda/ccclrt/algorithm/common.cuh | sed -n '1,120p'
echo
cat -n libcudacxx/include/cuda/__stream/stream.h | sed -n '1,120p'

Repository: NVIDIA/cccl

Length of output: 8730

important: thrust/testing/cuda/cccl_runtime_test_helper.cuh:48-57 thread cuda::device_ref through make_host_buffer() and random_samples_buffer() instead of hard-coding cuda::device_ref{0}. These helpers are used with streams and device buffers bound to current_test_device(), so non-zero-device runs can tie pinned host memory to the wrong device.

Source: MCP tools

}

template <typename T, typename RandomT = T>
[[nodiscard]] _CCCL_HOST_API inline auto
random_samples_buffer(cuda::stream_ref stream, cuda::std::size_t size, cuda::std::size_t first = 0)
{
auto result = make_host_buffer<T>(stream, size);
stream.sync();

const auto generator = unittest::generate_random_sample<RandomT>{};
for (cuda::std::size_t i = 0; i < size; ++i)
{
result[i] = static_cast<T>(generator(static_cast<unsigned int>(first + i)));
Comment on lines +59 to +65

@coderabbitai coderabbitai Bot Jul 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🚀 Performance & Scalability | 🟡 Minor | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Inspect the target file and locate the helper types it uses.
git ls-files thrust/testing/cuda/cccl_runtime_test_helper.cuh
printf '\n--- file outline ---\n'
ast-grep outline thrust/testing/cuda/cccl_runtime_test_helper.cuh --view expanded || true
printf '\n--- relevant lines ---\n'
cat -n thrust/testing/cuda/cccl_runtime_test_helper.cuh | sed -n '1,220p'

printf '\n--- search for make_host_buffer and buffer accessors ---\n'
rg -n "make_host_buffer|get_unsynchronized|unsynchronized|operator\\[" thrust/testing -g'*.cuh' -g'*.cu' -g'*.cpp' -g'*.h' || true

Repository: NVIDIA/cccl

Length of output: 5418

🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Focus on the helper implementation and any nearby buffer type definitions/usages.
cat -n thrust/testing/cuda/cccl_runtime_test_helper.cuh | sed -n '1,180p'

printf '\n--- search exact accessor names in repo ---\n'
rg -n "get_unsynchronized|unsynchronized" thrust libcudacxx cub c python -g'*.h' -g'*.hpp' -g'*.hh' -g'*.cuh' -g'*.cu' -g'*.cpp' || true

Repository: NVIDIA/cccl

Length of output: 9785

🌐 Web query:

CCCL libcudacxx runtime buffer get_unsynchronized operator[] synchronized accessor stream.sync

💡 Result:

In the NVIDIA CUDA Core Compute Libraries (CCCL), specifically within libcudacxx's runtime components, the buffer API provides typed, stream-ordered storage [1][2]. Key components regarding your query: get_unsynchronized The get_unsynchronized(const size_type n) method provides access to the n-th element of the buffer without performing any synchronization with the underlying CUDA stream [3]. Because it is "unsynchronized," the user is responsible for ensuring that the data is ready for access (e.g., via previous stream operations) before calling this method [3]. It is typically used when the caller has already accounted for necessary ordering or synchronization dependencies [3]. operator[] In many container-like abstractions within libcudacxx, operator[] is often mapped to or implemented alongside get_unsynchronized (or similar accessors) to provide convenient element access [3]. Users should consult the specific buffer class implementation in their version of CCCL, as this interface has evolved from earlier experimental versions (previously cudax::buffer) to the current cuda::buffer in libcudacxx [3]. synchronized accessor (stream.sync) The cuda::stream_ref (or cuda::stream) provides a sync() method, which performs a blocking synchronization of the wrapped CUDA stream, ensuring that all previously submitted work is complete [4][5]. When working with buffers, stream synchronization is managed explicitly [6]. For example: - The buffer itself does not automatically synchronize on element access [3]. - To ensure safe access to data processed by the GPU, one must call stream.sync() on the stream associated with that buffer's operations before accessing the data [4][3]. - Recent changes in the CCCL library have emphasized removing implicit synchronization from fundamental primitives (like set_stream) to give users better control over performance and synchronization overhead [6]. If you are using memory resources that are not natively stream-ordered, cuda::mr::synchronous_resource_adapter can be used to adapt them; this adapter will perform the necessary stream synchronization internally during allocation or deallocation if the underlying resource lacks stream-ordered support [7][8].

Citations:

important: After stream.sync(), use result.get_unsynchronized(i) in the loop; result[i] takes the synchronized path and adds an unnecessary sync per element.

Source: MCP tools

}

return result;
}

template <typename Buffer>
_CCCL_HOST_API inline void
assert_equal(cuda::stream_ref stream, Buffer& buffer, cuda::std::initializer_list<int> expected)
Expand All @@ -57,6 +83,19 @@ assert_equal(cuda::stream_ref stream, Buffer& buffer, cuda::std::initializer_lis
ASSERT_EQUAL(expected.begin()[i], actual[i]);
}
}

template <typename Buffer>
_CCCL_HOST_API inline void assert_filled(cuda::stream_ref stream, Buffer& buffer, typename Buffer::value_type expected)
{
std::vector<typename Buffer::value_type> actual(buffer.size());
cuda::copy_bytes(stream, buffer, actual);
stream.sync();

for (cuda::std::size_t i = 0; i < actual.size(); ++i)
{
ASSERT_EQUAL(expected, actual[i]);
}
}
} // namespace test_runtime

#define TEST_ASSERT_DEVICE(condition) ::test_runtime::assert_device((condition), #condition, __FILE__, __LINE__)
80 changes: 46 additions & 34 deletions thrust/testing/cuda/count.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,30 +1,37 @@
#include <thrust/count.h>
#include <thrust/execution_policy.h>

#include <cuda/buffer>
#include <cuda/cccl_runtime_test_helper.cuh>
#include <cuda/launch>
#include <cuda/stream>

#include <unittest/unittest.h>

#ifdef THRUST_TEST_DEVICE_SIDE
template <typename ExecutionPolicy, typename Iterator, typename T, typename Iterator2>
__global__ void count_kernel(ExecutionPolicy exec, Iterator first, Iterator last, T value, Iterator2 result)
struct count_kernel
{
*result = thrust::count(exec, first, last, value);
}
template <typename ExecutionPolicy, typename Input, typename T, typename Expected>
__device__ void operator()(ExecutionPolicy exec, Input data, T value, Expected expected) const
{
const auto result = thrust::count(exec, data.begin(), data.end(), value);
TEST_ASSERT_DEVICE(result == expected);
}
};

template <typename T, typename ExecutionPolicy>
void TestCountDevice(ExecutionPolicy exec, const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
thrust::device_vector<T> d_data = h_data;

thrust::device_vector<size_t> d_result(1);
const auto device = test_runtime::current_test_device();
cuda::stream stream{device};

size_t h_result = thrust::count(h_data.begin(), h_data.end(), T(5));
auto h_data = test_runtime::random_samples_buffer<T>(stream, n);
auto d_data = cuda::make_device_buffer<T>(stream, device, h_data);

count_kernel<<<1, 1>>>(exec, d_data.begin(), d_data.end(), T(5), d_result.begin());
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
const auto expected = thrust::count(h_data.begin(), h_data.end(), T{5});

ASSERT_EQUAL(h_result, d_result[0]);
cuda::launch(stream, test_runtime::single_thread_config(), count_kernel{}, exec, d_data, T{5}, expected);
stream.sync();
}

template <typename T>
Expand All @@ -41,12 +48,6 @@ void TestCountDeviceDevice(const size_t n)
}
DECLARE_VARIABLE_UNITTEST(TestCountDeviceDevice);

template <typename ExecutionPolicy, typename Iterator, typename Predicate, typename Iterator2>
__global__ void count_if_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Predicate pred, Iterator2 result)
{
*result = thrust::count_if(exec, first, last, pred);
}

template <typename T>
struct greater_than_five
{
Expand All @@ -56,20 +57,30 @@ struct greater_than_five
}
};

struct count_if_kernel
{
template <typename ExecutionPolicy, typename Input, typename Predicate, typename Expected>
__device__ void operator()(ExecutionPolicy exec, Input data, Predicate pred, Expected expected) const
{
const auto result = thrust::count_if(exec, data.begin(), data.end(), pred);
TEST_ASSERT_DEVICE(result == expected);
}
};

template <typename T, typename ExecutionPolicy>
void TestCountIfDevice(ExecutionPolicy exec, const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
thrust::device_vector<T> d_data = h_data;
const auto device = test_runtime::current_test_device();
cuda::stream stream{device};

thrust::device_vector<size_t> d_result(1);
auto h_data = test_runtime::random_samples_buffer<T>(stream, n);
auto d_data = cuda::make_device_buffer<T>(stream, device, h_data);

size_t h_result = thrust::count_if(h_data.begin(), h_data.end(), greater_than_five<T>());
count_if_kernel<<<1, 1>>>(exec, d_data.begin(), d_data.end(), greater_than_five<T>(), d_result.begin());
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
const auto expected = thrust::count_if(h_data.begin(), h_data.end(), greater_than_five<T>{});

ASSERT_EQUAL(h_result, d_result[0]);
cuda::launch(
stream, test_runtime::single_thread_config(), count_if_kernel{}, exec, d_data, greater_than_five<T>{}, expected);
stream.sync();
}

template <typename T>
Expand All @@ -89,15 +100,16 @@ DECLARE_VARIABLE_UNITTEST(TestCountIfDeviceDevice);

void TestCountCudaStreams()
{
thrust::device_vector<int> data{1, 1, 0, 0, 1};
const auto device = test_runtime::current_test_device();
cuda::stream stream{device};

cudaStream_t s;
cudaStreamCreate(&s);
auto data = cuda::make_device_buffer<int>(stream, device, {1, 1, 0, 0, 1});
auto policy = thrust::cuda::par.on(stream.get());

ASSERT_EQUAL(thrust::count(thrust::cuda::par.on(s), data.begin(), data.end(), 0), 2);
ASSERT_EQUAL(thrust::count(thrust::cuda::par.on(s), data.begin(), data.end(), 1), 3);
ASSERT_EQUAL(thrust::count(thrust::cuda::par.on(s), data.begin(), data.end(), 2), 0);
ASSERT_EQUAL(thrust::count(policy, data.begin(), data.end(), 0), 2);
ASSERT_EQUAL(thrust::count(policy, data.begin(), data.end(), 1), 3);
ASSERT_EQUAL(thrust::count(policy, data.begin(), data.end(), 2), 0);

cudaStreamDestroy(s);
stream.sync();
}
DECLARE_UNITTEST(TestCountCudaStreams);
147 changes: 80 additions & 67 deletions thrust/testing/cuda/equal.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,79 +2,93 @@
#include <thrust/execution_policy.h>
#include <thrust/functional.h>

#include <cuda/buffer>
#include <cuda/cccl_runtime_test_helper.cuh>
#include <cuda/launch>
#include <cuda/std/cstddef>
#include <cuda/stream>

#include <unittest/unittest.h>

#ifdef THRUST_TEST_DEVICE_SIDE
template <typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3>
__global__ void equal_kernel(ExecutionPolicy exec, Iterator1 first1, Iterator1 last1, Iterator2 first2, Iterator3 result)
struct equal_kernel
{
*result = thrust::equal(exec, first1, last1, first2);
}
template <typename ExecutionPolicy, typename Input1, typename Input2>
__device__ void
operator()(ExecutionPolicy exec, Input1 data1, Input2 data2, cuda::std::size_t size, bool expected) const
{
const auto result = thrust::equal(exec, data1.begin(), data1.begin() + size, data2.begin());
TEST_ASSERT_DEVICE(result == expected);
}
};

template <typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename BinaryPredicate, typename Iterator3>
__global__ void equal_kernel(
ExecutionPolicy exec, Iterator1 first1, Iterator1 last1, Iterator2 first2, BinaryPredicate pred, Iterator3 result)
struct equal_pred_kernel
{
*result = thrust::equal(exec, first1, last1, first2, pred);
}
template <typename ExecutionPolicy, typename Input1, typename Input2, typename BinaryPredicate>
__device__ void operator()(
ExecutionPolicy exec, Input1 data1, Input2 data2, cuda::std::size_t size, BinaryPredicate pred, bool expected) const
{
const auto result = thrust::equal(exec, data1.begin(), data1.begin() + size, data2.begin(), pred);
TEST_ASSERT_DEVICE(result == expected);
}
};

template <typename T, typename ExecutionPolicy>
void TestEqualDevice(ExecutionPolicy exec, const size_t n)
{
thrust::device_vector<T> d_data1 = unittest::random_samples<T>(n);
thrust::device_vector<T> d_data2 = unittest::random_samples<T>(n);
thrust::device_vector<bool> d_result(1, false);
const auto device = test_runtime::current_test_device();
cuda::stream stream{device};

// empty ranges
equal_kernel<<<1, 1>>>(exec, d_data1.begin(), d_data1.begin(), d_data1.begin(), d_result.begin());
auto h_data1 = test_runtime::random_samples_buffer<T>(stream, n);
auto h_data2 = test_runtime::random_samples_buffer<T>(stream, n, n);

if (n > 0)
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
h_data1[0] = T{0};
h_data2[0] = T{1};
}

ASSERT_EQUAL(d_result[0], true);
auto d_data1 = cuda::make_device_buffer<T>(stream, device, h_data1);
auto d_data2 = cuda::make_device_buffer<T>(stream, device, h_data2);

// symmetric cases
equal_kernel<<<1, 1>>>(exec, d_data1.begin(), d_data1.end(), d_data1.begin(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}
// empty ranges
cuda::launch(stream, test_runtime::single_thread_config(), equal_kernel{}, exec, d_data1, d_data1, 0, true);
stream.sync();

ASSERT_EQUAL(d_result[0], true);
// symmetric cases
cuda::launch(stream, test_runtime::single_thread_config(), equal_kernel{}, exec, d_data1, d_data1, n, true);
stream.sync();

if (n > 0)
{
d_data1[0] = 0;
d_data2[0] = 1;

// different vectors
equal_kernel<<<1, 1>>>(exec, d_data1.begin(), d_data1.end(), d_data2.begin(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}

ASSERT_EQUAL(d_result[0], false);
cuda::launch(stream, test_runtime::single_thread_config(), equal_kernel{}, exec, d_data1, d_data2, n, false);
stream.sync();

// different predicates
equal_kernel<<<1, 1>>>(
exec, d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), ::cuda::std::less<T>(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}

ASSERT_EQUAL(d_result[0], true);

equal_kernel<<<1, 1>>>(
exec, d_data1.begin(), d_data1.begin() + 1, d_data2.begin(), ::cuda::std::greater<T>(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}

ASSERT_EQUAL(d_result[0], false);
cuda::launch(
stream,
test_runtime::single_thread_config(),
equal_pred_kernel{},
exec,
d_data1,
d_data2,
1,
::cuda::std::less<T>{},
true);
stream.sync();

cuda::launch(
stream,
test_runtime::single_thread_config(),
equal_pred_kernel{},
exec,
d_data1,
d_data2,
1,
::cuda::std::greater<T>{},
false);
stream.sync();
}
}

Expand All @@ -95,26 +109,25 @@ DECLARE_VARIABLE_UNITTEST(TestEqualDeviceDevice);

void TestEqualCudaStreams()
{
thrust::device_vector<int> v1 = {5, 2, 0, 0, 0};
thrust::device_vector<int> v2 = {5, 2, 0, 6, 1};
const auto device = test_runtime::current_test_device();
cuda::stream stream{device};

cudaStream_t s;
cudaStreamCreate(&s);
auto v1 = cuda::make_device_buffer<int>(stream, device, {5, 2, 0, 0, 0});
auto v2 = cuda::make_device_buffer<int>(stream, device, {5, 2, 0, 6, 1});
auto policy = thrust::cuda::par.on(stream.get());

ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.end(), v1.begin()), true);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.end(), v2.begin()), false);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v2.begin(), v2.end(), v2.begin()), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.end(), v1.begin()), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.end(), v2.begin()), false);
ASSERT_EQUAL(thrust::equal(policy, v2.begin(), v2.end(), v2.begin()), true);

ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.begin() + 0, v1.begin()), true);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.begin() + 1, v1.begin()), true);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.begin() + 3, v2.begin()), true);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.begin() + 4, v2.begin()), false);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.begin() + 0, v1.begin()), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.begin() + 1, v1.begin()), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.begin() + 3, v2.begin()), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.begin() + 4, v2.begin()), false);

ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.end(), v2.begin(), ::cuda::std::less_equal<int>()),
true);
ASSERT_EQUAL(thrust::equal(thrust::cuda::par.on(s), v1.begin(), v1.end(), v2.begin(), ::cuda::std::greater<int>()),
false);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.end(), v2.begin(), ::cuda::std::less_equal<int>{}), true);
ASSERT_EQUAL(thrust::equal(policy, v1.begin(), v1.end(), v2.begin(), ::cuda::std::greater<int>{}), false);

cudaStreamDestroy(s);
stream.sync();
}
DECLARE_UNITTEST(TestEqualCudaStreams);
Loading
Loading