Skip to content

Commit

Permalink
Pre-commit clang-format
Browse files Browse the repository at this point in the history 
Signed-off-by: Navin Kumar <navink@nvidia.com>
  • Loading branch information
@NVnavkumar
NVnavkumar committed Nov 21, 2023
1 parent c3bfa14 commit a165262
Show file tree
Hide file tree
Showing 4 changed files with 265 additions and 266 deletions.
22 changes: 12 additions & 10 deletions src/main/cpp/src/GpuTimeZoneDBJni.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,32 +19,34 @@
extern "C" {

JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_GpuTimeZoneDB_convertTimestampColumnToUTC(
JNIEnv *env, jclass, jlong input_handle, jlong transitions_handle, jint tz_index) {
JNIEnv* env, jclass, jlong input_handle, jlong transitions_handle, jint tz_index)
{
JNI_NULL_CHECK(env, input_handle, "column is null", 0);
JNI_NULL_CHECK(env, transitions_handle, "column is null", 0);
try {
cudf::jni::auto_set_device(env);
auto const input = reinterpret_cast<cudf::column_view const *>(input_handle);
auto const transitions = reinterpret_cast<cudf::table_view const *>(transitions_handle);
auto const index = static_cast<cudf::size_type>(tz_index);
auto const input = reinterpret_cast<cudf::column_view const*>(input_handle);
auto const transitions = reinterpret_cast<cudf::table_view const*>(transitions_handle);
auto const index = static_cast<cudf::size_type>(tz_index);
return cudf::jni::ptr_as_jlong(
spark_rapids_jni::convert_timestamp_to_utc(*input, *transitions, index).release());
spark_rapids_jni::convert_timestamp_to_utc(*input, *transitions, index).release());
}
CATCH_STD(env, 0);
}

JNIEXPORT jlong JNICALL
Java_com_nvidia_spark_rapids_jni_GpuTimeZoneDB_convertUTCTimestampColumnToTimeZone(
JNIEnv *env, jclass, jlong input_handle, jlong transitions_handle, jint tz_index) {
JNIEnv* env, jclass, jlong input_handle, jlong transitions_handle, jint tz_index)
{
JNI_NULL_CHECK(env, input_handle, "column is null", 0);
JNI_NULL_CHECK(env, transitions_handle, "column is null", 0);
try {
cudf::jni::auto_set_device(env);
auto const input = reinterpret_cast<cudf::column_view const *>(input_handle);
auto const transitions = reinterpret_cast<cudf::table_view const *>(transitions_handle);
auto const index = static_cast<cudf::size_type>(tz_index);
auto const input = reinterpret_cast<cudf::column_view const*>(input_handle);
auto const transitions = reinterpret_cast<cudf::table_view const*>(transitions_handle);
auto const index = static_cast<cudf::size_type>(tz_index);
return cudf::jni::ptr_as_jlong(
spark_rapids_jni::convert_utc_timestamp_to_timezone(*input, *transitions, index).release());
spark_rapids_jni::convert_utc_timestamp_to_timezone(*input, *transitions, index).release());
}
CATCH_STD(env, 0);
}
Expand Down
107 changes: 62 additions & 45 deletions src/main/cpp/src/timezones.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -30,13 +30,13 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

using column = cudf::column;
using column_device_view = cudf::column_device_view;
using column_view = cudf::column_view;
using column = cudf::column;
using column_device_view = cudf::column_device_view;
using column_view = cudf::column_view;
using lists_column_device_view = cudf::detail::lists_column_device_view;
using size_type = cudf::size_type;
using struct_view = cudf::struct_view;
using table_view = cudf::table_view;
using size_type = cudf::size_type;
using struct_view = cudf::struct_view;
using table_view = cudf::table_view;

