Skip to content
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

Jump to bottom

Improve build time of libcudf iterator tests #9788

Merged
merged 2 commits into from
Dec 1, 2021
Merged

Improve build time of libcudf iterator tests #9788

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ffc87a8
Improve build time of libcudf iterator tests
davidwendt Nov 29, 2021
1055f21
Merge branch 'branch-22.02' into iterator-build-time
davidwendt Nov 30, 2021
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
17 changes: 14 additions & 3 deletions cpp/tests/iterator/iterator_tests.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,8 +18,8 @@
#include <cudf_test/column_wrapper.hpp>
#include <cudf_test/type_lists.hpp>

#include <cudf/detail/iterator.cuh> // include iterator header
#include <cudf/detail/utilities/transform_unary_functions.cuh> //for meanvar
#include <cudf/detail/iterator.cuh>
#include <cudf/detail/utilities/transform_unary_functions.cuh> // for meanvar
#include <cudf/detail/utilities/vector_factories.hpp>

#include <rmm/cuda_stream_view.hpp>
Expand All @@ -28,6 +28,7 @@

#include <thrust/equal.h>
#include <thrust/functional.h>
#include <thrust/logical.h>
#include <thrust/transform.h>

#include <cub/device/device_reduce.cuh>
Expand Down Expand Up @@ -83,7 +84,17 @@ struct IteratorTest : public cudf::test::BaseFixture {
EXPECT_EQ(thrust::distance(d_in, d_in_last), num_items);
auto dev_expected = cudf::detail::make_device_uvector_sync(expected);

bool result = thrust::equal(thrust::device, d_in, d_in_last, dev_expected.begin());
// using a temporary vector and calling transform and all_of separately is
// equivalent to thrust::equal but compiles ~3x faster
auto dev_results = rmm::device_uvector<bool>(num_items, rmm::cuda_stream_default);
thrust::transform(thrust::device,
d_in,
d_in_last,
dev_expected.begin(),
dev_results.begin(),
thrust::equal_to{});
auto result = thrust::all_of(
thrust::device, dev_results.begin(), dev_results.end(), thrust::identity<bool>{});
EXPECT_TRUE(result) << "thrust test";
}

Expand Down
37 changes: 18 additions & 19 deletions cpp/tests/iterator/optional_iterator_test_numeric.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,21 +50,15 @@ struct transformer_optional_meanvar {
}
};

struct sum_if_not_null {
template <typename T>
CUDA_HOST_DEVICE_CALLABLE thrust::optional<T> operator()(const thrust::optional<T>& lhs,
const thrust::optional<T>& rhs)
{
return lhs.value_or(T{0}) + rhs.value_or(T{0});
}
template <typename T>
struct optional_to_meanvar {
CUDA_HOST_DEVICE_CALLABLE T operator()(const thrust::optional<T>& v) { return v.value_or(T{0}); }
};

// TODO: enable this test also at __CUDACC_DEBUG__
// This test causes fatal compilation error only at device debug mode.
// Workaround: exclude this test only at device debug mode.
#if !defined(__CUDACC_DEBUG__)
// This test computes `count`, `sum`, `sum_of_squares` at a single reduction call.
// It would be useful for `var`, `std` operation
TYPED_TEST(NumericOptionalIteratorTest, mean_var_output)
{
using T = TypeParam;
Expand Down Expand Up @@ -104,22 +98,27 @@ TYPED_TEST(NumericOptionalIteratorTest, mean_var_output)
expected_value.value_squared = std::accumulate(
replaced_array.begin(), replaced_array.end(), T{0}, [](T acc, T i) { return acc + i * i; });

// std::cout << "expected <mixed_output> = " << expected_value << std::endl;

// GPU test
auto it_dev = d_col->optional_begin<T>(cudf::contains_nulls::YES{});
auto it_dev_squared = thrust::make_transform_iterator(it_dev, transformer);
auto result = thrust::reduce(it_dev_squared,
it_dev_squared + d_col->size(),
thrust::optional<T_output>{T_output{}},
sum_if_not_null{});

// this can be computed with a single reduce and without a temporary output vector
// but the approach increases the compile time by ~2x
auto results = rmm::device_uvector<T_output>(d_col->size(), rmm::cuda_stream_default);
thrust::transform(thrust::device,
it_dev_squared,
it_dev_squared + d_col->size(),
results.begin(),
optional_to_meanvar<T_output>{});
auto result = thrust::reduce(thrust::device, results.begin(), results.end(), T_output{});

if (not std::is_floating_point<T>()) {
EXPECT_EQ(expected_value, *result) << "optional iterator reduction sum";
EXPECT_EQ(expected_value, result) << "optional iterator reduction sum";
} else {
EXPECT_NEAR(expected_value.value, result->value, 1e-3) << "optional iterator reduction sum";
EXPECT_NEAR(expected_value.value_squared, result->value_squared, 1e-3)
EXPECT_NEAR(expected_value.value, result.value, 1e-3) << "optional iterator reduction sum";
EXPECT_NEAR(expected_value.value_squared, result.value_squared, 1e-3)
<< "optional iterator reduction sum squared";
EXPECT_EQ(expected_value.count, result->count) << "optional iterator reduction count";
EXPECT_EQ(expected_value.count, result.count) << "optional iterator reduction count";
}
}
#endif