Skip to content

Commit

Permalink
Merge branch 'branch-21.10' of github.com:rapidsai/cudf into bug-memo…
Browse files Browse the repository at this point in the history
…ry_errors_libcudf
  • Loading branch information
karthikeyann committed Aug 7, 2021
2 parents 45d213b + 115f3b6 commit 571bb5a
Show file tree
Hide file tree
Showing 140 changed files with 6,004 additions and 4,471 deletions.
3 changes: 2 additions & 1 deletion conda/recipes/libcudf/meta.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ test:
- test -f $PREFIX/lib/libcudftestutil.a
- test -f $PREFIX/include/cudf/aggregation.hpp
- test -f $PREFIX/include/cudf/ast/transform.hpp
- test -f $PREFIX/include/cudf/ast/detail/linearizer.hpp
- test -f $PREFIX/include/cudf/ast/detail/expression_parser.hpp
- test -f $PREFIX/include/cudf/ast/detail/operators.hpp
- test -f $PREFIX/include/cudf/ast/nodes.hpp
- test -f $PREFIX/include/cudf/ast/operators.hpp
Expand Down Expand Up @@ -102,6 +102,7 @@ test:
- test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp
- test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h
- test -f $PREFIX/include/cudf/detail/utilities/vector_factories.hpp
- test -f $PREFIX/include/cudf/detail/utilities/visitor_overload.hpp
- test -f $PREFIX/include/cudf/dictionary/detail/concatenate.hpp
- test -f $PREFIX/include/cudf/dictionary/detail/encode.hpp
- test -f $PREFIX/include/cudf/dictionary/detail/merge.hpp
Expand Down
3 changes: 2 additions & 1 deletion cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -151,7 +151,7 @@ add_library(cudf
src/aggregation/aggregation.cpp
src/aggregation/aggregation.cu
src/aggregation/result_cache.cpp
src/ast/linearizer.cpp
src/ast/expression_parser.cpp
src/ast/transform.cu
src/binaryop/binaryop.cpp
src/binaryop/compiled/binary_ops.cu
Expand Down Expand Up @@ -305,6 +305,7 @@ add_library(cudf
src/join/cross_join.cu
src/join/hash_join.cu
src/join/join.cu
src/join/join_utils.cu
src/join/semi_join.cu
src/lists/contains.cu
src/lists/combine/concatenate_list_elements.cu
Expand Down
2 changes: 2 additions & 0 deletions cpp/benchmarks/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@ target_link_libraries(cudf_datagen
GTest::gmock_main
GTest::gtest_main
benchmark::benchmark
nvbench::nvbench
Threads::Threads
cudf)

Expand Down Expand Up @@ -102,6 +103,7 @@ ConfigureBench(STREAM_COMPACTION_BENCH stream_compaction/drop_duplicates_benchma
###################################################################################################
# - join benchmark --------------------------------------------------------------------------------
ConfigureBench(JOIN_BENCH join/join_benchmark.cu join/conditional_join_benchmark.cu)
ConfigureNVBench(JOIN_NVBENCH join/join_nvbench.cu)

###################################################################################################
# - iterator benchmark ----------------------------------------------------------------------------
Expand Down
71 changes: 71 additions & 0 deletions cpp/benchmarks/fixture/rmm_pool_raii.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,71 @@
/*
* Copyright (c) 2021, 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 <rmm/mr/device/cuda_memory_resource.hpp>
#include <rmm/mr/device/owning_wrapper.hpp>
#include <rmm/mr/device/per_device_resource.hpp>
#include <rmm/mr/device/pool_memory_resource.hpp>

namespace cudf {

/**
* @brief An RAII class setting up RMM memory pool for `nvbench` benchmarks
*
* This is a temporary solution before templated fixtures tests are supported
* in `nvbench`. Similarly to `cudf::benchmark`, creating this RAII object in
* each benchmark will ensure that the RAPIDS Memory Manager pool mode is used
* in benchmarks, which eliminates memory allocation / deallocation performance
* overhead from the benchmark.
*
* Example:
*
* void my_benchmark(nvbench::state& state) {
* cudf::rmm_pool_raii pool_raii;
* state.exec([](nvbench::launch& launch) {
* // benchmark stuff
* });
* }
*
* NVBENCH_BENCH(my_benchmark);
*/
class rmm_pool_raii {
private:
// memory resource factory helpers
inline auto make_cuda() { return std::make_shared<rmm::mr::cuda_memory_resource>(); }

inline auto make_pool()
{
return rmm::mr::make_owning_wrapper<rmm::mr::pool_memory_resource>(make_cuda());
}

public:
rmm_pool_raii()
{
mr = make_pool();
rmm::mr::set_current_device_resource(mr.get()); // set default resource to pool
}

~rmm_pool_raii()
{
rmm::mr::set_current_device_resource(nullptr);
mr.reset();
}

private:
std::shared_ptr<rmm::mr::device_memory_resource> mr;
};

} // namespace cudf
111 changes: 11 additions & 100 deletions cpp/benchmarks/join/conditional_join_benchmark.cu
Original file line number Diff line number Diff line change
Expand Up @@ -14,106 +14,12 @@
* limitations under the License.
*/

#include <benchmark/benchmark.h>

#include <thrust/iterator/counting_iterator.h>

#include <cudf/ast/nodes.hpp>
#include <cudf/ast/operators.hpp>
#include <cudf/column/column_factories.hpp>
#include <cudf/join.hpp>
#include <cudf/table/table.hpp>
#include <cudf/table/table_view.hpp>
#include <cudf/utilities/error.hpp>
#include <cudf_test/base_fixture.hpp>
#include <cudf_test/column_wrapper.hpp>

#include <fixture/benchmark_fixture.hpp>
#include <synchronization/synchronization.hpp>

#include <vector>

#include "generate_input_tables.cuh"
#include <benchmarks/join/join_benchmark_common.hpp>

template <typename key_type, typename payload_type>
class ConditionalJoin : public cudf::benchmark {
};

template <typename key_type, typename payload_type, bool Nullable, typename Join>
static void BM_join(benchmark::State& state, Join JoinFunc)
{
const cudf::size_type build_table_size{(cudf::size_type)state.range(0)};
const cudf::size_type probe_table_size{(cudf::size_type)state.range(1)};
const cudf::size_type rand_max_val{build_table_size * 2};
const double selectivity = 0.3;
const bool is_build_table_key_unique = true;

// Generate build and probe tables
cudf::test::UniformRandomGenerator<cudf::size_type> rand_gen(0, build_table_size);
auto build_random_null_mask = [&rand_gen](int size) {
if (Nullable) {
// roughly 25% nulls
auto validity = thrust::make_transform_iterator(
thrust::make_counting_iterator(0),
[&rand_gen](auto i) { return (rand_gen.generate() & 3) == 0; });
return cudf::test::detail::make_null_mask(validity, validity + size);
} else {
return cudf::create_null_mask(size, cudf::mask_state::UNINITIALIZED);
}
};

std::unique_ptr<cudf::column> build_key_column = [&]() {
return Nullable ? cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<key_type>()),
build_table_size,
build_random_null_mask(build_table_size))
: cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<key_type>()),
build_table_size);
}();
std::unique_ptr<cudf::column> probe_key_column = [&]() {
return Nullable ? cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<key_type>()),
probe_table_size,
build_random_null_mask(probe_table_size))
: cudf::make_numeric_column(cudf::data_type(cudf::type_to_id<key_type>()),
probe_table_size);
}();

