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

Fix nvcc-imposed UB in constexpr functions #17534

Merged
merged 7 commits into from
Dec 9, 2024
Merged

Fix nvcc-imposed UB in constexpr functions #17534

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
bd15c0c
time utils
vuule Dec 5, 2024
bc383ab
RLE stream size
vuule Dec 5, 2024
7beb6b4
Merge branch 'branch-25.02' into remove-constexpr-with-cudf_fail
vuule Dec 5, 2024
e745be8
style
vuule Dec 6, 2024
001f3be
Merge branch 'remove-constexpr-with-cudf_fail' of https://github.com/…
vuule Dec 6, 2024
3bbfdba
instances in include
vuule Dec 6, 2024
db95935
Update Copyright date in device_operators.cuh
mythrocks Dec 7, 2024
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
18 changes: 17 additions & 1 deletion cpp/include/cudf/detail/utilities/device_operators.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2019-2023, NVIDIA CORPORATION.
* Copyright (c) 2019-2024, 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 @@ -83,7 +83,11 @@ struct DeviceSum {
template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
static constexpr T identity()
{
#ifndef __CUDA_ARCH__
mythrocks marked this conversation as resolved.
Show resolved Hide resolved
CUDF_FAIL("fixed_point does not yet support device operator identity");
#else
CUDF_UNREACHABLE("fixed_point does not yet support device operator identity");
#endif
return T{};
}
};
Expand Down Expand Up @@ -141,7 +145,11 @@ struct DeviceMin {
template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
static constexpr T identity()
{
#ifndef __CUDA_ARCH__
CUDF_FAIL("fixed_point does not yet support DeviceMin identity");
#else
CUDF_UNREACHABLE("fixed_point does not yet support DeviceMin identity");
#endif
return cuda::std::numeric_limits<T>::max();
}

Expand Down Expand Up @@ -189,7 +197,11 @@ struct DeviceMax {
template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
static constexpr T identity()
{
#ifndef __CUDA_ARCH__
CUDF_FAIL("fixed_point does not yet support DeviceMax identity");
#else
CUDF_UNREACHABLE("fixed_point does not yet support DeviceMax identity");
#endif
return cuda::std::numeric_limits<T>::lowest();
}

Expand Down Expand Up @@ -225,7 +237,11 @@ struct DeviceProduct {
template <typename T, std::enable_if_t<cudf::is_fixed_point<T>()>* = nullptr>
static constexpr T identity()
{
#ifndef __CUDA_ARCH__
CUDF_FAIL("fixed_point does not yet support DeviceProduct identity");
#else
CUDF_UNREACHABLE("fixed_point does not yet support DeviceProduct identity");
#endif
return T{1, numeric::scale_type{0}};
}
};
Expand Down
2 changes: 2 additions & 0 deletions cpp/include/cudf/utilities/span.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -417,7 +417,9 @@ class base_2dspan {
constexpr base_2dspan(RowType<T, dynamic_extent> flat_view, size_t columns)
: _flat{flat_view}, _size{columns == 0 ? 0 : flat_view.size() / columns, columns}
{
#ifndef __CUDA_ARCH__
CUDF_EXPECTS(_size.first * _size.second == flat_view.size(), "Invalid 2D span size");
#endif
}

/**
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/orc/writer_impl.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -506,7 +506,7 @@ size_t max_varint_size()
return cudf::util::div_rounding_up_unsafe(sizeof(T) * 8, 7);
}

constexpr size_t RLE_stream_size(TypeKind kind, size_t count)
size_t RLE_stream_size(TypeKind kind, size_t count)
{
using cudf::util::div_rounding_up_unsafe;
constexpr auto byte_rle_max_len = 128;
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/io/utilities/time_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2021, NVIDIA CORPORATION.
* Copyright (c) 2021-2024, 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 @@ -32,7 +32,7 @@ static const __device__ __constant__ int32_t powers_of_ten[10] = {

struct get_period {
template <typename T>
constexpr int32_t operator()()
int32_t operator()()
{
if constexpr (is_chrono<T>()) { return T::period::den; }
CUDF_FAIL("Invalid, non chrono type");
Expand All @@ -42,7 +42,7 @@ struct get_period {
/**
* @brief Function that translates cuDF time unit to clock frequency
*/
constexpr int32_t to_clockrate(type_id timestamp_type_id)
inline int32_t to_clockrate(type_id timestamp_type_id)
{
return timestamp_type_id == type_id::EMPTY
? 0
Expand Down
Loading