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.

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

Add support for org.apache.spark.sql.catalyst.expressions.Bin #2760

Open
wants to merge 6 commits into
base: branch-25.02
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from 2 commits
Commits
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
1 change: 1 addition & 0 deletions src/main/cpp/CMakeLists.txt
ustcfy marked this conversation as resolved.
Show resolved Hide resolved
Original file line number Diff line number Diff line change
Expand Up @@ -211,6 +211,7 @@ add_library(
src/case_when.cu
src/cast_decimal_to_string.cu
src/cast_float_to_string.cu
src/cast_long_to_binary_string.cu
src/cast_string.cu
src/cast_string_to_float.cu
src/datetime_rebase.cu
Expand Down
17 changes: 16 additions & 1 deletion src/main/cpp/src/CastStringJni.cpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -156,6 +156,21 @@ JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_CastStrings_fromDecimal
CATCH_CAST_EXCEPTION(env, 0);
}

JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_CastStrings_fromLongToBinary(
JNIEnv* env, jclass, jlong input_column)
{
JNI_NULL_CHECK(env, input_column, "input column is null", 0);

try {
cudf::jni::auto_set_device(env);

auto const& cv = *reinterpret_cast<cudf::column_view const*>(input_column);
return cudf::jni::release_as_jlong(
spark_rapids_jni::long_to_binary_string(cv, cudf::get_default_stream()));
}
CATCH_CAST_EXCEPTION(env, 0);
}

JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_CastStrings_toIntegersWithBase(
JNIEnv* env, jclass, jlong input_column, jint base, jboolean ansi_enabled, jint j_dtype)
{
Expand Down
120 changes: 120 additions & 0 deletions src/main/cpp/src/cast_long_to_binary_string.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,120 @@
/*
* Copyright (c) 2025, 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 "cast_string.hpp"

#include <cudf/column/column_device_view.cuh>
#include <cudf/column/column_factories.hpp>
#include <cudf/detail/null_mask.hpp>
#include <cudf/detail/nvtx/ranges.hpp>
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Useless include (not sure)?

#include <cudf/strings/detail/strings_children.cuh>
#include <cudf/utilities/default_stream.hpp>
#include <cudf/utilities/type_dispatcher.hpp>

#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

namespace spark_rapids_jni {

namespace detail {
namespace {

template <typename LongType>
ustcfy marked this conversation as resolved.
Show resolved Hide resolved
struct long_to_binary_string_fn {
cudf::column_device_view d_longs;
cudf::size_type* d_sizes;
char* d_chars;
cudf::detail::input_offsetalator d_offsets;

__device__ cudf::size_type compute_output_size(LongType value)
{
auto const size = 64 - __clzll(value);
// If the value is 0, the size should be 1
return size > 0 ? size : 1;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: we can first check if the value == 0 to save a __clzll call in this case.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But wouldn't that introduce thread divergence? I understand how that can be an advantage on a CPU, but I don't really see it on a GPU.

Copy link
Collaborator Author

@ustcfy ustcfy Jan 14, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The PTX translated from the code is as follows:

auto const size = 64 - __clzll(value);
return size > 0 ? size : 1;

// PTX
clz.b64         %r1, %rd1;           // %r1: number of leading zeros in %rd1 (value)
mov.u32         %r2, 64;             // %r2: constant value 64
sub.s32         %r3, %r2, %r1;       // %r3: size = 64 - %r1
max.s32         %r4, %r3, 1;         // %r4: return max(size, 1)
if (!value) return 1;
return 64 - __clzll(value);

// PTX
setp.eq.s64     %p1, %rd1, 0;        // %p1: set to true if %rd1 (value) is zero
clz.b64         %r1, %rd1;           // %r1: number of leading zeros in %rd1 (value)
mov.u32         %r2, 64;             // %r2: constant value 64
sub.s32         %r3, %r2, %r1;       // %r3: size = 64 - %r1
selp.b32        %r4, 1, %r3, %p1;    // %r4: if %p1 is true, return 1; otherwise, return %r3 (final size)

Copy link
Collaborator

@thirtiseven thirtiseven Jan 15, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But wouldn't that introduce thread divergence? I understand how that can be an advantage on a CPU, but I don't really see it on a GPU.

Yes, I was thinking there must be an if else to check size == 0 case so we can put it earlier than clz to save some calls without introducing new branch. But it looks like the compiler will optimize the size > 0 ? size : 1; to a max.s32 so it's branch less then the original approach looks better in anyway.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have now changed it to this: return max(64 - __clzll(value), 1);

}

__device__ void long_to_binary_string(cudf::size_type idx)
{
auto const value = d_longs.element<LongType>(idx);
char* d_buffer = d_chars + d_offsets[idx];
for (auto i = d_sizes[idx] - 1; i >= 0; --i) {
*d_buffer++ = value & (1LL << i) ? '1' : '0';
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

*d_buffer++ = '0' + ((value & (1LL << i)) >> i); perhaps this approach is more efficient since it avoids branching, which might degrade performance on GPUs with warp divergence.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But I am not sure if it is a good practice which is really effective. I would like to hear your opinions on this issue @res-life @ttnghia .

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe *d_buffer++ = '0' + ((value & (1LL << i)) != 0);? It will be (very slightly) cheaper and easier to read.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@thirtiseven
Yes, this one is also a Branch-Free expression since the compiler shall use setne instruction avoids branching by directly setting a register based on the zero flag (ZF):

cmp rax, 0
setne al
and al, 1
add eax, 48

The corresponding codes of my alternative would be translated into:

sar     rax, cl
add     rax, 48

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But I am not sure if it is a good practice which is really effective. I would like to hear your opinions on this issue @res-life @ttnghia .

Yes, I think this approach is more efficient.
You may conduct a benchmark test to double confirm.

}
}

__device__ void operator()(cudf::size_type idx)
{
if (d_longs.is_null(idx)) {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

NIT: Just some nice-to-have improvement, use constexpr if instead of if and add an extra template variable nullable for this functor. Because we already knew whether the column_view is nullable or NOT.

if (d_chars == nullptr) { d_sizes[idx] = 0; }
return;
}
if (d_chars != nullptr) {
long_to_binary_string(idx);
} else {
d_sizes[idx] = compute_output_size(d_longs.element<LongType>(idx));
}
}
};

struct dispatch_long_to_binary_string_fn {
template <typename LongType, CUDF_ENABLE_IF(std::is_same_v<LongType, std::int64_t>)>
std::unique_ptr<cudf::column> operator()(cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr) const
{
auto const d_column = cudf::column_device_view::create(input, stream);

auto [offsets, chars] = cudf::strings::detail::make_strings_children(
long_to_binary_string_fn<LongType>{*d_column}, input.size(), stream, mr);

return cudf::make_strings_column(input.size(),
std::move(offsets),
chars.release(),
input.null_count(),
cudf::detail::copy_bitmask(input, stream, mr));
}

template <typename LongType, CUDF_ENABLE_IF(not std::is_same_v<LongType, std::int64_t>)>
std::unique_ptr<cudf::column> operator()(cudf::column_view const&,
rmm::cuda_stream_view,
rmm::device_async_resource_ref) const
{
CUDF_FAIL("Values for long_to_binary_string function must be a long type.");
}
};

} // namespace

std::unique_ptr<cudf::column> long_to_binary_string(cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
{
if (input.is_empty()) return cudf::make_empty_column(cudf::type_id::STRING);
return type_dispatcher(input.type(), dispatch_long_to_binary_string_fn{}, input, stream, mr);
}

} // namespace detail

// external API
std::unique_ptr<cudf::column> long_to_binary_string(cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
{
CUDF_FUNC_RANGE();
ustcfy marked this conversation as resolved.
Show resolved Hide resolved
return detail::long_to_binary_string(input, stream, mr);
}

} // namespace spark_rapids_jni
7 changes: 6 additions & 1 deletion src/main/cpp/src/cast_string.hpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -133,4 +133,9 @@ std::unique_ptr<cudf::column> decimal_to_non_ansi_string(
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

std::unique_ptr<cudf::column> long_to_binary_string(
cudf::column_view const& input,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

} // namespace spark_rapids_jni
5 changes: 4 additions & 1 deletion src/main/cpp/tests/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#=============================================================================
# Copyright (c) 2022-2024, NVIDIA CORPORATION.
# Copyright (c) 2022-2025, NVIDIA CORPORATION.
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -58,6 +58,9 @@ ConfigureTest(FORMAT_FLOAT
ConfigureTest(CAST_FLOAT_TO_STRING
cast_float_to_string.cpp)

ConfigureTest(CAST_LONG_TO_BINARY_STRING
cast_long_to_binary_string.cpp)

ConfigureTest(DATETIME_REBASE
datetime_rebase.cpp)

Expand Down
42 changes: 42 additions & 0 deletions src/main/cpp/tests/cast_long_to_binary_string.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
/*
* Copyright (c) 2025, 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 <cudf_test/base_fixture.hpp>
#include <cudf_test/column_wrapper.hpp>

#include <rmm/device_uvector.hpp>

#include <cast_string.hpp>

#include <limits>

using namespace cudf;

constexpr cudf::test::debug_output_level verbosity{cudf::test::debug_output_level::FIRST_ERROR};

struct LongToBinaryStringTests : public cudf::test::BaseFixture {};

TEST_F(LongToBinaryStringTests, FromLongToBinary)
{
auto const longs = cudf::test::fixed_width_column_wrapper<int64_t>{0L, 1L, 10L, -0L, -1L};
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could we add some edge case like null, LONG_MAX, LONG_MIN?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do we have a test case at plugin level to make sure Bin(13.3) returns 1101 ?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do we have a test case at plugin level to make sure Bin(13.3) returns 1101 ?

I tested it locally, and Bin(13.3) indeed returns 1101. I will soon submit the plugin PR.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is really a Spark test not a good test for binary ops. Spark only accepts a Long as the input to bin

https://github.com/apache/spark/blob/3569e768e657d4e28ee7520808ec910cdff2b099/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/expressions/mathExpressions.scala#L1010

So any floating point input gets a cast to long inserted in before bin is called. So that test is really a test that Spark is doing the right thing. Even then it would probably be something that we would want to put in the integration tests if we did test it at all.

Note that you can also pass in a string as an input and it will still try to cast it to a long before calling bin.


auto results = spark_rapids_jni::long_to_binary_string(longs, cudf::get_default_stream());

auto const expected = cudf::test::strings_column_wrapper{
"0", "1", "1010", "0", "1111111111111111111111111111111111111111111111111111111111111111"};
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.


CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected, verbosity);
}
7 changes: 6 additions & 1 deletion src/main/java/com/nvidia/spark/rapids/jni/CastStrings.java
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2023, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -111,6 +111,10 @@ public static ColumnVector fromDecimal(ColumnView cv) {
return new ColumnVector(fromDecimal(cv.getNativeView()));
}

public static ColumnVector fromLongToBinary(ColumnView cv) {
return new ColumnVector(fromLongToBinary(cv.getNativeView()));
}

/**
* Convert a string column to a given floating-point type column.
*
Expand Down Expand Up @@ -160,6 +164,7 @@ private static native long toDecimal(long nativeColumnView, boolean ansi_enabled
private static native long fromDecimal(long nativeColumnView);
private static native long fromFloatWithFormat(long nativeColumnView, int digits);
private static native long fromFloat(long nativeColumnView);
private static native long fromLongToBinary(long nativeColumnView);
private static native long toIntegersWithBase(long nativeColumnView, int base,
boolean ansiEnabled, int dtype);
private static native long fromIntegersWithBase(long nativeColumnView, int base);
Expand Down
Loading