Skip to content

Commit e1329ca

Browse files
dhruvilmehtadhruvilmehta
authored
Unified equal_to_value functions (#8114)
1 parent 6c3033c commit e1329ca

File tree

7 files changed

+29
-124
lines changed

7 files changed

+29
-124
lines changed

‎libcudacxx/benchmarks/bench/all_of/basic.cu‎

Lines changed: 3 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -10,27 +10,13 @@
1010

1111
#include <thrust/device_vector.h>
1212

13+
#include <cuda/functional>
1314
#include <cuda/memory_pool>
1415
#include <cuda/std/__pstl_algorithm>
1516
#include <cuda/stream>
1617

1718
#include "nvbench_helper.cuh"
1819

19-
template <class T>
20-
struct equal_to_val
21-
{
22-
T val_;
23-
24-
constexpr equal_to_val(const T& val) noexcept
25-
: val_(val)
26-
{}
27-
28-
__device__ constexpr bool operator()(const T& val) const noexcept
29-
{
30-
return val == val_;
31-
}
32-
};
33-
3420
template <typename T>
3521
static void basic(nvbench::state& state, nvbench::type_list<T>)
3622
{
@@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
5137

5238
state.exec(
5339
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
54-
do_not_optimize(cuda::std::all_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
40+
do_not_optimize(
41+
cuda::std::all_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
5542
});
5643
}
5744

‎libcudacxx/benchmarks/bench/any_of/basic.cu‎

Lines changed: 3 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -10,27 +10,13 @@
1010

1111
#include <thrust/device_vector.h>
1212

13+
#include <cuda/functional>
1314
#include <cuda/memory_pool>
1415
#include <cuda/std/__pstl_algorithm>
1516
#include <cuda/stream>
1617

1718
#include "nvbench_helper.cuh"
1819

19-
template <class T>
20-
struct equal_to_val
21-
{
22-
T val_;
23-
24-
constexpr equal_to_val(const T& val) noexcept
25-
: val_(val)
26-
{}
27-
28-
__device__ constexpr bool operator()(const T& val) const noexcept
29-
{
30-
return val == val_;
31-
}
32-
};
33-
3420
template <typename T>
3521
static void basic(nvbench::state& state, nvbench::type_list<T>)
3622
{
@@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
5137

5238
state.exec(
5339
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
54-
do_not_optimize(cuda::std::any_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
40+
do_not_optimize(
41+
cuda::std::any_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
5542
});
5643
}
5744

‎libcudacxx/benchmarks/bench/none_of/basic.cu‎

Lines changed: 3 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -10,27 +10,13 @@
1010

1111
#include <thrust/device_vector.h>
1212

13+
#include <cuda/functional>
1314
#include <cuda/memory_pool>
1415
#include <cuda/std/__pstl_algorithm>
1516
#include <cuda/stream>
1617

1718
#include "nvbench_helper.cuh"
1819

19-
template <class T>
20-
struct equal_to_val
21-
{
22-
T val_;
23-
24-
constexpr equal_to_val(const T& val) noexcept
25-
: val_(val)
26-
{}
27-
28-
__device__ constexpr bool operator()(const T& val) const noexcept
29-
{
30-
return val == val_;
31-
}
32-
};
33-
3420
template <typename T>
3521
static void basic(nvbench::state& state, nvbench::type_list<T>)
3622
{
@@ -51,7 +37,8 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
5137

5238
state.exec(
5339
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
54-
do_not_optimize(cuda::std::none_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
40+
do_not_optimize(
41+
cuda::std::none_of(cuda_policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
5542
});
5643
}
5744

‎thrust/benchmarks/bench/all_of/basic.cu‎

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -5,26 +5,12 @@
55
#include <thrust/fill.h>
66
#include <thrust/logical.h>
77

8+
#include <cuda/functional>
89
#include <cuda/memory_pool>
910
#include <cuda/stream>
1011

1112
#include "nvbench_helper.cuh"
1213

13-
template <class T>
14-
struct equal_to_val
15-
{
16-
T val_;
17-
18-
constexpr equal_to_val(const T& val) noexcept
19-
: val_(val)
20-
{}
21-
22-
__device__ constexpr bool operator()(const T& val) const noexcept
23-
{
24-
return val == val_;
25-
}
26-
};
27-
2814
template <typename T>
2915
static void basic(nvbench::state& state, nvbench::type_list<T>)
3016
{
@@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
4329

4430
caching_allocator_t alloc{};
4531

46-
state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
47-
[&](nvbench::launch& launch) {
48-
do_not_optimize(thrust::all_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
49-
});
32+
state.exec(
33+
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
34+
do_not_optimize(thrust::all_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
35+
});
5036
}
5137

5238
NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))

‎thrust/benchmarks/bench/any_of/basic.cu‎

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -5,26 +5,12 @@
55
#include <thrust/fill.h>
66
#include <thrust/logical.h>
77

8+
#include <cuda/functional>
89
#include <cuda/memory_pool>
910
#include <cuda/stream>
1011

1112
#include "nvbench_helper.cuh"
1213

13-
template <class T>
14-
struct equal_to_val
15-
{
16-
T val_;
17-
18-
constexpr equal_to_val(const T& val) noexcept
19-
: val_(val)
20-
{}
21-
22-
__device__ constexpr bool operator()(const T& val) const noexcept
23-
{
24-
return val == val_;
25-
}
26-
};
27-
2814
template <typename T>
2915
static void basic(nvbench::state& state, nvbench::type_list<T>)
3016
{
@@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
4329

4430
caching_allocator_t alloc{};
4531

46-
state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
47-
[&](nvbench::launch& launch) {
48-
do_not_optimize(thrust::any_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
49-
});
32+
state.exec(
33+
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
34+
do_not_optimize(thrust::any_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
35+
});
5036
}
5137

5238
NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))

‎thrust/benchmarks/bench/find/basic.cu‎

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -4,26 +4,12 @@
44
#include <thrust/device_vector.h>
55
#include <thrust/find.h>
66

7+
#include <cuda/functional>
78
#include <cuda/memory_pool>
89
#include <cuda/stream>
910

1011
#include "nvbench_helper.cuh"
1112

12-
struct equal_to_val
13-
{
14-
size_t val_;
15-
16-
constexpr equal_to_val(const size_t val) noexcept
17-
: val_(val)
18-
{}
19-
20-
template <class T>
21-
__device__ constexpr bool operator()(const T& val) const noexcept
22-
{
23-
return val == val_;
24-
}
25-
};
26-
2713
template <typename T>
2814
static void basic(nvbench::state& state, nvbench::type_list<T>)
2915
{
@@ -42,10 +28,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
4228

4329
caching_allocator_t alloc{};
4430

45-
state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
46-
[&](nvbench::launch& launch) {
47-
do_not_optimize(thrust::find(policy(alloc, launch), dinput.begin(), dinput.end(), val));
48-
});
31+
state.exec(
32+
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
33+
do_not_optimize(thrust::find(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
34+
});
4935
}
5036

5137
NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))

