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

Use cuda::proclaim_return_type on device lambdas #1662

Merged
merged 15 commits into from
Dec 21, 2023
Merged

Use cuda::proclaim_return_type on device lambdas #1662

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
9076ad5
adding proclaim_return_type to device lambdas
hyperbolic2346 Dec 19, 2023
6784570
clang-format
hyperbolic2346 Dec 19, 2023
8fb716c
No cuda::proclaim_return_type on non-device lambda
hyperbolic2346 Dec 19, 2023
44d3720
Adding Mithun's changes for CCCL 2
hyperbolic2346 Dec 19, 2023
32f65e2
Merge remote-tracking branch 'upstream/branch-24.02' into mwilson/pro…
hyperbolic2346 Dec 20, 2023
7d12f40
linting
hyperbolic2346 Dec 20, 2023
4691bfb
updating return type
hyperbolic2346 Dec 20, 2023
11a872e
Update src/main/cpp/CMakeLists.txt
hyperbolic2346 Dec 20, 2023
a87f3e4
Update jni
ttnghia Dec 21, 2023
3208e7d
Apply suggestions from code review
hyperbolic2346 Dec 21, 2023
8c6487b
Fix styles
ttnghia Dec 21, 2023
380c218
linting
hyperbolic2346 Dec 21, 2023
d8b62d7
Update submodule manually
ttnghia Dec 21, 2023
ae3561a
Fix header
ttnghia Dec 21, 2023
348299a
Merge branch 'remove_row_conversion' into mwilson/proclaim_return_type
ttnghia Dec 21, 2023
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
7 changes: 2 additions & 5 deletions src/main/cpp/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,11 +94,8 @@ include(cmake/Modules/ConfigureCUDA.cmake) # set other CUDA compilation flags
# ##################################################################################################
# * dependencies ----------------------------------------------------------------------------------

# find libcu++
include(${rapids-cmake-dir}/cpm/libcudacxx.cmake)

# find thrust/cub
include(${CUDF_DIR}/cpp/cmake/thirdparty/get_thrust.cmake)
# find CCCL
include(${CUDF_DIR}/cpp/cmake/thirdparty/get_cccl.cmake)
ttnghia marked this conversation as resolved.
Show resolved Hide resolved

# JNI
find_package(JNI REQUIRED)
Expand Down
19 changes: 11 additions & 8 deletions src/main/cpp/src/bloom_filter.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,8 @@

#include <thrust/logical.h>

#include <cuda/functional>

#include <byteswap.h>

namespace spark_rapids_jni {
Expand Down Expand Up @@ -316,14 +318,15 @@ std::unique_ptr<cudf::list_scalar> bloom_filter_merge(cudf::column_view const& b
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(0) + num_words,
dst,
[src, num_buffers = bloom_filters.size(), stride = buf_size] __device__(
cudf::size_type word_index) {
cudf::bitmask_type out = (reinterpret_cast<cudf::bitmask_type const*>(src))[word_index];
for (auto idx = 1; idx < num_buffers; idx++) {
out |= (reinterpret_cast<cudf::bitmask_type const*>(src + idx * stride))[word_index];
}
return out;
});
cuda::proclaim_return_type<cudf::bitmask_type>(
[src, num_buffers = bloom_filters.size(), stride = buf_size] __device__(
cudf::size_type word_index) {
cudf::bitmask_type out = (reinterpret_cast<cudf::bitmask_type const*>(src))[word_index];
for (auto idx = 1; idx < num_buffers; idx++) {
out |= (reinterpret_cast<cudf::bitmask_type const*>(src + idx * stride))[word_index];
}
return out;
}));

// create the 1-row list column and move it into a scalar.
return std::make_unique<cudf::list_scalar>(
Expand Down
192 changes: 99 additions & 93 deletions src/main/cpp/src/datetime_rebase.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,8 @@
#include <thrust/iterator/counting_iterator.h>
#include <thrust/transform.h>

#include <cuda/functional>

namespace {

// Convert a date in Julian calendar to the number of days since epoch.
Expand Down Expand Up @@ -73,28 +75,29 @@ std::unique_ptr<cudf::column> gregorian_to_julian_days(cudf::column_view const&
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
output->mutable_view().begin<cudf::timestamp_D>(),
[d_input = input.begin<cudf::timestamp_D>()] __device__(auto const idx) {
auto constexpr julian_end = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{4}};
auto constexpr gregorian_start = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{15}};

auto const days_ts = d_input[idx].time_since_epoch().count();
auto const days_since_epoch = cuda::std::chrono::sys_days(cudf::duration_D{days_ts});

// Convert the input into local date in Proleptic Gregorian calendar.
auto const ymd = cuda::std::chrono::year_month_day(days_since_epoch);
if (ymd > julian_end && ymd < gregorian_start) {
// This is the same as rebasing from the local date given at `gregorian_start`.
return cudf::timestamp_D{cudf::duration_D{-141427}};
}

// No change since this time.
if (ymd >= gregorian_start) { return d_input[idx]; }

// Reinterpret year/month/day as in Julian calendar then compute the days since epoch.
return cudf::timestamp_D{cudf::duration_D{days_from_julian(ymd)}};
});
cuda::proclaim_return_type<cudf::timestamp_D>(
[d_input = input.begin<cudf::timestamp_D>()] __device__(auto const idx) {
auto constexpr julian_end = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{4}};
auto constexpr gregorian_start = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{15}};