generate_input_tables<key_type, cudf::size_type>(
build_key_column->mutable_view().data<key_type>(),
build_table_size,
probe_key_column->mutable_view().data<key_type>(),
probe_table_size,
selectivity,
rand_max_val,
is_build_table_key_unique);

auto payload_data_it = thrust::make_counting_iterator(0);
cudf::test::fixed_width_column_wrapper<payload_type> build_payload_column(
payload_data_it, payload_data_it + build_table_size);

cudf::test::fixed_width_column_wrapper<payload_type> probe_payload_column(
payload_data_it, payload_data_it + probe_table_size);

CHECK_CUDA(0);

cudf::table_view build_table({build_key_column->view(), build_payload_column});
cudf::table_view probe_table({probe_key_column->view(), probe_payload_column});

// Benchmark the inner join operation

for (auto _ : state) {
cuda_event_timer raii(state, true, rmm::cuda_stream_default);

// Common column references.
const auto col_ref_left_0 = cudf::ast::column_reference(0);
const auto col_ref_right_0 = cudf::ast::column_reference(0, cudf::ast::table_reference::RIGHT);
auto left_zero_eq_right_zero =
cudf::ast::expression(cudf::ast::ast_operator::EQUAL, col_ref_left_0, col_ref_right_0);

auto result =
JoinFunc(probe_table, build_table, left_zero_eq_right_zero, cudf::null_equality::UNEQUAL);
}
}

