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

Reduce cudf library size #7583

Merged
merged 2 commits into from
Mar 23, 2021
+61 −54
Merged

Reduce cudf library size #7583

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4b24c11
compress cuda fatbinaries to reduce file size
robertmaynard Mar 12, 2021
1943e4e
replace release_assert with cudf_assert
robertmaynard Mar 12, 2021
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
3 changes: 3 additions & 0 deletions cpp/cmake/Modules/ConfigureCUDA.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,9 @@ if(DISABLE_DEPRECATION_WARNING)
list(APPEND CUDF_CUDA_FLAGS -Xcompiler=-Wno-deprecated-declarations)
endif()

# make sure we produce smallest binary size
list(APPEND CUDF_CUDA_FLAGS -Xfatbin=-compress-all)

# Option to enable line info in CUDA device compilation to allow introspection when profiling / memchecking
if(CUDA_ENABLE_LINEINFO)
list(APPEND CUDF_CUDA_FLAGS -lineinfo)
Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cudf/ast/detail/operators.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -187,7 +187,7 @@ CUDA_HOST_DEVICE_CALLABLE constexpr void ast_operator_dispatcher(ast_operator op
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid operator.");
#else
release_assert(false && "Invalid operator.");
cudf_assert(false && "Invalid operator.");
#endif
break;
}
Expand Down Expand Up @@ -784,7 +784,7 @@ struct double_dispatch_binary_operator_types {
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid binary operation.");
#else
release_assert(false && "Invalid binary operation.");
cudf_assert(false && "Invalid binary operation.");
#endif
}
};
Expand Down Expand Up @@ -819,7 +819,7 @@ struct single_dispatch_binary_operator_types {
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid binary operation.");
#else
release_assert(false && "Invalid binary operation.");
cudf_assert(false && "Invalid binary operation.");
#endif
}
};
Expand Down Expand Up @@ -924,7 +924,7 @@ struct dispatch_unary_operator_types {
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid unary operation.");
#else
release_assert(false && "Invalid unary operation.");
cudf_assert(false && "Invalid unary operation.");
#endif
}
};
Expand Down Expand Up @@ -996,7 +996,7 @@ struct return_type_functor {
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid binary operation. Return type cannot be determined.");
#else
release_assert(false && "Invalid binary operation. Return type cannot be determined.");
cudf_assert(false && "Invalid binary operation. Return type cannot be determined.");
#endif
}

Expand Down Expand Up @@ -1024,7 +1024,7 @@ struct return_type_functor {
#ifndef __CUDA_ARCH__
CUDF_FAIL("Invalid unary operation. Return type cannot be determined.");
#else
release_assert(false && "Invalid unary operation. Return type cannot be determined.");
cudf_assert(false && "Invalid unary operation. Return type cannot be determined.");
#endif
}
};
Expand Down
8 changes: 4 additions & 4 deletions cpp/include/cudf/ast/detail/transform.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ struct unary_row_output : public row_output {
Input input,
detail::device_data_reference output) const
{
release_assert(false && "Invalid unary dispatch operator for the provided input.");
cudf_assert(false && "Invalid unary dispatch operator for the provided input.");
}
};

Expand Down Expand Up @@ -116,7 +116,7 @@ struct binary_row_output : public row_output {
RHS rhs,
detail::device_data_reference output) const
{
release_assert(false && "Invalid binary dispatch operator for the provided input.");
cudf_assert(false && "Invalid binary dispatch operator for the provided input.");
}
};

Expand Down Expand Up @@ -239,7 +239,7 @@ struct row_evaluator {
detail::device_data_reference rhs,
detail::device_data_reference output) const
{
release_assert(false && "Invalid binary dispatch operator for the provided input.");
cudf_assert(false && "Invalid binary dispatch operator for the provided input.");
}