auto const days_ts = d_input[idx].time_since_epoch().count();
auto const days_since_epoch = cuda::std::chrono::sys_days(cudf::duration_D{days_ts});

// Convert the input into local date in Proleptic Gregorian calendar.
auto const ymd = cuda::std::chrono::year_month_day(days_since_epoch);
if (ymd > julian_end && ymd < gregorian_start) {
// This is the same as rebasing from the local date given at `gregorian_start`.
return cudf::timestamp_D{cudf::duration_D{-141427}};
}

// No change since this time.
if (ymd >= gregorian_start) { return d_input[idx]; }

// Reinterpret year/month/day as in Julian calendar then compute the days since epoch.
return cudf::timestamp_D{cudf::duration_D{days_from_julian(ymd)}};
}));

return output;
}
Expand Down Expand Up @@ -142,19 +145,20 @@ std::unique_ptr<cudf::column> julian_to_gregorian_days(cudf::column_view const&
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
output->mutable_view().begin<cudf::timestamp_D>(),
[d_input = input.begin<cudf::timestamp_D>()] __device__(auto const idx) {
auto const days_ts = d_input[idx].time_since_epoch().count();
if (days_ts >= -141427) { // Gregorian start day
return d_input[idx];
}

// Reinterpret year/month/day as in Gregorian calendar then compute the days
// since epoch.
auto const ymd = julian_from_days(days_ts);
auto const result =
cuda::std::chrono::local_days{ymd}.time_since_epoch().count();
return cudf::timestamp_D{cudf::duration_D{result}};
});
cuda::proclaim_return_type<cudf::timestamp_D>(
[d_input = input.begin<cudf::timestamp_D>()] __device__(auto const idx) {
auto const days_ts = d_input[idx].time_since_epoch().count();
if (days_ts >= -141427) { // Gregorian start day
return d_input[idx];
}

// Reinterpret year/month/day as in Gregorian calendar then compute the days
// since epoch.
auto const ymd = julian_from_days(days_ts);
auto const result =
cuda::std::chrono::local_days{ymd}.time_since_epoch().count();
return cudf::timestamp_D{cudf::duration_D{result}};
}));

return output;
}
Expand Down Expand Up @@ -242,39 +246,40 @@ std::unique_ptr<cudf::column> gregorian_to_julian_micros(cudf::column_view const
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
output->mutable_view().begin<cudf::timestamp_us>(),
[d_input = input.begin<cudf::timestamp_us>()] __device__(auto const idx) {
// This timestamp corresponds to October 15th, 1582 UTC.
// After this day, there is no difference in microsecond values between Gregorian
// and Julian calendars.
int64_t constexpr last_switch_gregorian_ts = -12219292800000000L;

auto const micros_ts = d_input[idx].time_since_epoch().count();
if (micros_ts >= last_switch_gregorian_ts) { return d_input[idx]; }

// Convert the input into local date-time in Proleptic Gregorian calendar.
auto const days_since_epoch = cuda::std::chrono::sys_days(static_cast<cudf::duration_D>(
cuda::std::chrono::floor<cuda::std::chrono::days>(cudf::duration_us(micros_ts))));
auto const ymd = cuda::std::chrono::year_month_day(days_since_epoch);
auto const timeparts = get_time_components(micros_ts);

auto constexpr julian_end = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{4}};
auto constexpr gregorian_start = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{15}};

// Reinterpret the local date-time as in Julian calendar and compute microseconds since
// the epoch from that Julian local date-time.
// If the input date is outside of both calendars, consider it as it is a local date
// given at `gregorian_start` (-141427 Julian days since epoch).
auto const julian_days =
(ymd > julian_end && ymd < gregorian_start) ? -141427 : days_from_julian(ymd);
int64_t result = (julian_days * 24L * 3600L) + (timeparts.hour * 3600L) +
(timeparts.minute * 60L) + timeparts.second;
result *= MICROS_PER_SECOND; // to microseconds
result += timeparts.subsecond;