#define CONDITIONAL_INNER_JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \
BENCHMARK_TEMPLATE_DEFINE_F(ConditionalJoin, name, key_type, payload_type) \
(::benchmark::State & st) \
Expand All @@ -124,7 +30,8 @@ static void BM_join(benchmark::State& state, Join JoinFunc)
cudf::null_equality compare_nulls) { \
return cudf::conditional_inner_join(left, right, binary_pred, compare_nulls); \
}; \
BM_join<key_type, payload_type, nullable>(st, join); \
constexpr bool is_conditional = true; \
BM_join<key_type, payload_type, nullable, is_conditional>(st, join); \
}

CONDITIONAL_INNER_JOIN_BENCHMARK_DEFINE(conditional_inner_join_32bit, int32_t, int32_t, false);
Expand All @@ -142,7 +49,8 @@ CONDITIONAL_INNER_JOIN_BENCHMARK_DEFINE(conditional_inner_join_64bit_nulls, int6
cudf::null_equality compare_nulls) { \
return cudf::conditional_left_join(left, right, binary_pred, compare_nulls); \
}; \
BM_join<key_type, payload_type, nullable>(st, join); \
constexpr bool is_conditional = true; \
BM_join<key_type, payload_type, nullable, is_conditional>(st, join); \
}

CONDITIONAL_LEFT_JOIN_BENCHMARK_DEFINE(conditional_left_join_32bit, int32_t, int32_t, false);
Expand All @@ -160,7 +68,8 @@ CONDITIONAL_LEFT_JOIN_BENCHMARK_DEFINE(conditional_left_join_64bit_nulls, int64_
cudf::null_equality compare_nulls) { \
return cudf::conditional_inner_join(left, right, binary_pred, compare_nulls); \
}; \
BM_join<key_type, payload_type, nullable>(st, join); \
constexpr bool is_conditional = true; \
BM_join<key_type, payload_type, nullable, is_conditional>(st, join); \
}

CONDITIONAL_FULL_JOIN_BENCHMARK_DEFINE(conditional_full_join_32bit, int32_t, int32_t, false);
Expand All @@ -178,7 +87,8 @@ CONDITIONAL_FULL_JOIN_BENCHMARK_DEFINE(conditional_full_join_64bit_nulls, int64_
cudf::null_equality compare_nulls) { \
return cudf::conditional_left_anti_join(left, right, binary_pred, compare_nulls); \
}; \
BM_join<key_type, payload_type, nullable>(st, join); \
constexpr bool is_conditional = true; \
BM_join<key_type, payload_type, nullable, is_conditional>(st, join); \
}

CONDITIONAL_LEFT_ANTI_JOIN_BENCHMARK_DEFINE(conditional_left_anti_join_32bit,
Expand Down Expand Up @@ -208,7 +118,8 @@ CONDITIONAL_LEFT_ANTI_JOIN_BENCHMARK_DEFINE(conditional_left_anti_join_64bit_nul
cudf::null_equality compare_nulls) { \
return cudf::conditional_left_semi_join(left, right, binary_pred, compare_nulls); \
}; \
BM_join<key_type, payload_type, nullable>(st, join); \
constexpr bool is_conditional = true; \
BM_join<key_type, payload_type, nullable, is_conditional>(st, join); \
}

CONDITIONAL_LEFT_SEMI_JOIN_BENCHMARK_DEFINE(conditional_left_semi_join_32bit,
Expand Down
Loading

0 comments on commit 571bb5a

Please sign in to comment.