‎thrust/benchmarks/bench/none_of/basic.cu‎

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -5,26 +5,12 @@
55
#include <thrust/fill.h>
66
#include <thrust/logical.h>
77

8+
#include <cuda/functional>
89
#include <cuda/memory_pool>
910
#include <cuda/stream>
1011

1112
#include "nvbench_helper.cuh"
1213

13-
template <class T>
14-
struct equal_to_val
15-
{
16-
T val_;
17-
18-
constexpr equal_to_val(const T& val) noexcept
19-
: val_(val)
20-
{}
21-
22-
__device__ constexpr bool operator()(const T& val) const noexcept
23-
{
24-
return val == val_;
25-
}
26-
};
27-
2814
template <typename T>
2915
static void basic(nvbench::state& state, nvbench::type_list<T>)
3016
{
@@ -43,10 +29,10 @@ static void basic(nvbench::state& state, nvbench::type_list<T>)
4329

4430
caching_allocator_t alloc{};
4531

46-
state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync,
47-
[&](nvbench::launch& launch) {
48-
do_not_optimize(thrust::none_of(policy(alloc, launch), dinput.begin(), dinput.end(), equal_to_val{val}));
49-
});
32+
state.exec(
33+
nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch | nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
34+
do_not_optimize(thrust::none_of(policy(alloc, launch), dinput.begin(), dinput.end(), cuda::equal_to_value{val}));
35+
});
5036
}
5137

5238
NVBENCH_BENCH_TYPES(basic, NVBENCH_TYPE_AXES(fundamental_types))

0 commit comments

Comments
 (0)