return cudf::timestamp_us{cudf::duration_us{result}};
});
cuda::proclaim_return_type<cudf::timestamp_us>(
[d_input = input.begin<cudf::timestamp_us>()] __device__(auto const idx) {
// This timestamp corresponds to October 15th, 1582 UTC.
// After this day, there is no difference in microsecond values between Gregorian
// and Julian calendars.
int64_t constexpr last_switch_gregorian_ts = -12219292800000000L;

auto const micros_ts = d_input[idx].time_since_epoch().count();
if (micros_ts >= last_switch_gregorian_ts) { return d_input[idx]; }

// Convert the input into local date-time in Proleptic Gregorian calendar.
auto const days_since_epoch = cuda::std::chrono::sys_days(static_cast<cudf::duration_D>(
cuda::std::chrono::floor<cuda::std::chrono::days>(cudf::duration_us(micros_ts))));
auto const ymd = cuda::std::chrono::year_month_day(days_since_epoch);
auto const timeparts = get_time_components(micros_ts);

auto constexpr julian_end = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{4}};
auto constexpr gregorian_start = cuda::std::chrono::year_month_day{
cuda::std::chrono::year{1582}, cuda::std::chrono::month{10}, cuda::std::chrono::day{15}};

// Reinterpret the local date-time as in Julian calendar and compute microseconds since
// the epoch from that Julian local date-time.
// If the input date is outside of both calendars, consider it as it is a local date
// given at `gregorian_start` (-141427 Julian days since epoch).
auto const julian_days =
(ymd > julian_end && ymd < gregorian_start) ? -141427 : days_from_julian(ymd);
int64_t result = (julian_days * 24L * 3600L) + (timeparts.hour * 3600L) +
(timeparts.minute * 60L) + timeparts.second;
result *= MICROS_PER_SECOND; // to microseconds
result += timeparts.subsecond;

return cudf::timestamp_us{cudf::duration_us{result}};
}));

return output;
}
Expand Down Expand Up @@ -304,31 +309,32 @@ std::unique_ptr<cudf::column> julian_to_gregorian_micros(cudf::column_view const
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
output->mutable_view().begin<cudf::timestamp_us>(),
[d_input = input.begin<cudf::timestamp_us>()] __device__(auto const idx) {
// This timestamp corresponds to October 15th, 1582 UTC.
// After this day, there is no difference in microsecond values between Gregorian
// and Julian calendars.
int64_t constexpr last_switch_gregorian_ts = -12219292800000000L;

auto const micros_ts = d_input[idx].time_since_epoch().count();
if (micros_ts >= last_switch_gregorian_ts) { return d_input[idx]; }

// Convert the input into local date-time in Julian calendar.
auto const days_since_epoch = cuda::std::chrono::sys_days(static_cast<cudf::duration_D>(
cuda::std::chrono::floor<cuda::std::chrono::days>(cudf::duration_us(micros_ts))));
auto const ymd = julian_from_days(days_since_epoch.time_since_epoch().count());
auto const timeparts = get_time_components(micros_ts);

// Reinterpret the local date-time as in Gregorian calendar and compute microseconds since
// the epoch from that Gregorian local date-time.
auto const gregorian_days = cuda::std::chrono::local_days(ymd).time_since_epoch().count();
int64_t result = (gregorian_days * 24L * 3600L) + (timeparts.hour * 3600L) +
(timeparts.minute * 60L) + timeparts.second;
result *= MICROS_PER_SECOND; // to microseconds
result += timeparts.subsecond;

return cudf::timestamp_us{cudf::duration_us{result}};
});
cuda::proclaim_return_type<cudf::timestamp_us>(
[d_input = input.begin<cudf::timestamp_us>()] __device__(auto const idx) {
// This timestamp corresponds to October 15th, 1582 UTC.
// After this day, there is no difference in microsecond values between Gregorian
// and Julian calendars.
int64_t constexpr last_switch_gregorian_ts = -12219292800000000L;

auto const micros_ts = d_input[idx].time_since_epoch().count();
if (micros_ts >= last_switch_gregorian_ts) { return d_input[idx]; }

// Convert the input into local date-time in Julian calendar.
auto const days_since_epoch = cuda::std::chrono::sys_days(static_cast<cudf::duration_D>(
cuda::std::chrono::floor<cuda::std::chrono::days>(cudf::duration_us(micros_ts))));
auto const ymd = julian_from_days(days_since_epoch.time_since_epoch().count());
auto const timeparts = get_time_components(micros_ts);

// Reinterpret the local date-time as in Gregorian calendar and compute microseconds since
// the epoch from that Gregorian local date-time.
auto const gregorian_days = cuda::std::chrono::local_days(ymd).time_since_epoch().count();
int64_t result = (gregorian_days * 24L * 3600L) + (timeparts.hour * 3600L) +
(timeparts.minute * 60L) + timeparts.second;
result *= MICROS_PER_SECOND; // to microseconds
result += timeparts.subsecond;

return cudf::timestamp_us{cudf::duration_us{result}};
}));

return output;
}
Expand Down
Loading