diff --git a/cpp/src/sqls/sqls_ops.cu b/cpp/src/sqls/sqls_ops.cu index 245115d49ed8..3a8a6c56cd54 100644 --- a/cpp/src/sqls/sqls_ops.cu +++ b/cpp/src/sqls/sqls_ops.cu @@ -1101,7 +1101,7 @@ gdf_error gdf_group_by_single(int ncols, // # columns return GDF_DATASET_EMPTY; } for (int i = 0; i < ncols; ++i) { - GDF_REQUIRE(!cols[i]->valid || !cols[i]->null_count, GDF_VALIDITY_UNSUPPORTED); + GDF_REQUIRE(!cols[i]->valid || !cols[i]->null_count, GDF_VALIDITY_UNSUPPORTED); } GDF_REQUIRE(!col_agg->valid || !col_agg->null_count, GDF_VALIDITY_UNSUPPORTED); diff --git a/cpp/src/utilities/bit_util.cuh b/cpp/src/utilities/bit_util.cuh index fd0c7afe8eee..e42133dbae84 100644 --- a/cpp/src/utilities/bit_util.cuh +++ b/cpp/src/utilities/bit_util.cuh @@ -17,28 +17,10 @@ */ #pragma once +#include "cudf_utils.h" namespace gdf { namespace util { -static constexpr int ValidSize = 32; -using ValidType = uint32_t; - - -// Instead of this function, use gdf_get_num_chars_bitmask from gdf/utils.h -//__host__ __device__ __forceinline__ -// size_t -// valid_size(size_t column_length) -//{ -// const size_t n_ints = (column_length / ValidSize) + ((column_length % ValidSize) ? 1 : 0); -// return n_ints * sizeof(ValidType); -//} - -// Instead of this function, use gdf_is_valid from gdf/utils.h -///__host__ __device__ __forceinline__ bool get_bit(const gdf_valid_type* const bits, size_t i) -///{ -/// return bits == nullptr? true : bits[i >> size_t(3)] & (1 << (i & size_t(7))); -///} - __host__ __device__ __forceinline__ uint8_t byte_bitmask(size_t i) @@ -70,6 +52,19 @@ __host__ __device__ __forceinline__ size_t last_byte_index(size_t column_size) return (column_size + 8 - 1) / 8; } + +__host__ __forceinline__ size_t null_count(uint8_t* const bits, size_t column_size) +{ + size_t count = 0; + for(size_t i = 0; i < column_size; i++) + { + if (gdf_is_valid(bits, i) == false) + count++; + } + return count; +} + + static inline std::string chartobin(gdf_valid_type c, size_t size = 8) { std::string bin; diff --git a/cpp/tests/groupby/groupby_test.cu b/cpp/tests/groupby/groupby_test.cu index 00482f5344f9..fb32a2fdf9e9 100644 --- a/cpp/tests/groupby/groupby_test.cu +++ b/cpp/tests/groupby/groupby_test.cu @@ -36,6 +36,12 @@ #include "test_parameters.cuh" #include "groupby_test_helpers.cuh" +// See this header for all valid handling +#include "valid_vectors.h" + +#include +#include + // A new instance of this class will be created for each *TEST(GroupTest, ...) // Put all repeated setup and validation stuff here template @@ -104,7 +110,7 @@ struct GroupTest : public GdfTest { template gdf_col_pointer create_gdf_column(std::vector const & host_vector, - const gdf_size_type n_count = 0) + const gdf_size_type n_count = 0, gdf_valid_type* host_valid = nullptr) { // Deduce the type and set the gdf_dtype accordingly gdf_dtype gdf_col_type = N_GDF_TYPES; @@ -127,10 +133,16 @@ struct GroupTest : public GdfTest { // Allocate device storage for gdf_column and copy contents from host_vector EXPECT_EQ(RMM_ALLOC(&(the_column->data), host_vector.size() * sizeof(col_type), 0), RMM_SUCCESS); EXPECT_EQ(cudaMemcpy(the_column->data, host_vector.data(), host_vector.size() * sizeof(col_type), cudaMemcpyHostToDevice), cudaSuccess); - - int valid_size = gdf_get_num_chars_bitmask(host_vector.size()); - EXPECT_EQ(RMM_ALLOC((void**)&(the_column->valid), valid_size, 0), RMM_SUCCESS); - EXPECT_EQ(cudaMemset(the_column->valid, 0xff, valid_size), cudaSuccess); + if(host_valid != nullptr) { + auto valid_size = gdf_get_num_chars_bitmask(host_vector.size()); + EXPECT_EQ(RMM_ALLOC((void**)&(the_column->valid), valid_size, 0), RMM_SUCCESS); + EXPECT_EQ(cudaMemcpy(the_column->valid, host_valid, valid_size, cudaMemcpyHostToDevice), cudaSuccess); + } + else { + int valid_size = gdf_get_num_chars_bitmask(host_vector.size()); + EXPECT_EQ(RMM_ALLOC((void**)&(the_column->valid), valid_size, 0), RMM_SUCCESS); + EXPECT_EQ(cudaMemset(the_column->valid, 0xff, valid_size), cudaSuccess); + } // Fill the gdf_column members the_column->null_count = n_count; @@ -458,7 +470,220 @@ TYPED_TEST(GroupTest, EmptyInput) // Create a new derived class from JoinTest so we can do a new Typed Test set of tests template struct GroupValidTest : public GroupTest -{ }; +{ + + // multi_column_t is a tuple of vectors. The number of vectors in the tuple + // determines the number of columns to be grouped, and the value_type of each + // vector determiens the data type of the column + using multi_column_t = typename test_parameters::multi_column_t; + + //output_t is the output type of the aggregation column + using output_t = typename test_parameters::output_type; + + //map_t is used for reference solution + using map_t = typename test_parameters::ref_map_type; + + //tuple_t is tuple of datatypes associated with each column to be grouped + using tuple_t = typename test_parameters::tuple_t; + + using gdf_col_pointer = typename std::unique_ptr>; + + //contains input valid generated for gdf calculation and reference solution + std::vector input_key_valids; + + //contains the input valid aggregation column + host_valid_pointer input_value_valid; + + //contains grouped by column valid output of the gdf groupby call + std::vector output_key_valid; + + //contains the aggregated output column valid + host_valid_pointer output_value_valid; + + /* --------------------------------------------------------------------------*/ + /** + * @Synopsis Creates a unique_ptr that wraps a gdf_column structure intialized with a host vector + * + * @Param host_vector The host vector whose data is used to initialize the gdf_column + * + * @Returns A unique_ptr wrapping the new gdf_column + */ + /* ----------------------------------------------------------------------------*/ + // Compile time recursion to convert each vector in a tuple of vectors into + // a gdf_column and append it to a vector of gdf_columns + template + inline typename std::enable_if::type + convert_tuple_to_gdf_columns(std::vector &gdf_columns,std::tuple...>& t, std::vector& valids) + { + //bottom of compile-time recursion + //purposely empty... + } + + template + inline typename std::enable_if::type + convert_tuple_to_gdf_columns(std::vector &gdf_columns,std::tuple...>& t, std::vector& valids) + { + // Creates a gdf_column for the current vector and pushes it onto + // the vector of gdf_columns + if (valids.size() == 0) { + gdf_columns.push_back(this->create_gdf_column(std::get(t))); + } + else{ + auto valid = valids[I].get(); + auto column_size = std::get(t).size(); + auto null_count = gdf::util::null_count(valid, column_size); + gdf_columns.push_back(this->create_gdf_column(std::get(t), null_count, valid)); + } + + //recurse to next vector in tuple + convert_tuple_to_gdf_columns(gdf_columns, t, valids); + } + + // Converts a tuple of host vectors into a vector of gdf_columns + std::vector + initialize_gdf_columns(multi_column_t host_columns, std::vector& output_key_valid) + { + std::vector gdf_columns; + convert_tuple_to_gdf_columns(gdf_columns, host_columns, output_key_valid); + return gdf_columns; + } + + /* --------------------------------------------------------------------------*/ + /** + * @Synopsis Initializes key columns and aggregation column for gdf group by call + * + * @Param key_count The number of unique keys + * @Param value_per_key The number of times a random aggregation value is generated for a key + * @Param max_key The maximum value of the key columns + * @Param max_val The maximum value of aggregation column + * @Param print Optionally print the keys and aggregation columns for debugging + */ + /* ----------------------------------------------------------------------------*/ + void create_input_with_nulls(const size_t key_count, const size_t value_per_key, + const size_t max_key, const size_t max_val, + bool print = false) { + size_t shuffle_seed = rand(); + initialize_keys(this->input_key, key_count, value_per_key, max_key, shuffle_seed); + initialize_values(this->input_value, key_count, value_per_key, max_val, shuffle_seed); + + // Init valids + auto column_length = this->input_value.size(); + auto n_tuples = std::tuple_size::value; + for (size_t i = 0; i < n_tuples; i++) + input_key_valids.push_back(create_and_init_valid(column_length)); + input_value_valid = create_and_init_valid(column_length); + + this->gdf_input_key_columns = initialize_gdf_columns(this->input_key, input_key_valids); + + auto null_count = gdf::util::null_count(input_value_valid.get(), column_length); + this->gdf_input_value_column = this->create_gdf_column(this->input_value, null_count, input_value_valid.get()); + + // Fill vector of raw pointers to gdf_columns + for(auto const& c : this->gdf_input_key_columns){ + this->gdf_raw_input_key_columns.push_back(c.get()); + } + this->gdf_raw_input_val_column = this->gdf_input_value_column.get(); + + if(print) + { + std::cout << "Key column(s) created. Size: " << std::get<0>(this->input_key).size() << std::endl; + print_tuple_vector(this->input_key, input_key_valids); + + std::cout << "Value column(s) created. Size: " << this->input_value.size() << std::endl; + print_vector(this->input_value, input_value_valid.get()); + std::cout << std::endl; + + } + } + + + void create_gdf_output_buffers_with_nulls(const size_t key_count, const size_t value_per_key) { + initialize_keys(this->output_key, key_count, value_per_key, 0, 0, false); + initialize_values(this->output_value, key_count, value_per_key, 0, 0); + + // Init Valids + auto column_length = this->output_value.size(); + auto n_tuples = std::tuple_size::value; + for (size_t i = 0; i < n_tuples; i++) + output_key_valid.push_back(create_and_init_valid(column_length)); + output_value_valid = create_and_init_valid(column_length); + + this->gdf_output_key_columns = initialize_gdf_columns(this->output_key, output_key_valid); + auto null_count = gdf::util::null_count(output_value_valid.get(), column_length); + + auto str = gdf::util::gdf_valid_to_str(output_value_valid.get(), column_length); + + this->gdf_output_value_column = this->create_gdf_column(this->output_value, null_count, output_value_valid.get()); + for(auto const& c : this->gdf_output_key_columns){ + this->gdf_raw_output_key_columns.push_back(c.get()); + } + this->gdf_raw_output_val_column = this->gdf_output_value_column.get(); + } + + map_t + compute_reference_solution_with_nulls(void) { + map_t key_val_map; + + auto get_input_key_valids = [](std::vector& valids, int i) { + std::basic_string valid_key; + std::for_each(valids.begin(), + valids.end(), + [&i, &valid_key](host_valid_pointer& valid){ + bool b1 = gdf_is_valid(valid.get(), i); + valid_key.push_back(b1); + }); + return valid_key; + }; + + if (test_parameters::op != agg_op::AVG) { + AggOp agg; + for (size_t i = 0; i < this->input_value.size(); ++i) { + bool valid_value = gdf_is_valid(this->input_value_valid.get(), i); + auto valid_key = get_input_key_valids(this->input_key_valids, i); + auto l_key = extractKeyWithNulls(this->input_key, valid_key, i); + + if (valid_value) { + auto sch = key_val_map.find(l_key); + if (sch != key_val_map.end()) { + key_val_map[l_key] = agg(sch->second, this->input_value[i]); + } else { + key_val_map[l_key] = agg(this->input_value[i]); + } + } + } + } else { + std::map counters; + AggOp agg; + for (size_t i = 0; i < this->input_value.size(); ++i) { + bool valid_value = gdf_is_valid(this->input_value_valid.get(), i); + auto valid_key = get_input_key_valids(this->input_key_valids, i); + auto l_key = extractKeyWithNulls(this->input_key, valid_key, i); + if(valid_value) + counters[l_key]++; + if (valid_value) { + auto sch = key_val_map.find(l_key); + if (sch != key_val_map.end()) { + key_val_map[l_key] = agg(sch->second, this->input_value[i]); + } else { + key_val_map[l_key] = agg(this->input_value[i]); + } + } + } + for (auto& e : key_val_map) { + e.second = e.second/counters[e.first]; + } + } + return key_val_map; + } + void print_reference_solution(map_t & dicc) { + std::stringstream ss; + for(auto &iter : dicc) { + print_tuple_value(ss, iter.first); + ss << " => " << iter.second << std::endl; + } + std::cout << ss.str() << std::endl; + } +}; TYPED_TEST_CASE(GroupValidTest, ValidTestImplementations); @@ -473,3 +698,21 @@ TYPED_TEST(GroupValidTest, ReportValidMaskError) this->compute_gdf_result(GDF_VALIDITY_UNSUPPORTED); } +TYPED_TEST(GroupValidTest, GroupbyValidExampleTest) +{ + const size_t num_keys = 4; + const size_t num_values_per_key = 4; + const size_t max_key = num_keys * 1; + const size_t max_val = 10; + + this->create_input_with_nulls(num_keys, num_values_per_key, max_key, max_val, true); + auto reference_map = this->compute_reference_solution_with_nulls(); + this->print_reference_solution(reference_map); + this->create_gdf_output_buffers_with_nulls(num_keys, num_values_per_key); + + // @todo: merge with hash-based or sort-based solution with null support + // this->compute_gdf_result(); + + // @todo: compare gdf solution and the reference solution both with possible nulls + // this->compare_gdf_result(reference_map); +} diff --git a/cpp/tests/groupby/groupby_test_helpers.cuh b/cpp/tests/groupby/groupby_test_helpers.cuh index 32331224ea9d..cb2b77191f11 100644 --- a/cpp/tests/groupby/groupby_test_helpers.cuh +++ b/cpp/tests/groupby/groupby_test_helpers.cuh @@ -1,6 +1,9 @@ #pragma once #include "test_parameters.cuh" #include +#include +#include +#include "valid_vectors.h" //Terminating call //Extract the value of the Ith element of a tuple of vectors keys @@ -30,6 +33,61 @@ extractKey(std::tuple...>& keys, const size_t index) { return key; } + + +//Terminating call +//Extract the value of the Ith element of a tuple of vectors keys +//at index location and store it in the Ith element of the tuple key +template +inline typename std::enable_if::type +extractWithNulls(const std::tuple...>& keys, std::basic_string &valids, const size_t index, std::tuple &key) { + +} + +//Extract the value (replace nulls with MAX_VALUE of ValueType) of the Ith element of a tuple of vectors keys +//at index location and store it in the Ith element of the tuple key +template +inline typename std::enable_if::type +extractWithNulls(const std::tuple...>& keys, std::basic_string &valids, const size_t index, std::tuple &key) { + using key_type = typename std::decay(keys))>::type; + std::get(key) = valids[I] ? std::get(keys)[index] : std::numeric_limits::max(); + extractWithNulls(keys, valids, index, key); +} + + +//Extract a tuple of values (replace nulls with MAX_VALUE of ValueType) from a tuple of vector for a given index +//keys Tuple of vectors of types Keys +//index Location of the value to be extracted in each vector +template +std::tuple +extractKeyWithNulls(std::tuple...>& keys, std::basic_string &valids, const size_t index) { + std::tuple key; + extractWithNulls(keys, valids, index, key); + return key; +} + + + +//Terminating call +//Extract the value of the Ith element of a tuple of vectors keys +//at index location and store it in the Ith element of the tuple key +template +inline typename std::enable_if::type +print_tuple_value(std::stringstream &ss, const std::tuple &item) { + +} + +//Extract the value (replace nulls with MAX_VALUE of ValueType) of the Ith element of a tuple of vectors keys +//at index location and store it in the Ith element of the tuple key +template +inline typename std::enable_if::type +print_tuple_value(std::stringstream &ss, const std::tuple &item) { + using val_type = typename std::decay(item))>::type; + auto val = std::get(item) == std::numeric_limits::max()? "@" : std::to_string(std::get(item)); + ss << val << "|" ; + print_tuple_value(ss, item); +} + //Struct to generate random values of type T template struct RandomValues { @@ -231,34 +289,54 @@ void copy_output( copy_gdf_column(group_by_output_value, output_value); } + +// +// custom functions +// + + // Prints a vector template -void print_vector(std::vector& v) +void print_vector(std::vector& v, gdf_valid_type* valid = nullptr) { - std::copy(v.begin(), v.end(), std::ostream_iterator(std::cout, ", ")); + auto functor = [&valid, &v] (int index) -> std::string { + if ( gdf_is_valid(valid, index)) { + if (sizeof(v[index]) == 1) + return std::to_string((int)v[index]); + return std::to_string(v[index]); + } + return std::string("@"); + }; + std::vector indexes(v.size()); + std::iota (std::begin(indexes), std::end(indexes), 0); + std::transform(indexes.begin(), indexes.end(), std::ostream_iterator(std::cout, ", "), functor); } template inline typename std::enable_if::type -print_tuple_vector(std::tuple...>& t) +print_tuple_vector(std::tuple...>& t, const std::vector& valids = {}) { - //bottom of compile-time recursion - //purposely empty... + //bottom of compile-time recursion + //purposely empty... } //compile time recursion to print a tuple of vectors template inline typename std::enable_if::type -print_tuple_vector(std::tuple...>& t) +print_tuple_vector(std::tuple...>& t, const std::vector& valids = {}) { - // print the current vector: - print_vector(std::get(t)); - std::cout << std::endl; + // print the current vector: + if (valids.size() == 0) + print_vector(std::get(t), nullptr); + else + print_vector(std::get(t), valids[I].get()); + std::cout << std::endl; - //recurse to next vector in tuple - print_tuple_vector(t); + //recurse to next vector in tuple + print_tuple_vector(t, valids); } + //print a tuple recursively. Terminating empty call. template inline typename std::enable_if::type diff --git a/cpp/tests/groupby/test_parameters.cuh b/cpp/tests/groupby/test_parameters.cuh index 020c6b69c049..b164a65045f9 100644 --- a/cpp/tests/groupby/test_parameters.cuh +++ b/cpp/tests/groupby/test_parameters.cuh @@ -154,5 +154,7 @@ typedef ::testing::Types< > Implementations; typedef ::testing::Types< - TestParameters< agg_op::AVG, HASH, VTuple, int32_t> + TestParameters< agg_op::AVG, HASH, VTuple, int32_t>, + TestParameters< agg_op::SUM, HASH, VTuple, int64_t >, + TestParameters< agg_op::MIN, HASH, VTuple, int64_t > > ValidTestImplementations; diff --git a/cpp/tests/groupby/valid_vectors.h b/cpp/tests/groupby/valid_vectors.h new file mode 100644 index 000000000000..9fb87632dd22 --- /dev/null +++ b/cpp/tests/groupby/valid_vectors.h @@ -0,0 +1,84 @@ +/* +* Copyright (c) 2018, 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. +*/ +#pragma once +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// host_valid_pointer is a wrapper for gdf_valid_type* with custom deleter +using host_valid_pointer = typename std::shared_ptr; + +// Create a valid pointer and init randomly the last half column +static inline host_valid_pointer create_and_init_valid(size_t length) +{ + auto deleter = [](gdf_valid_type* valid) { delete[] valid; }; + auto n_bytes = gdf_get_num_chars_bitmask(length); + auto valid_bits = new gdf_valid_type[n_bytes]; + + for (size_t i = 0; i < length; ++i) { + if (i < length / 2 || std::rand() % 2 == 1) { + gdf::util::turn_bit_on(valid_bits, i); + } else { + gdf::util::turn_bit_off(valid_bits, i); + } + } + return host_valid_pointer{ valid_bits, deleter }; +} + + +// Create a valid pointer and init randomly the last half column +static inline +void init_valid(std::uint8_t * valid_bits, + const std::int64_t start_offset, + size_t length) +{ + gdf_valid_type current_byte; + size_t byte_offset = start_offset / 8; + size_t bit_offset = start_offset % 8; + + if (length > 0) { + current_byte = valid_bits[byte_offset]; + } + for (size_t i = 0; i < length; ++i) { + if (i < length / 2 || std::rand() % 2 == 1) { + current_byte |= gdf::util::byte_bitmask(bit_offset); + } else { + current_byte &= gdf::util::flipped_bitmask(bit_offset); + } + ++bit_offset; + valid_bits[byte_offset] = current_byte; + if (bit_offset == 8) { + bit_offset = 0; + ++byte_offset; + current_byte = valid_bits[byte_offset]; + } + } +} + +// Initialize valids +static inline void initialize_valids(std::vector& valids, size_t size, size_t length) +{ + for (size_t i = 0; i < size; ++i) { + valids.push_back(create_and_init_valid(length)); + } +} +