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

Remove scalar_to_column_view #11241

Draft
wants to merge 3 commits into
base: branch-22.08
Choose a base branch
from
Draft

Remove scalar_to_column_view #11241

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7e7de2b
debug first pass
rwlee Jul 6, 2022
06570b0
Merge remote-tracking branch 'pub/branch-22.08' into rwlee/remove_s2cv
rwlee Jul 7, 2022
e0095ad
line of failure pinpointed
rwlee Jul 7, 2022
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
132 changes: 66 additions & 66 deletions cpp/src/binaryop/compiled/binary_ops.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,71 +36,71 @@ namespace binops {
namespace compiled {

namespace {
/**
* @brief Converts scalar to column_view with single element.
*
* @return pair with column_view and column containing any auxiliary data to create column_view from
* scalar
*/
struct scalar_as_column_view {
using return_type = typename std::pair<column_view, std::unique_ptr<column>>;
template <typename T, std::enable_if_t<(is_fixed_width<T>())>* = nullptr>
return_type operator()(scalar const& s, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)
{
auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
auto col_v =
column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data());
return std::pair{col_v, std::unique_ptr<column>(nullptr)};
}
template <typename T, std::enable_if_t<(!is_fixed_width<T>())>* = nullptr>
return_type operator()(scalar const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)
{
CUDF_FAIL("Unsupported type");
}
};
// specialization for cudf::string_view
template <>
scalar_as_column_view::return_type scalar_as_column_view::operator()<cudf::string_view>(
scalar const& s, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr)
{
using T = cudf::string_view;
auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
// /**
// * @brief Converts scalar to column_view with single element.
// *
// * @return pair with column_view and column containing any auxiliary data to create column_view from
// * scalar
// */
// struct scalar_as_column_view {
// using return_type = typename std::pair<column_view, std::unique_ptr<column>>;
// template <typename T, std::enable_if_t<(is_fixed_width<T>())>* = nullptr>
// return_type operator()(scalar const& s, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)
// {
// auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
// auto col_v =
// column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data());
// return std::pair{col_v, std::unique_ptr<column>(nullptr)};
// }
// template <typename T, std::enable_if_t<(!is_fixed_width<T>())>* = nullptr>
// return_type operator()(scalar const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)
// {
// CUDF_FAIL("Unsupported type");
// }
// };
// // specialization for cudf::string_view
// template <>
// scalar_as_column_view::return_type scalar_as_column_view::operator()<cudf::string_view>(
// scalar const& s, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr)
// {
// using T = cudf::string_view;
// auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));

// build offsets column from the string size
auto offsets_transformer_itr =
thrust::make_constant_iterator<size_type>(h_scalar_type_view.size());
auto offsets_column = strings::detail::make_offsets_child_column(
offsets_transformer_itr, offsets_transformer_itr + 1, stream, mr);
// // build offsets column from the string size
// auto offsets_transformer_itr =
// thrust::make_constant_iterator<size_type>(h_scalar_type_view.size());
// auto offsets_column = strings::detail::make_offsets_child_column(
// offsets_transformer_itr, offsets_transformer_itr + 1, stream, mr);

auto chars_column_v =
column_view(data_type{type_id::INT8}, h_scalar_type_view.size(), h_scalar_type_view.data());
// Construct string column_view
auto col_v = column_view(s.type(),
1,
nullptr,
(bitmask_type const*)s.validity_data(),
cudf::UNKNOWN_NULL_COUNT,
0,
{offsets_column->view(), chars_column_v});
return std::pair{col_v, std::move(offsets_column)};
}
// auto chars_column_v =
// column_view(data_type{type_id::INT8}, h_scalar_type_view.size(), h_scalar_type_view.data());
// // Construct string column_view
// auto col_v = column_view(s.type(),
// 1,
// nullptr,
// (bitmask_type const*)s.validity_data(),
// cudf::UNKNOWN_NULL_COUNT,
// 0,
// {offsets_column->view(), chars_column_v});
// return std::pair{col_v, std::move(offsets_column)};
// }

/**
* @brief Converts scalar to column_view with single element.
*
* @param scal scalar to convert
* @param stream CUDA stream used for device memory operations and kernel launches.
* @param mr Device memory resource used to allocate the returned column's device memory
* @return pair with column_view and column containing any auxiliary data to create
* column_view from scalar
*/
auto scalar_to_column_view(
scalar const& scal,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
{
return type_dispatcher(scal.type(), scalar_as_column_view{}, scal, stream, mr);
}
// /**
// * @brief Converts scalar to column_view with single element.
// *
// * @param scal scalar to convert
// * @param stream CUDA stream used for device memory operations and kernel launches.
// * @param mr Device memory resource used to allocate the returned column's device memory
// * @return pair with column_view and column containing any auxiliary data to create
// * column_view from scalar
// */
// auto scalar_to_column_view(
// scalar const& scal,
// rmm::cuda_stream_view stream,
// rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
// {
// return type_dispatcher(scal.type(), scalar_as_column_view{}, scal, stream, mr);
// }

// This functor does the actual comparison between string column value and a scalar string
// or between two string column values using a comparator
Expand Down Expand Up @@ -362,8 +362,8 @@ void binary_operation(mutable_column_view& out,
binary_operator op,
rmm::cuda_stream_view stream)
{
auto [lhsv, aux] = scalar_to_column_view(lhs, stream);
operator_dispatcher(out, lhsv, rhs, true, false, op, stream);
auto lhsc = make_column_from_scalar(lhs, 1, stream);
operator_dispatcher(out, lhsc->view(), rhs, true, false, op, stream);
}
// vector_scalar
void binary_operation(mutable_column_view& out,
Expand All @@ -372,8 +372,8 @@ void binary_operation(mutable_column_view& out,
binary_operator op,
rmm::cuda_stream_view stream)
{
auto [rhsv, aux] = scalar_to_column_view(rhs, stream);
operator_dispatcher(out, lhs, rhsv, false, true, op, stream);
auto rhsc = make_column_from_scalar(rhs, 1, stream);
operator_dispatcher(out, lhs, rhsc->view(), false, true, op, stream);
}
} // namespace compiled
} // namespace binops
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/binaryop/compiled/binary_ops.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -123,8 +123,8 @@ struct ops_wrapper {
// To suppress nvcc warning
return std::invoke_result_t<BinaryOperator, TypeCommon, TypeCommon>{};
}();
if constexpr (is_bool_result<BinaryOperator, TypeCommon, TypeCommon>())
out.element<decltype(result)>(i) = result;
if constexpr (is_bool_result<BinaryOperator, TypeCommon, TypeCommon>()){}
// out.element<decltype(result)>(i) = result;
else
type_dispatcher(out.type(), typed_casted_writer<decltype(result)>{}, i, out, result);
}
Expand Down
3 changes: 3 additions & 0 deletions cpp/src/binaryop/compiled/equality_ops.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,9 @@ void dispatch_equality_op(mutable_column_view& out,
auto rhsd = column_device_view::create(rhs, stream);
if (common_dtype) {
if (op == binary_operator::EQUAL) {
// cudf::test::print(out);
// cudf::test::print(lhs);
// cudf::test::print(rhs);
for_each(stream,
out.size(),
binary_op_device_dispatcher<ops::Equal>{
Expand Down
4 changes: 2 additions & 2 deletions cpp/tests/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -191,8 +191,8 @@ ConfigureTest(
binaryop/binop-verify-input-test.cpp
binaryop/binop-null-test.cpp
binaryop/binop-compiled-test.cpp
binaryop/binop-compiled-fixed_point-test.cpp
binaryop/binop-generic-ptx-test.cpp
# binaryop/binop-compiled-fixed_point-test.cpp
# binaryop/binop-generic-ptx-test.cpp
)

# ##################################################################################################
Expand Down
20 changes: 12 additions & 8 deletions cpp/tests/binaryop/assert-binops.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,10 +91,12 @@ void ASSERT_BINOP(column_view const& out,
auto out_data = out_h.first;

ASSERT_EQ(out_data.size(), rhs_data.size());
for (size_t i = 0; i < out_data.size(); ++i) {
auto lhs = out_data[i];
auto rhs = (TypeOut)(op(lhs_h, rhs_data[i]));
ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs);
if(lhs.is_valid()) {
for (size_t i = 0; i < out_data.size(); ++i) {
auto lhs = out_data[i];
auto rhs = (TypeOut)(op(lhs_h, rhs_data[i]));
ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs);
}
}

if (rhs.nullable()) {
Expand Down Expand Up @@ -138,10 +140,12 @@ void ASSERT_BINOP(column_view const& out,
auto out_data = out_h.first;

ASSERT_EQ(out_data.size(), lhs_data.size());
for (size_t i = 0; i < out_data.size(); ++i) {
auto lhs = out_data[i];
auto rhs = (TypeOut)(op(lhs_data[i], rhs_h));
ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs);
if(rhs.is_valid()) {
for (size_t i = 0; i < out_data.size(); ++i) {
auto lhs = out_data[i];
auto rhs = (TypeOut)(op(lhs_data[i], rhs_h));
ASSERT_TRUE(value_comparator(lhs, rhs)) << stringify_out_values{}(i, lhs, rhs);
}
}

if (lhs.nullable()) {
Expand Down