Skip to content

Unified equal_to_value functions #8114

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

Merged
davebayer merged 6 commits into from
Mar 20, 2026
+29 −124
Merged

Unified equal_to_value functions #8114

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
906a059
removed duplicate equal_to_val implementation and unified into equal_…
dhruvilmehta Mar 19, 2026
3d5fbdf
fix formatting
dhruvilmehta Mar 19, 2026
50a272b
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Mar 19, 2026
2e5d7fc
Apply suggestions from code review
oleksandr-pavlyk Mar 19, 2026
c2c21b2
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Mar 19, 2026
b51dccd
Use public header
miscco Mar 20, 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
19 changes: 3 additions & 16 deletions libcudacxx/benchmarks/bench/all_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,27 +10,13 @@

#include <thrust/device_vector.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/std/__pstl_algorithm>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(cuda::std::all_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
do_not_optimize(
cuda::std::all_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

Expand Down
19 changes: 3 additions & 16 deletions libcudacxx/benchmarks/bench/any_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,27 +10,13 @@

#include <thrust/device_vector.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/std/__pstl_algorithm>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(cuda::std::any_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
do_not_optimize(
cuda::std::any_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

Expand Down
19 changes: 3 additions & 16 deletions libcudacxx/benchmarks/bench/none_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,27 +10,13 @@

#include <thrust/device_vector.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/std/__pstl_algorithm>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(cuda::std::none_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
do_not_optimize(
cuda::std::none_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

Expand Down
24 changes: 5 additions & 19 deletions thrust/benchmarks/bench/all_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,26 +5,12 @@
#include <thrust/fill.h>
#include <thrust/logical.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

caching_allocator_t alloc{};

state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
[&](nvbench::launch& launch) {
do_not_optimize(thrust::all_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
});
state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(thrust::all_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))
Expand Down
24 changes: 5 additions & 19 deletions thrust/benchmarks/bench/any_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,26 +5,12 @@
#include <thrust/fill.h>
#include <thrust/logical.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

caching_allocator_t alloc{};

state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
[&](nvbench::launch& launch) {
do_not_optimize(thrust::any_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
});
state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(thrust::any_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))
Expand Down
23 changes: 4 additions & 19 deletions thrust/benchmarks/bench/find/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,21 +16,6 @@

#include "nvbench_helper.cuh"

struct equal_to_val
{
size_t val_;

constexpr equal_to_val(const size_t val) noexcept
: val_(val)
{}

template <class T>
__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -49,10 +34,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

caching_allocator_t alloc{};

state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
[&](nvbench::launch& launch) {
do_not_optimize(thrust::find(policy(alloc, launch), dinput.begin(), dinput.end(), val));
});
state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(thrust::find(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))
Expand Down
24 changes: 5 additions & 19 deletions thrust/benchmarks/bench/none_of/basic.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,26 +5,12 @@
#include <thrust/fill.h>
#include <thrust/logical.h>

#include <cuda/__functional/equal_to_value.h>
#include <cuda/memory_pool>
#include <cuda/stream>

#include "nvbench_helper.cuh"

template <class T>
struct equal_to_val
{
T val_;

constexpr equal_to_val(const T& val) noexcept
: val_(val)
{}

__device__ constexpr bool operator()(const T& val) const noexcept
{
return val == val_;
}
};

template <typename T>
static void basic(nvbench::state& state, nvbench::type_list<T>)
{
Expand All @@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)

caching_allocator_t alloc{};

state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
[&](nvbench::launch& launch) {
do_not_optimize(thrust::none_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
});
state.exec(
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
do_not_optimize(thrust::none_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
});
}

NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))
Expand Down
Loading