private:
Expand Down Expand Up @@ -311,7 +311,7 @@ __device__ void evaluate_row_expression(detail::row_evaluator const& evaluator,
output,
op);
} else {
release_assert(false && "Invalid operator arity.");
cudf_assert(false && "Invalid operator arity.");
}
}
}
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cudf/column/column_device_view.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -774,7 +774,7 @@ struct index_element_fn {
std::is_unsigned<IndexType>::value)>* = nullptr>
__device__ size_type operator()(Args&&... args)
{
release_assert(false and "dictionary indices must be an unsigned integral type");
cudf_assert(false and "dictionary indices must be an unsigned integral type");
return 0;
}
};
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/aggregation/aggregation.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,8 +19,8 @@
#include <cudf/aggregation.hpp>
#include <cudf/column/column_device_view.cuh>
#include <cudf/detail/aggregation/aggregation.hpp>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/detail/utilities/device_atomics.cuh>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/dictionary/dictionary_column_view.hpp>
#include <cudf/table/table_device_view.cuh>

Expand Down Expand Up @@ -103,7 +103,7 @@ struct update_target_element {
column_device_view source,
size_type source_index) const noexcept
{
release_assert(false and "Invalid source type and aggregation combination.");
cudf_assert(false and "Invalid source type and aggregation combination.");
}
};

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/aggregation/aggregation.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#pragma once

