diff --git a/CHANGELOG.md b/CHANGELOG.md index 44a5bbbc5cb..0a5f619a2e9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -114,6 +114,7 @@ - PR #6833 Use settings.xml if existing for internal build - PR #6835 Move template param to member var to improve compile of hash/groupby.cu - PR #6837 Avoid gather when copying strings view from start of strings column +- PR #6859 Move align_ptr_for_type() from cuda.cuh to alignment.hpp ## Bug Fixes diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index c7aeb04c31b..12fedc066a6 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -85,6 +85,7 @@ test: - test -f $PREFIX/include/cudf/detail/transform.hpp - test -f $PREFIX/include/cudf/detail/transpose.hpp - test -f $PREFIX/include/cudf/detail/unary.hpp + - test -f $PREFIX/include/cudf/detail/utilities/alignment.hpp - test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp - test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h - test -f $PREFIX/include/cudf/dictionary/detail/concatenate.hpp diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 69517aa8d27..237242f5346 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -17,7 +17,7 @@ #include #include -#include +#include #include #include #include diff --git a/cpp/include/cudf/detail/utilities/alignment.hpp b/cpp/include/cudf/detail/utilities/alignment.hpp new file mode 100644 index 00000000000..e52032fe104 --- /dev/null +++ b/cpp/include/cudf/detail/utilities/alignment.hpp @@ -0,0 +1,46 @@ +/* + * Copyright (c) 2020, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace cudf { +namespace detail { + +/** + * @brief Returns the aligned address for holding array of type T in pre-allocated memory. + * + * @tparam T The data type to align upon. + * + * @param destination pointer to pre-allocated contiguous storage to store type T. + * @return Pointer of type T, aligned to alignment of type T. + */ +template +T* align_ptr_for_type(void* destination) +{ + constexpr std::size_t bytes_needed{sizeof(T)}; + constexpr std::size_t alignment{alignof(T)}; + + // pad the allocation for aligning the first pointer + auto padded_bytes_needed = bytes_needed + (alignment - 1); + // std::align captures last argument by reference and modifies it, but we don't want it modified + return reinterpret_cast( + std::align(alignment, bytes_needed, destination, padded_bytes_needed)); +} + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/detail/utilities/cuda.cuh b/cpp/include/cudf/detail/utilities/cuda.cuh index 141caa1fc8c..33c61414a1c 100644 --- a/cpp/include/cudf/detail/utilities/cuda.cuh +++ b/cpp/include/cudf/detail/utilities/cuda.cuh @@ -174,23 +174,5 @@ void device_single_thread(Functor functor, rmm::cuda_stream_view stream = rmm::c single_thread_kernel<<<1, 1, 0, stream.value()>>>(functor); } -/** - * @brief Returns the aligned address for holding array of type T in pre-allocated memory - * @param destination pointer to pre-allocated contiguous storage to store type T. - * @return Pointer of type T, aligned to alignment of type T. - */ -template -T* align_ptr_for_type(void* destination) -{ - constexpr std::size_t bytes_needed{sizeof(T)}; - constexpr std::size_t alignment{alignof(T)}; - - // pad the allocation for aligning the first pointer - auto padded_bytes_needed = bytes_needed + (alignment - 1); - // std::align captures last argument by reference and modifies it, but we don't want it modified - return reinterpret_cast( - std::align(alignment, bytes_needed, destination, padded_bytes_needed)); -} - } // namespace detail } // namespace cudf