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

Force inlining to improve AST performance #9530

Merged
merged 6 commits into from
Nov 12, 2021
Merged

Force inlining to improve AST performance #9530

Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
dd7d860
Force inlining of the type dispatcher and a couple of class methods o…
vyasr Oct 26, 2021
9b1805f
clang-format.
vyasr Oct 26, 2021
f663915
Define a new macro for forceinlining.
vyasr Oct 30, 2021
fbc687f
Move macro to types.hpp for global visibility and add a device-only v…
vyasr Nov 3, 2021
7bef3b7
Rename macros to shorter versions.
vyasr Nov 11, 2021
1733fd4
Fix typo.
vyasr Nov 11, 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
13 changes: 10 additions & 3 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 @@ -26,6 +26,14 @@

#include <string>

// Define a new macro analogous to CUDA_HOST_DEVICE_CALLABLE that maximizes the
// chances of inlining when applied to a callable that is called on the GPU.
#ifdef __CUDACC__
#define CUDA_HOST_DEVICE_FORCEINLINE_CALLABLE __host__ __device__ __forceinline__
vyasr marked this conversation as resolved.
Show resolved Hide resolved
#else
#define CUDA_HOST_DEVICE_FORCEINLINE_CALLABLE inline
#endif

/**
* @file
* @brief Defines the mapping between `cudf::type_id` runtime type information
Expand Down Expand Up @@ -411,9 +419,8 @@ using scalar_device_type_t = typename type_to_scalar_type_impl<T>::ScalarDeviceT
template <template <cudf::type_id> typename IdTypeMap = id_to_type_impl,
typename Functor,
typename... Ts>
__host__ __device__ __forceinline__ constexpr decltype(auto) type_dispatcher(cudf::data_type dtype,
Functor f,
Ts&&... args)
CUDA_HOST_DEVICE_FORCEINLINE_CALLABLE constexpr decltype(auto) type_dispatcher(
cudf::data_type dtype, Functor f, Ts&&... args)
vyasr marked this conversation as resolved.
Show resolved Hide resolved
{
switch (dtype.id()) {
case type_id::BOOL8:
Expand Down