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

[REVIEW] duration type unary casts and binary ops #5394

Merged
merged 78 commits into from
Jul 22, 2020
Merged
Show file tree
Hide file tree
Changes from 54 commits
Commits
Show all changes
78 commits
Select commit Hold shift + click to select a range
2053a09
Merge branch 'fea-duration_types' of https://github.com/karthikeyann/…
sriramch Jun 4, 2020
be19e6b
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 4, 2020
b9b3c39
- support for duration types and unary casts
sriramch Jun 4, 2020
4f9a5fa
- changelog entries
sriramch Jun 4, 2020
989fe6c
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 5, 2020
97860f0
Merge branch 'fea-duration_types' of https://github.com/karthikeyann/…
sriramch Jun 5, 2020
ab108a7
- enable tests with duration types
sriramch Jun 8, 2020
49cc59f
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 8, 2020
8fca5d5
Merge branch 'fea-duration_types' of https://github.com/karthikeyann/…
sriramch Jun 8, 2020
805f538
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 8, 2020
b7c2781
- make the timestamp constructor explicit
sriramch Jun 9, 2020
aac5ebb
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 9, 2020
4f15a5a
- initial set of changes to support binop with chrono types
sriramch Jun 9, 2020
1d72ee4
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 9, 2020
d398736
- add column printer for duration types (to aid debugging)
sriramch Jun 9, 2020
3c57ab9
- remove timestamp constructors that can be inherited
sriramch Jun 10, 2020
1c9323a
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 10, 2020
3659135
- fix build issues
sriramch Jun 10, 2020
26efd47
- disallow conversion between numeric and timestamps
sriramch Jun 12, 2020
63a5d53
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 12, 2020
81a872f
- binops for duration/timestamp types
sriramch Jun 12, 2020
6dc6b89
- create factory methods for fixed with column wrapper type that can …
sriramch Jun 16, 2020
220b4b8
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 16, 2020
f596e55
- add a fixed width column wrapper constructor with input iterator an…
sriramch Jun 16, 2020
1bd2c5a
- make the tests use the fixed width column wrapper factory methods i…
sriramch Jun 18, 2020
9dbd5cd
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 18, 2020
de7e262
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 19, 2020
5ee73a0
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 20, 2020
05d6571
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 22, 2020
bfcec07
- localize create fixed width column wrapper in a single place
sriramch Jun 22, 2020
ce374cb
- incorporate review comments
sriramch Jun 22, 2020
57af9ea
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 22, 2020
8bdceab
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 23, 2020
e5d55b9
- instead of new factory methods that perform type conversions while …
sriramch Jun 23, 2020
b7d580a
- remove fixed width column wrapper factory methods, and explicitly p…
sriramch Jun 24, 2020
3fee810
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 24, 2020
c7b8778
- make source type to be int32
sriramch Jun 24, 2020
6969939
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 25, 2020
04029c9
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 25, 2020
c5e6f02
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 27, 2020
65544bd
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 29, 2020
ad8df2e
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 29, 2020
35df63a
- reserve host_vector, in case the types do not have a default constr…
sriramch Jun 29, 2020
3bc0c64
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jun 29, 2020
d16a95b
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 1, 2020
803e1cc
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 2, 2020
bed8374
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 3, 2020
0449a2a
- guard the specialization to is_trivially_copyable for duration type…
sriramch Jul 6, 2020
240041f
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 8, 2020
61e2101
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 9, 2020
73c1221
- incorporate review comments
sriramch Jul 9, 2020
3021455
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 9, 2020
4207233
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 10, 2020
43cff5e
- fix comment for duration type
sriramch Jul 10, 2020
183c95b
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 10, 2020
0b5b6a4
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 10, 2020
a5b7363
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 11, 2020
2732be4
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 13, 2020
83e2a35
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 13, 2020
16bdb08
- review comments
sriramch Jul 13, 2020
b78a57c
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 14, 2020
cfe0b95
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 14, 2020
1ae05be
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 14, 2020
351b7f3
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 15, 2020
8d760e9
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 15, 2020
a5062e2
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 15, 2020
0d2b0d2
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 16, 2020
2408894
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 16, 2020
a93f300
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 16, 2020
e180913
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 16, 2020
b936244
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 17, 2020
332e2d0
- remove specialization of `is_trivially_copyable`
sriramch Jul 17, 2020
dc9ec82
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 17, 2020
75ed6fe
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 18, 2020
d9ee191
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 20, 2020
6c57c59
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 21, 2020
ae0a31d
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 21, 2020
b7de3c9
Merge branch 'branch-0.15' of https://github.com/rapidsai/cudf into u…
sriramch Jul 21, 2020
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 CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
- PR #5356 Use `size_type` instead of `scalar` in `cudf::repeat`.
- PR #5397 Add internal implementation of nested loop equijoins.
- PR #5303 Add slice_strings functionality using delimiter string
- PR #5394 Enable cast and binops with duration types (builds on PR 5359)
- PR #5301 Add Java bindings for `zfill`
- PR #5359 Add duration types
- PR #5364 Validate array interface during buffer construction
Expand Down
4 changes: 4 additions & 0 deletions cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -621,6 +621,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit
${CMAKE_BINARY_DIR}/include/bit.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit
Expand All @@ -641,6 +642,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.hpp
${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/utilities/bit.hpp
${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/timestamps.hpp
${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/wrappers/durations.hpp
${LIBCUDACXX_INCLUDE_DIR}/details/__config
${LIBCUDACXX_INCLUDE_DIR}/simt/limits
${LIBCUDACXX_INCLUDE_DIR}/simt/cfloat
Expand All @@ -664,6 +666,7 @@ add_custom_command(WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/utilities/bit.hpp > ${CMAKE_BINARY_DIR}/include/bit.hpp.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify ../src/rolling/rolling_jit_detail.hpp > ${CMAKE_BINARY_DIR}/include/rolling_jit_detail.hpp.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/timestamps.hpp > ${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/wrappers/durations.hpp > ${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/details/__config libcudacxx_details_config > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/limits libcudacxx_simt_limits > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
COMMAND ${CMAKE_BINARY_DIR}/stringify ${LIBCUDACXX_INCLUDE_DIR}/simt/cfloat libcudacxx_simt_cfloat > ${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit
Expand All @@ -687,6 +690,7 @@ add_custom_target(stringify_run DEPENDS
${CMAKE_BINARY_DIR}/include/jit/types.hpp.jit
${CMAKE_BINARY_DIR}/include/bit.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/timestamps.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/durations.hpp.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/details/__config.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/limits.jit
${CMAKE_BINARY_DIR}/include/jit/libcudacxx/simt/cfloat.jit
Expand Down
57 changes: 53 additions & 4 deletions cpp/include/cudf/utilities/traits.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,8 @@

#include <type_traits>

#include <cuda_runtime_api.h>

namespace cudf {

/**
Expand Down Expand Up @@ -365,16 +367,39 @@ constexpr inline bool is_duration(data_type type)
}

/**
* @brief Indicates whether the type `T` is a cudf chrono type.
* @brief Indicates whether the type `T` is a chrono type.
*
* @tparam T The type to verify
* @return true `T` is a chrono type
* @return false `T` is not a chrono type
* @return true `T` is a duration or a timestamp type
* @return false `T` is neither a duration nor a timestamp type
**/
template <typename T>
constexpr inline bool is_chrono()
{
return cudf::is_timestamp<T>() || cudf::is_duration<T>();
return is_duration<T>() || is_timestamp<T>();
}

struct is_chrono_impl {
template <typename T>
bool operator()()
{
return is_chrono<T>();
}
};

/**
* @brief Indicates whether `type` is a chrono `data_type`.
*
* Chrono types include cudf timestamp types, which represent a point in time, and cudf
* duration types that represent a time interval.
*
* @param type The `data_type` to verify
* @return true `type` is a chrono type
* @return false `type` is not a chrono type
**/
constexpr inline bool is_chrono(data_type type)
{
return cudf::type_dispatcher(type, is_chrono_impl{});
}

/**
Expand Down Expand Up @@ -502,4 +527,28 @@ constexpr inline bool is_nested(data_type type)
}

/** @} */

template <typename From, typename To>
struct is_convertible : std::is_convertible<From, To> {
};

// This will ensure that timestamps can be promoted to a higher precision. Presently, they can't
// do that due to nvcc/gcc compiler issues
template <typename Duration1, typename Duration2>
struct is_convertible<cudf::detail::timestamp<Duration1>, cudf::detail::timestamp<Duration2>>
: std::is_convertible<typename cudf::detail::time_point<Duration1>::duration,
typename cudf::detail::time_point<Duration2>::duration> {
};

} // namespace cudf

#if (CUDART_VERSION == 10000)
namespace std {
// This specialization is required due to an nvcc bug. This works in 10.1.
// We can remove this once the build pipeline removes 10.0.
template <typename Duration>
struct is_trivially_copyable<cudf::detail::timestamp<Duration>>
Copy link
Contributor Author

@sriramch sriramch Jun 18, 2020

Choose a reason for hiding this comment

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

i'll remove this before the merge; i'm keeping this only to get the build pipeline to pass.

Copy link
Contributor

Choose a reason for hiding this comment

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

Is this left in so the CUDA 10.0 CI job passes? Are we expecting to remove CUDA 10.0 support before merging this PR? If not, should we guard this definition with an #ifdef so it's only include when the CUDA version is strictly 10.0?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

i have now guarded this such that this specialization is applicable only for 10.0 builds.

Copy link
Member

Choose a reason for hiding this comment

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

I would not take this out until we are well past 10.0, just in case someone chooses to keep using CUDA 10.0 "at their own risk". It's one thing to not support a CUDA version, but to knowingly block it is another...

Copy link
Contributor

Choose a reason for hiding this comment

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

I would not take this out until we are well past 10.0, just in case someone chooses to keep using CUDA 10.0 "at their own risk". It's one thing to not support a CUDA version, but to knowingly block it is another...

It is undefined behavior to include this specialization. We shouldn't be including it in a released build whatsoever.

Copy link
Contributor Author

@sriramch sriramch Jul 17, 2020

Choose a reason for hiding this comment

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

there is a guard to enable this only for cuda 10.0. if this is removed, then 10.0 ci builds will fail. i'll remove this if 10.0 support is removed from ci.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

it looks like we are no longer building on 10.0. i'll remove this code block and check to be sure.

: std::is_trivially_copyable<cudf::detail::time_point<Duration>> {
};
} // namespace std
sriramch marked this conversation as resolved.
Show resolved Hide resolved
#endif
27 changes: 15 additions & 12 deletions cpp/include/cudf/wrappers/timestamps.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -32,23 +32,26 @@ namespace cudf {
namespace detail {
// TODO: Use chrono::utc_clock when available in libcu++?
template <class Duration>
using time_point = simt::std::chrono::time_point<simt::std::chrono::system_clock, Duration>;
using time_point = simt::std::chrono::sys_time<Duration>;

template <class Duration>
struct timestamp : time_point<Duration> {
// Bring over base class constructors and make them visible here
sriramch marked this conversation as resolved.
Show resolved Hide resolved
using time_point<Duration>::time_point;

// This is needed as __shared__ objects of this type can't be assigned in device code
Copy link
Contributor

Choose a reason for hiding this comment

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

I believe this is only a problem with simt::std::chrono and not an issue with cuda::std::chrono, yes?

Copy link
Contributor Author

@sriramch sriramch Jun 22, 2020

Choose a reason for hiding this comment

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

i think this also shows up with cuda::std::chrono. assignments of timestamps to the __shared__ device variable will trigger this

// when the initializer list constructs subobjects with values, which is what std::time_point
// does.
constexpr timestamp() : time_point<Duration>(Duration()){};
constexpr timestamp(Duration d) : time_point<Duration>(d){};

// Implicitly convert a tick count into a timestamp
// TODO: is this still needed and is this operation even meaningful? The duration units
// cannot be inferred from this
sriramch marked this conversation as resolved.
Show resolved Hide resolved
constexpr timestamp(typename Duration::rep r) : time_point<Duration>(Duration(r)){};
/**
* @brief Constructs a new timestamp by copying the contents of another
* `time_point` and converting its duration value if necessary.
*
* @param other The `timestamp` to copy
*/
// TODO: This explict truncation is intended?
template <class FromDuration>
inline constexpr explicit timestamp(time_point<FromDuration> const& other)
: time_point<Duration>(simt::std::chrono::duration_cast<Duration>(other.time_since_epoch())){};

// The inherited copy constructor will hide the auto generated copy constructor;
// hence, explicitly define and delegate
constexpr timestamp(const time_point<Duration>& other) : time_point<Duration>(other) {}
};
} // namespace detail

Expand Down
9 changes: 7 additions & 2 deletions cpp/src/binaryop/binaryop.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@
#include "compiled/binary_ops.hpp"

#include <bit.hpp.jit>
#include <durations.hpp.jit>
#include <jit/common_headers.hpp>
#include <string>
#include <timestamps.hpp.jit>
Expand Down Expand Up @@ -69,8 +70,12 @@ namespace jit {

const std::string hash = "prog_binop";

const std::vector<std::string> header_names{
"operation.h", "traits.h", cudf_types_hpp, cudf_utilities_bit_hpp, cudf_wrappers_timestamps_hpp};
const std::vector<std::string> header_names{"operation.h",
"traits.h",
cudf_types_hpp,
cudf_utilities_bit_hpp,
cudf_wrappers_timestamps_hpp,
cudf_wrappers_durations_hpp};

std::istream* headers_code(std::string filename, std::iostream& stream)
{
Expand Down
1 change: 1 addition & 0 deletions cpp/src/binaryop/jit/code/kernel.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,7 @@ const char* kernel =
#include <simt/limits>
#include <cudf/utilities/bit.hpp>
#include <cudf/wrappers/timestamps.hpp>
#include <cudf/wrappers/durations.hpp>
#include "operation.h"

template <typename TypeOut, typename TypeLhs, typename TypeRhs, typename TypeOpe>
Expand Down
92 changes: 62 additions & 30 deletions cpp/src/binaryop/jit/code/operation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -31,11 +31,20 @@ const char* operation =
using namespace simt::std;

struct Add {
// Disallow sum of timestamps with any other type (including itself)
// Allow sum between chronos only when both input and output types
// are chronos. Unsupported combinations will fail to compile
template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(!is_timestamp_v<TypeOut> &&
!is_timestamp_v<TypeLhs> &&
!is_timestamp_v<TypeRhs>)>* = nullptr>
enable_if_t<(is_chrono_v<TypeOut> &&
is_chrono_v<TypeLhs> &&
is_chrono_v<TypeRhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return x + y;
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(!is_chrono_v<TypeOut> ||
!is_chrono_v<TypeLhs> ||
!is_chrono_v<TypeRhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(x) + static_cast<TypeCommon>(y));
Expand All @@ -45,11 +54,20 @@ const char* operation =
using RAdd = Add;

struct Sub {
// Disallow difference of timestamps with any other type (including itself)
// Allow difference between chronos only when both input and output types
// are chronos. Unsupported combinations will fail to compile
template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(is_chrono_v<TypeOut> &&
is_chrono_v<TypeLhs> &&
is_chrono_v<TypeRhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return x - y;
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(!is_timestamp_v<TypeOut> &&
!is_timestamp_v<TypeLhs> &&
!is_timestamp_v<TypeRhs>)>* = nullptr>
enable_if_t<(!is_chrono_v<TypeOut> ||
!is_chrono_v<TypeLhs> ||
!is_chrono_v<TypeRhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(x) - static_cast<TypeCommon>(y));
Expand All @@ -64,28 +82,54 @@ const char* operation =
};

struct Mul {
template <typename TypeOut, typename TypeLhs, typename TypeRhs>
template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(!is_duration_v<TypeOut>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(x) * static_cast<TypeCommon>(y));
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(is_duration_v<TypeOut>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return DurationProduct<TypeOut>(x, y);
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(is_duration_v<TypeLhs> && is_integral_v<TypeRhs>) ||
(is_integral_v<TypeLhs> && is_duration_v<TypeRhs>)>* = nullptr>
static TypeOut DurationProduct(TypeLhs x, TypeRhs y) {
return x * y;
}
};

using RMul = Mul;

struct Div {
template <typename TypeOut, typename TypeLhs, typename TypeRhs>
template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(!is_duration_v<TypeLhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(x) / static_cast<TypeCommon>(y));
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(is_duration_v<TypeLhs>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return DurationDivide<TypeOut>(x, y);
}

template <typename TypeOut, typename TypeLhs, typename TypeRhs,
enable_if_t<(is_integral_v<TypeRhs> || is_duration_v<TypeRhs>)>* = nullptr>
static TypeOut DurationDivide(TypeLhs x, TypeRhs y) {
return x / y;
}
};

struct RDiv {
template <typename TypeOut, typename TypeLhs, typename TypeRhs>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(y) / static_cast<TypeCommon>(x));
return Div::operate<TypeOut, TypeRhs, TypeLhs>(y, x);
}
};

Expand Down Expand Up @@ -142,32 +186,20 @@ const char* operation =
static TypeOut operate(TypeLhs x, TypeRhs y) {
return static_cast<TypeOut>(fmod(static_cast<double>(x), static_cast<double>(y)));
}
};

struct RMod {
template <typename TypeOut,
typename TypeLhs,
typename TypeRhs,
enable_if_t<(is_integral_v<typename common_type<TypeOut, TypeLhs, TypeRhs>::type>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
using TypeCommon = typename common_type<TypeOut, TypeLhs, TypeRhs>::type;
return static_cast<TypeOut>(static_cast<TypeCommon>(y) % static_cast<TypeCommon>(x));
}

template <typename TypeOut,
typename TypeLhs,
typename TypeRhs,
enable_if_t<(isFloat<typename common_type<TypeOut, TypeLhs, TypeRhs>::type>)>* = nullptr>
enable_if_t<(is_duration_v<TypeLhs> && is_duration_v<TypeOut>)>* = nullptr>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return static_cast<TypeOut>(fmodf(static_cast<float>(y), static_cast<float>(x)));
return x % y;
}
};

template <typename TypeOut,
typename TypeLhs,
typename TypeRhs,
enable_if_t<(isDouble<typename common_type<TypeOut, TypeLhs, TypeRhs>::type>)>* = nullptr>
struct RMod {
template <typename TypeOut, typename TypeLhs, typename TypeRhs>
static TypeOut operate(TypeLhs x, TypeRhs y) {
return static_cast<TypeOut>(fmod(static_cast<double>(y), static_cast<double>(x)));
return Mod::operate<TypeOut, TypeRhs, TypeLhs>(y, x);
}
};

Expand Down
11 changes: 11 additions & 0 deletions cpp/src/binaryop/jit/code/traits.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,17 @@ const char* traits =
simt::std::is_same<cudf::timestamp_us, T>::value ||
simt::std::is_same<cudf::timestamp_ns, T>::value;

template <typename T>
constexpr bool is_duration_v =
simt::std::is_same<cudf::duration_D, T>::value ||
simt::std::is_same<cudf::duration_s, T>::value ||
simt::std::is_same<cudf::duration_ms, T>::value ||
simt::std::is_same<cudf::duration_us, T>::value ||
simt::std::is_same<cudf::duration_ns, T>::value;

template <typename T>
constexpr bool is_chrono_v = is_timestamp_v<T> || is_duration_v<T>;

template <>
constexpr bool isFloat<float> = true;

Expand Down
2 changes: 1 addition & 1 deletion cpp/src/reductions/simple.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -76,7 +76,7 @@ struct result_type_dispatcher {
// - same dtypes (including cudf::wrappers)
// - any arithmetic dtype to any arithmetic dtype
// - bool to/from any arithmetic dtype
return std::is_convertible<ElementType, ResultType>::value &&
return cudf::is_convertible<ElementType, ResultType>::value &&
(std::is_arithmetic<ResultType>::value ||
std::is_same<Op, cudf::reduction::op::min>::value ||
std::is_same<Op, cudf::reduction::op::max>::value) &&
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/rolling/rolling_detail.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -101,7 +101,7 @@ struct rolling_store_output_functor<_T, true> {
template <typename T = _T, std::enable_if_t<cudf::is_timestamp<T>()> * = nullptr>
CUDA_HOST_DEVICE_CALLABLE void operator()(T &out, T &val, size_type count)
{
out = val.time_since_epoch() / count;
out = static_cast<T>(val.time_since_epoch() / count);
sriramch marked this conversation as resolved.
Show resolved Hide resolved
}
};
} // namespace detail
Expand Down
Loading