namespace {

Expand All @@ -46,7 +46,8 @@ namespace {
// the offset.
// To transition from UTC: do a binary search on the utcInstant child column and add
// the offset.
template <typename timestamp_type> struct convert_timestamp_tz_functor {
template <typename timestamp_type>
struct convert_timestamp_tz_functor {
using duration_type = typename timestamp_type::duration;

// The list column of transitions to figure out the correct offset
Expand All @@ -63,87 +64,103 @@ template <typename timestamp_type> struct convert_timestamp_tz_functor {
* @param timestamp input timestamp
*
*/
__device__ timestamp_type operator()(timestamp_type const &timestamp) const {
__device__ timestamp_type operator()(timestamp_type const& timestamp) const
{
auto const utc_instants = transitions.child().child(0);
auto const tz_instants = transitions.child().child(1);
auto const utc_offsets = transitions.child().child(2);
auto const tz_instants = transitions.child().child(1);
auto const utc_offsets = transitions.child().child(2);

auto const epoch_seconds = static_cast<int64_t>(
cuda::std::chrono::duration_cast<cudf::duration_s>(timestamp.time_since_epoch()).count());
cuda::std::chrono::duration_cast<cudf::duration_s>(timestamp.time_since_epoch()).count());
auto const tz_transitions = cudf::list_device_view{transitions, tz_index};
auto const list_size = tz_transitions.size();
auto const list_size = tz_transitions.size();

auto const transition_times = cudf::device_span<int64_t const>(
(to_utc ? tz_instants : utc_instants).data<int64_t>() + tz_transitions.element_offset(0),
static_cast<size_t>(list_size));
(to_utc ? tz_instants : utc_instants).data<int64_t>() + tz_transitions.element_offset(0),
static_cast<size_t>(list_size));

auto const it = thrust::upper_bound(thrust::seq, transition_times.begin(),
transition_times.end(), epoch_seconds);
auto const idx = static_cast<size_type>(thrust::distance(transition_times.begin(), it));
auto const it = thrust::upper_bound(
thrust::seq, transition_times.begin(), transition_times.end(), epoch_seconds);
auto const idx = static_cast<size_type>(thrust::distance(transition_times.begin(), it));
auto const list_offset = tz_transitions.element_offset(idx - 1);
auto const utc_offset = cuda::std::chrono::duration_cast<duration_type>(
cudf::duration_s{static_cast<int64_t>(utc_offsets.element<int32_t>(list_offset))});
auto const utc_offset = cuda::std::chrono::duration_cast<duration_type>(
cudf::duration_s{static_cast<int64_t>(utc_offsets.element<int32_t>(list_offset))});
return to_utc ? timestamp - utc_offset : timestamp + utc_offset;
}
};

template <typename timestamp_type>
auto convert_timestamp_tz(column_view const &input, table_view const &transitions,
size_type tz_index, bool to_utc, rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource *mr) {
auto convert_timestamp_tz(column_view const& input,
table_view const& transitions,
size_type tz_index,
bool to_utc,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
// get the fixed transitions
auto const ft_cdv_ptr = column_device_view::create(transitions.column(0), stream);
auto const ft_cdv_ptr = column_device_view::create(transitions.column(0), stream);
auto const fixed_transitions = lists_column_device_view{*ft_cdv_ptr};

auto results = cudf::make_timestamp_column(input.type(), input.size(),
auto results = cudf::make_timestamp_column(input.type(),
input.size(),
cudf::detail::copy_bitmask(input, stream, mr),
input.null_count(), stream, mr);
input.null_count(),
stream,
mr);

thrust::transform(
rmm::exec_policy(stream), input.begin<timestamp_type>(), input.end<timestamp_type>(),
results->mutable_view().begin<timestamp_type>(),
convert_timestamp_tz_functor<timestamp_type>{fixed_transitions, tz_index, to_utc});
rmm::exec_policy(stream),
input.begin<timestamp_type>(),
input.end<timestamp_type>(),
results->mutable_view().begin<timestamp_type>(),
convert_timestamp_tz_functor<timestamp_type>{fixed_transitions, tz_index, to_utc});

return results;
}

} // namespace
} // namespace

namespace spark_rapids_jni {

std::unique_ptr<column> convert_timestamp(column_view const &input, table_view const &transitions,
size_type tz_index, bool to_utc,
std::unique_ptr<column> convert_timestamp(column_view const& input,
table_view const& transitions,
size_type tz_index,
bool to_utc,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource *mr) {
rmm::mr::device_memory_resource* mr)
{
auto const type = input.type().id();

switch (type) {
case cudf::type_id::TIMESTAMP_SECONDS:
return convert_timestamp_tz<cudf::timestamp_s>(input, transitions, tz_index, to_utc, stream,
mr);
return convert_timestamp_tz<cudf::timestamp_s>(
input, transitions, tz_index, to_utc, stream, mr);
case cudf::type_id::TIMESTAMP_MILLISECONDS:
return convert_timestamp_tz<cudf::timestamp_ms>(input, transitions, tz_index, to_utc, stream,
mr);
return convert_timestamp_tz<cudf::timestamp_ms>(
input, transitions, tz_index, to_utc, stream, mr);
case cudf::type_id::TIMESTAMP_MICROSECONDS:
return convert_timestamp_tz<cudf::timestamp_us>(input, transitions, tz_index, to_utc, stream,
mr);
return convert_timestamp_tz<cudf::timestamp_us>(
input, transitions, tz_index, to_utc, stream, mr);
default: CUDF_FAIL("Unsupported timestamp unit for timezone conversion");
}
}

std::unique_ptr<column> convert_timestamp_to_utc(column_view const &input,
table_view const &transitions, size_type tz_index,
std::unique_ptr<column> convert_timestamp_to_utc(column_view const& input,
table_view const& transitions,
size_type tz_index,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource *mr) {
rmm::mr::device_memory_resource* mr)
{
return convert_timestamp(input, transitions, tz_index, true, stream, mr);
}

std::unique_ptr<column> convert_utc_timestamp_to_timezone(column_view const &input,
table_view const &transitions,
std::unique_ptr<column> convert_utc_timestamp_to_timezone(column_view const& input,
table_view const& transitions,
size_type tz_index,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource *mr) {
rmm::mr::device_memory_resource* mr)
{
return convert_timestamp(input, transitions, tz_index, false, stream, mr);
}

} // namespace spark_rapids_jni
} // namespace spark_rapids_jni
19 changes: 11 additions & 8 deletions src/main/cpp/src/timezones.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,6 @@

#include <cstddef>


namespace spark_rapids_jni {

/**
Expand All @@ -40,9 +39,11 @@ namespace spark_rapids_jni {
* @param mr Device memory resource used to allocate the returned timestamp column's memory
*/
std::unique_ptr<cudf::column> convert_timestamp_to_utc(
cudf::column_view const &input, cudf::table_view const &transitions, cudf::size_type tz_index,
rmm::cuda_stream_view stream = cudf::get_default_stream(),
rmm::mr::device_memory_resource *mr = rmm::mr::get_current_device_resource());
cudf::column_view const& input,
cudf::table_view const& transitions,
cudf::size_type tz_index,
rmm::cuda_stream_view stream = cudf::get_default_stream(),
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
* @brief Convert input column timestamps in UTC to specified timezone
Expand All @@ -59,8 +60,10 @@ std::unique_ptr<cudf::column> convert_timestamp_to_utc(
* @param mr Device memory resource used to allocate the returned timestamp column's memory
*/
std::unique_ptr<cudf::column> convert_utc_timestamp_to_timezone(
cudf::column_view const &input, cudf::table_view const &transitions, cudf::size_type tz_index,
rmm::cuda_stream_view stream = cudf::get_default_stream(),
rmm::mr::device_memory_resource *mr = rmm::mr::get_current_device_resource());
cudf::column_view const& input,
cudf::table_view const& transitions,
cudf::size_type tz_index,
rmm::cuda_stream_view stream = cudf::get_default_stream(),
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

} // namespace spark_rapids_jni
} // namespace spark_rapids_jni
Loading

0 comments on commit a165262

Please sign in to comment.