#include <cudf/aggregation.hpp>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/types.hpp>
#include <cudf/utilities/error.hpp>
#include <cudf/utilities/traits.hpp>
Expand Down Expand Up @@ -627,7 +627,7 @@ CUDA_HOST_DEVICE_CALLABLE decltype(auto) aggregation_dispatcher(aggregation::Kin
#ifndef __CUDA_ARCH__
CUDF_FAIL("Unsupported aggregation.");
#else
release_assert(false && "Unsupported aggregation.");
cudf_assert(false && "Unsupported aggregation.");

// The following code will never be reached, but the compiler generates a
// warning if there isn't a return value.
Expand Down
2 changes: 1 addition & 1 deletion cpp/include/cudf/detail/gather.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,8 +18,8 @@
#include <cudf/detail/copy.hpp>
#include <cudf/detail/indexalator.cuh>
#include <cudf/detail/null_mask.hpp>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/detail/utilities/cuda.cuh>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/valid_if.cuh>
#include <cudf/dictionary/dictionary_column_view.hpp>
#include <cudf/dictionary/dictionary_factories.hpp>
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/indexalator.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -268,7 +268,7 @@ struct input_indexalator : base_indexalator<input_indexalator> {
template <typename T, std::enable_if_t<not is_index_type<T>()>* = nullptr>
__device__ size_type operator()(void const* tp)
{
release_assert(false and "only index types are supported");
cudf_assert(false and "only index types are supported");
return 0;
}
};
Expand Down Expand Up @@ -366,7 +366,7 @@ struct output_indexalator : base_indexalator<output_indexalator> {
template <typename T, std::enable_if_t<not is_index_type<T>()>* = nullptr>
__device__ void operator()(void* tp, size_type const value)
{
release_assert(false and "only index types are supported");
cudf_assert(false and "only index types are supported");
}
};

Expand Down
6 changes: 3 additions & 3 deletions .../cudf/detail/utilities/release_assert.cuh → cpp/include/cudf/detail/utilities/assert.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,11 +27,11 @@
*
* Relies on the `__PRETTY_FUNCTION__` macro which is specific to GCC and Clang.
*/
#if defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))
#if !defined(NDEBUG) && defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))
#define __ASSERT_STR_HELPER(x) #x
#define release_assert(e) \
#define cudf_assert(e) \
((e) ? static_cast<void>(0) \
: __assert_fail(__ASSERT_STR_HELPER(e), __FILE__, __LINE__, __PRETTY_FUNCTION__))
#else
#define release_assert(e) (static_cast<void>(0))
#define cudf_assert(e) (static_cast<void>(0))
#endif
14 changes: 7 additions & 7 deletions cpp/include/cudf/detail/utilities/hash_functions.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
#pragma once

#include <cudf/column/column_device_view.cuh>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/strings/string_view.cuh>
#include <hash/hash_constants.hpp>

Expand Down Expand Up @@ -154,7 +154,7 @@ struct MD5ListHasher {
size_type offset_end,
md5_intermediate_data* hash_state) const
{
release_assert(false && "MD5 Unsupported chrono type column");
cudf_assert(false && "MD5 Unsupported chrono type column");
}

template <typename T, std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
Expand All @@ -163,7 +163,7 @@ struct MD5ListHasher {
size_type offset_end,
md5_intermediate_data* hash_state) const
{
release_assert(false && "MD5 Unsupported non-fixed-width type column");
cudf_assert(false && "MD5 Unsupported non-fixed-width type column");
}

template <typename T, std::enable_if_t<is_floating_point<T>()>* = nullptr>
Expand Down Expand Up @@ -273,15 +273,15 @@ struct MD5Hash {
size_type row_index,
md5_intermediate_data* hash_state) const
{
release_assert(false && "MD5 Unsupported chrono type column");
cudf_assert(false && "MD5 Unsupported chrono type column");
}

template <typename T, std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
void __device__ operator()(column_device_view col,
size_type row_index,
md5_intermediate_data* hash_state) const
{
release_assert(false && "MD5 Unsupported non-fixed-width type column");
cudf_assert(false && "MD5 Unsupported non-fixed-width type column");
}

template <typename T, std::enable_if_t<is_floating_point<T>()>* = nullptr>
Expand Down Expand Up @@ -344,7 +344,7 @@ void CUDA_DEVICE_CALLABLE MD5Hash::operator()<list_view>(column_device_view col,
column_device_view offsets = col.child(offsets_column_index);
column_device_view data = col.child(data_column_index);

if (data.type().id() == type_id::LIST) release_assert(false && "Nested list unsupported");
if (data.type().id() == type_id::LIST) cudf_assert(false && "Nested list unsupported");

cudf::type_dispatcher(data.type(),
MD5ListHasher{},
Expand Down Expand Up @@ -724,7 +724,7 @@ struct IdentityHash {
CUDA_HOST_DEVICE_CALLABLE std::enable_if_t<!std::is_arithmetic<Key>::value, return_type>
operator()(Key const& key) const
{
release_assert(false && "IdentityHash does not support this data type");
cudf_assert(false && "IdentityHash does not support this data type");
return 0;
}

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/fixed_point/fixed_point.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@

#pragma once

#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/types.hpp>

// Note: The <cuda/std/*> versions are used in order for Jitify to work with our fixed_point type.
Expand Down Expand Up @@ -91,7 +91,7 @@ template <typename Rep,
is_supported_representation_type<Rep>())>* = nullptr>
CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
{
release_assert(exponent >= 0 && "integer exponentiation with negative exponent is not possible.");
cudf_assert(exponent >= 0 && "integer exponentiation with negative exponent is not possible.");
if (exponent == 0) return static_cast<Rep>(1);
auto extra = static_cast<Rep>(1);
auto square = static_cast<Rep>(Base);
Expand Down
14 changes: 7 additions & 7 deletions cpp/include/cudf/lists/list_device_view.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,12 +37,12 @@ class list_device_view {
: lists_column(lists_column), _row_index(row_index)
{
column_device_view const& offsets = lists_column.offsets();
release_assert(row_index >= 0 && row_index < lists_column.size() &&
row_index < offsets.size() && "row_index out of bounds");
cudf_assert(row_index >= 0 && row_index < lists_column.size() && row_index < offsets.size() &&
"row_index out of bounds");

begin_offset = offsets.element<size_type>(row_index);
release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() &&
"begin_offset out of bounds.");
cudf_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() &&
"begin_offset out of bounds.");
_size = offsets.element<size_type>(row_index + 1) - begin_offset;
}

Expand Down Expand Up @@ -71,7 +71,7 @@ class list_device_view {
*/
CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const
{
release_assert(idx >= 0 && idx < size() && "idx out of bounds");
cudf_assert(idx >= 0 && idx < size() && "idx out of bounds");
return begin_offset + idx;
}

Expand All @@ -93,7 +93,7 @@ class list_device_view {
*/
CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const
{
release_assert(idx >= 0 && idx < size() && "Index out of bounds.");
cudf_assert(idx >= 0 && idx < size() && "Index out of bounds.");
auto element_offset = begin_offset + idx;
return lists_column.child().is_null(element_offset);
}
Expand Down Expand Up @@ -294,7 +294,7 @@ struct list_size_functor {
CUDA_HOST_DEVICE_CALLABLE list_size_functor(column_device_view const& d_col) : d_column(d_col)
{
#if defined(__CUDA_ARCH__)
release_assert(d_col.type().id() == type_id::LIST && "Only list type column is supported");
cudf_assert(d_col.type().id() == type_id::LIST && "Only list type column is supported");
#else
CUDF_EXPECTS(d_col.type().id() == type_id::LIST, "Only list type column is supported");
#endif
Expand Down
6 changes: 3 additions & 3 deletions cpp/include/cudf/table/row_operators.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,8 +17,8 @@
#pragma once

#include <cudf/column/column_device_view.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/detail/utilities/hash_functions.cuh>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/sorting.hpp>
#include <cudf/table/table_device_view.cuh>
#include <cudf/utilities/traits.hpp>
Expand Down Expand Up @@ -190,7 +190,7 @@ class element_equality_comparator {
std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* = nullptr>
__device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index)
{
release_assert(false && "Attempted to compare elements of uncomparable types.");
cudf_assert(false && "Attempted to compare elements of uncomparable types.");
return false;
}

Expand Down Expand Up @@ -291,7 +291,7 @@ class element_relational_comparator {
std::enable_if_t<not cudf::is_relationally_comparable<Element, Element>()>* = nullptr>
__device__ weak_ordering operator()(size_type lhs_element_index, size_type rhs_element_index)
{
release_assert(false && "Attempted to compare elements of uncomparable types.");
cudf_assert(false && "Attempted to compare elements of uncomparable types.");
return weak_ordering::LESS;
}

Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/utilities/type_dispatcher.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@

#pragma once

#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/fixed_point/fixed_point.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/error.hpp>
Expand Down Expand Up @@ -501,7 +501,7 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty
#ifndef __CUDA_ARCH__
CUDF_FAIL("Unsupported type_id.");
#else
release_assert(false && "Unsupported type_id.");
cudf_assert(false && "Unsupported type_id.");

// The following code will never be reached, but the compiler generates a
// warning if there isn't a return value.
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/groupby/hash/multi_pass_kernels.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,8 +20,8 @@
#include <cudf/aggregation.hpp>
#include <cudf/column/column_device_view.cuh>
#include <cudf/detail/aggregation/aggregation.hpp>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/detail/utilities/device_atomics.cuh>
#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/table/table_device_view.cuh>
#include <cudf/utilities/type_dispatcher.hpp>

Expand Down Expand Up @@ -65,7 +65,7 @@ struct var_hash_functor {
size_type source_index,
size_type target_index) noexcept
{
release_assert(false and "Invalid source type for std, var aggregation combination.");
cudf_assert(false and "Invalid source type for std, var aggregation combination.");
}

template <typename Source>
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/page_data.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
#include <io/utilities/block_utils.cuh>
#include <io/utilities/column_buffer.hpp>

#include <cudf/detail/utilities/release_assert.cuh>
#include <cudf/detail/utilities/assert.cuh>
#include <cudf/utilities/bit.hpp>

#include <rmm/cuda_stream_view.hpp>
Expand Down
Loading