diff --git a/CHANGELOG.md b/CHANGELOG.md index ab6a1b62..7e8dc6c8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -12,6 +12,8 @@ Full documentation for hipCUB is available at [https://rocm.docs.amd.com/project * Added `--emulation` option for `rtest.py` * Unit tests can be run with `[--emulation|-e|--test|-t]=` * Added `DeviceSelect::FlaggedIf` and its inplace overload. +* Added CUB macros missing from hipCUB: `HIPCUB_MAX`, `HIPCUB_MIN`, `HIPCUB_QUOTIENT_FLOOR`, `HIPCUB_QUOTIENT_CEILING`, `HIPCUB_ROUND_UP_NEAREST` and `HIPCUB_ROUND_DOWN_NEAREST`. +* Added `hipcub::AliasTemporaries` function for CUB parity. ### Changed * Changed the subset of tests that are run for smoke tests such that the smoke test will complete with faster run-time and to never exceed 2GB of vram usage. Use `python rtest.py [--emulation|-e|--test|-t]=smoke` to run these tests. diff --git a/hipcub/include/hipcub/backend/cub/hipcub.hpp b/hipcub/include/hipcub/backend/cub/hipcub.hpp index a6e627cf..1c5a8a3e 100644 --- a/hipcub/include/hipcub/backend/cub/hipcub.hpp +++ b/hipcub/include/hipcub/backend/cub/hipcub.hpp @@ -112,5 +112,7 @@ // These functions must be wrapped so they return // hipError_t instead of cudaError_t #include "util_allocator.hpp" +#include "util_device.hpp" +#include "util_temporary_storage.hpp" #endif // HIPCUB_CUB_HIPCUB_HPP_ diff --git a/hipcub/include/hipcub/backend/cub/util_device.hpp b/hipcub/include/hipcub/backend/cub/util_device.hpp new file mode 100644 index 00000000..06683269 --- /dev/null +++ b/hipcub/include/hipcub/backend/cub/util_device.hpp @@ -0,0 +1,39 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_CUB_UTIL_DEVICE_HPP_ +#define HIPCUB_CUB_UTIL_DEVICE_HPP_ + +#include "../../config.hpp" + +#include "util_temporary_storage.hpp" + +#include "cub/util_device.cuh" + +#endif // HIPCUB_CUB_UTIL_DEVICE_HPP_ diff --git a/hipcub/include/hipcub/backend/cub/util_macro.hpp b/hipcub/include/hipcub/backend/cub/util_macro.hpp new file mode 100644 index 00000000..ce62e2a0 --- /dev/null +++ b/hipcub/include/hipcub/backend/cub/util_macro.hpp @@ -0,0 +1,71 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_CUB_MACRO_HPP_ +#define HIPCUB_CUB_MACRO_HPP_ + +#include "../../config.hpp" + +#include "cub/util_macro.cuh" + +BEGIN_HIPCUB_NAMESPACE + +#ifndef HIPCUB_MAX + /// Select maximum(a, b) + #define HIPCUB_MAX(a, b) CUB_MAX(a, b) +#endif + +#ifndef HIPCUB_MIN + /// Select minimum(a, b) + #define HIPCUB_MIN(a, b) CUB_MIN(a, b) +#endif + +#ifndef HIPCUB_QUOTIENT_FLOOR + /// Quotient of x/y rounded down to nearest integer + #define HIPCUB_QUOTIENT_FLOOR(x, y) CUB_QUOTIENT_FLOOR(x, y) +#endif + +#ifndef HIPCUB_QUOTIENT_CEILING + /// Quotient of x/y rounded up to nearest integer + #define HIPCUB_QUOTIENT_CEILING(x, y) CUB_QUOTIENT_CEILING(x, y) +#endif + +#ifndef HIPCUB_ROUND_UP_NEAREST + /// x rounded up to the nearest multiple of y + #define HIPCUB_ROUND_UP_NEAREST(x, y) CUB_ROUND_UP_NEAREST(x, y) +#endif + +#ifndef HIPCUB_ROUND_DOWN_NEAREST + /// x rounded down to the nearest multiple of y + #define HIPCUB_ROUND_DOWN_NEAREST(x, y) CUB_ROUND_DOWN_NEAREST(x, y) +#endif + +END_HIPCUB_NAMESPACE + +#endif // HIPCUB_CUB_MACRO_HPP_ diff --git a/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp b/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp new file mode 100644 index 00000000..e73865a4 --- /dev/null +++ b/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp @@ -0,0 +1,71 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_CUB_UTIL_TEMPORARY_STORAGE_HPP_ +#define HIPCUB_CUB_UTIL_TEMPORARY_STORAGE_HPP_ + +#include "../../config.hpp" + +#include + +BEGIN_HIPCUB_NAMESPACE + +/// \brief Alias temporaries to externally-allocated device storage (or simply return the amount of storage needed). +/// \tparam ALLOCATIONS - The number of allocations that are needed. +/// \param d_temp_storage - [in] Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. +/// \param temp_storage_bytes - [in,out] Size in bytes of \t d_temp_storage allocation. +/// \param allocations - [out] Pointers to device allocations needed. +/// \param allocation_sizes - [in] Sizes in bytes of device allocations needed. +template +HIPCUB_HOST_DEVICE +HIPCUB_FORCEINLINE hipError_t AliasTemporaries(void* d_temp_storage, + size_t& temp_storage_bytes, + void* (&allocations)[ALLOCATIONS], + size_t (&allocation_sizes)[ALLOCATIONS]) +{ + cudaError_t error = ::cub::AliasTemporaries(d_temp_storage, + temp_storage_bytes, + allocations, + allocation_sizes); + + if(cudaSuccess == error) + { + return hipSuccess; + } + else if(cudaErrorInvalidValue == error) + { + return hipErrorInvalidValue; + } + + return hipErrorUnknown; +} + +END_HIPCUB_NAMESPACE + +#endif // HIPCUB_CUB_UTIL_TEMPORARY_STORAGE_HPP_ diff --git a/hipcub/include/hipcub/backend/rocprim/hipcub.hpp b/hipcub/include/hipcub/backend/rocprim/hipcub.hpp index d06f968b..5c3f5a56 100644 --- a/hipcub/include/hipcub/backend/rocprim/hipcub.hpp +++ b/hipcub/include/hipcub/backend/rocprim/hipcub.hpp @@ -102,7 +102,10 @@ // Util #include "tuple.hpp" #include "util_allocator.hpp" +#include "util_device.hpp" +#include "util_macro.hpp" #include "util_ptx.hpp" +#include "util_temporary_storage.hpp" #include "util_type.hpp" #endif // HIPCUB_ROCPRIM_HIPCUB_HPP_ diff --git a/hipcub/include/hipcub/backend/rocprim/util_device.hpp b/hipcub/include/hipcub/backend/rocprim/util_device.hpp new file mode 100644 index 00000000..25f1aca7 --- /dev/null +++ b/hipcub/include/hipcub/backend/rocprim/util_device.hpp @@ -0,0 +1,37 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_ROCPRIM_UTIL_DEVICE_HPP_ +#define HIPCUB_ROCPRIM_UTIL_DEVICE_HPP_ + +#include "../../config.hpp" + +#include "util_temporary_storage.hpp" + +#endif // HIPCUB_ROCPRIM_UTIL_DEVICE_HPP_ diff --git a/hipcub/include/hipcub/backend/rocprim/util_macro.hpp b/hipcub/include/hipcub/backend/rocprim/util_macro.hpp new file mode 100644 index 00000000..8c9f33c4 --- /dev/null +++ b/hipcub/include/hipcub/backend/rocprim/util_macro.hpp @@ -0,0 +1,98 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_ROCPRIM_MACRO_HPP_ +#define HIPCUB_ROCPRIM_MACRO_HPP_ + +#include "../../config.hpp" + +#include + +BEGIN_HIPCUB_NAMESPACE + +/** + * \addtogroup UtilModule + * @{ + */ + +#ifndef DOXYGEN_SHOULD_SKIP_THIS + #define HIPCUB_PREVENT_MACRO_SUBSTITUTION +template +constexpr __host__ __device__ +auto min HIPCUB_PREVENT_MACRO_SUBSTITUTION(T&& t, U&& u) + -> decltype(t < u ? std::forward(t) : std::forward(u)) +{ + return t < u ? std::forward(t) : std::forward(u); +} + +template +constexpr __host__ __device__ +auto max HIPCUB_PREVENT_MACRO_SUBSTITUTION(T&& t, U&& u) + -> decltype(t < u ? std::forward(u) : std::forward(t)) +{ + return t < u ? std::forward(u) : std::forward(t); +} + #undef HIPCUB_PREVENT_MACRO_SUBSTITUTION +#endif + +#ifndef HIPCUB_MAX + /// Select maximum(a, b) + #define HIPCUB_MAX(a, b) (((b) > (a)) ? (b) : (a)) +#endif + +#ifndef HIPCUB_MIN + /// Select minimum(a, b) + #define HIPCUB_MIN(a, b) (((b) < (a)) ? (b) : (a)) +#endif + +#ifndef HIPCUB_QUOTIENT_FLOOR + /// Quotient of x/y rounded down to nearest integer + #define HIPCUB_QUOTIENT_FLOOR(x, y) ((x) / (y)) +#endif + +#ifndef HIPCUB_QUOTIENT_CEILING + /// Quotient of x/y rounded up to nearest integer + #define HIPCUB_QUOTIENT_CEILING(x, y) (((x) + (y)-1) / (y)) +#endif + +#ifndef HIPCUB_ROUND_UP_NEAREST + /// x rounded up to the nearest multiple of y + #define HIPCUB_ROUND_UP_NEAREST(x, y) (HIPCUB_QUOTIENT_CEILING(x, y) * y) +#endif + +#ifndef HIPCUB_ROUND_DOWN_NEAREST + /// x rounded down to the nearest multiple of y + #define HIPCUB_ROUND_DOWN_NEAREST(x, y) (((x) / (y)) * y) +#endif + +/** @} */ // end group UtilModule + +END_HIPCUB_NAMESPACE + +#endif // HIPCUB_ROCPRIM_MACRO_HPP_ diff --git a/hipcub/include/hipcub/backend/rocprim/util_temporary_storage.hpp b/hipcub/include/hipcub/backend/rocprim/util_temporary_storage.hpp new file mode 100644 index 00000000..ed6d63f8 --- /dev/null +++ b/hipcub/include/hipcub/backend/rocprim/util_temporary_storage.hpp @@ -0,0 +1,88 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_ROCPRIM_UTIL_TEMPORARY_STORAGE_HPP_ +#define HIPCUB_ROCPRIM_UTIL_TEMPORARY_STORAGE_HPP_ + +#include "../../config.hpp" + +#include + +#include + +BEGIN_HIPCUB_NAMESPACE + +namespace detail +{ +// Base case: When N == 0 +template +HIPCUB_HOST_DEVICE +typename std::enable_if::type generate_partition(void* d_temp_storage, + size_t& temp_storage_bytes, + Generator /*gen*/, + Ts... args) +{ + return rocprim::detail::temp_storage::partition( + d_temp_storage, + temp_storage_bytes, + rocprim::detail::temp_storage::linear_partition(args...)); +} + +// Recursive case: When N > 0 +template +HIPCUB_HOST_DEVICE +typename std::enable_if<(N > 0), hipError_t>::type + generate_partition(void* d_temp_storage, size_t& temp_storage_bytes, Generator gen, Ts... args) +{ + return generate_partition(d_temp_storage, temp_storage_bytes, gen, gen(N - 1), args...); +} +} // namespace detail + +/// \brief Alias temporaries to externally-allocated device storage (or simply return the amount of storage needed). +/// \tparam ALLOCATIONS - The number of allocations that are needed. +/// \param d_temp_storage - [in] Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. +/// \param temp_storage_bytes - [in,out] Size in bytes of \t d_temp_storage allocation. +/// \param allocations - [out] Pointers to device allocations needed. +/// \param allocation_sizes - [in] Sizes in bytes of device allocations needed. +template +HIPCUB_HOST_DEVICE +HIPCUB_FORCEINLINE hipError_t AliasTemporaries(void* d_temp_storage, + size_t& temp_storage_bytes, + void* (&allocations)[ALLOCATIONS], + size_t (&allocation_sizes)[ALLOCATIONS]) +{ + auto generator = [&](int i) + { return rocprim::detail::temp_storage::make_partition(&allocations[i], allocation_sizes[i]); }; + + return detail::generate_partition(d_temp_storage, temp_storage_bytes, generator); +} + +END_HIPCUB_NAMESPACE + +#endif // HIPCUB_ROCPRIM_UTIL_TEMPORARY_STORAGE_HPP_ diff --git a/hipcub/include/hipcub/util_device.hpp b/hipcub/include/hipcub/util_device.hpp new file mode 100644 index 00000000..2e75212e --- /dev/null +++ b/hipcub/include/hipcub/util_device.hpp @@ -0,0 +1,39 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_UTIL_DEVICE_HPP_ +#define HIPCUB_UTIL_DEVICE_HPP_ + +#ifdef __HIP_PLATFORM_AMD__ + #include "backend/rocprim/util_device.hpp" +#elif defined(__HIP_PLATFORM_NVIDIA__) + #include "backend/cub/util_device.hpp" +#endif + +#endif // HIPCUB_UTIL_DEVICE_HPP_ diff --git a/hipcub/include/hipcub/util_macro.hpp b/hipcub/include/hipcub/util_macro.hpp new file mode 100644 index 00000000..fdcf92dc --- /dev/null +++ b/hipcub/include/hipcub/util_macro.hpp @@ -0,0 +1,37 @@ +/****************************************************************************** + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_UTIL_MACRO_HPP_ +#define HIPCUB_UTIL_MACRO_HPP_ + +#ifdef __HIP_PLATFORM_AMD__ + #include "backend/rocprim/util_macro.hpp" +#elif defined(__HIP_PLATFORM_NVIDIA__) + #include "backend/cub/util_macro.hpp" +#endif + +#endif // HIPCUB_UTIL_MACRO_HPP_ diff --git a/hipcub/include/hipcub/util_temporary_storage.hpp b/hipcub/include/hipcub/util_temporary_storage.hpp new file mode 100644 index 00000000..e36087a4 --- /dev/null +++ b/hipcub/include/hipcub/util_temporary_storage.hpp @@ -0,0 +1,39 @@ +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2024, NVIDIA CORPORATION. All rights reserved. + * Modifications Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of the NVIDIA CORPORATION nor the + * names of its contributors may be used to endorse or promote products + * derived from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE + * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY + * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES + * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND + * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + * + ******************************************************************************/ + +#ifndef HIPCUB_UTIL_TEMPORARY_STORAGE_HPP_ +#define HIPCUB_UTIL_TEMPORARY_STORAGE_HPP_ + +#ifdef __HIP_PLATFORM_AMD__ + #include "backend/rocprim/util_temporary_storage.hpp" +#elif defined(__HIP_PLATFORM_NVIDIA__) + #include "backend/cub/util_temporary_storage.hpp" +#endif + +#endif // HIPCUB_UTIL_TEMPORARY_STORAGE_HPP_ diff --git a/test/hipcub/CMakeLists.txt b/test/hipcub/CMakeLists.txt index f3243ccd..e5db0448 100644 --- a/test/hipcub/CMakeLists.txt +++ b/test/hipcub/CMakeLists.txt @@ -222,6 +222,7 @@ add_hipcub_test("hipcub.DeviceSpmv" test_hipcub_device_spmv.cpp) add_hipcub_test("hipcub.DevicePartition" test_hipcub_device_partition.cpp) add_hipcub_test("hipcub.Grid" test_hipcub_grid.cpp) add_hipcub_test("hipcub.UtilPtx" test_hipcub_util_ptx.cpp) +add_hipcub_test("hipcub.UtilDevice" test_hipcub_util_device.cpp) add_hipcub_test("hipcub.Vector" test_hipcub_vector.cpp) add_hipcub_test("hipcub.WarpExchange" test_hipcub_warp_exchange.cpp) add_hipcub_test("hipcub.WarpLoad" test_hipcub_warp_load.cpp) diff --git a/test/hipcub/test_hipcub_util_device.cpp b/test/hipcub/test_hipcub_util_device.cpp new file mode 100644 index 00000000..eeccd247 --- /dev/null +++ b/test/hipcub/test_hipcub_util_device.cpp @@ -0,0 +1,129 @@ +// MIT License +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "common_test_header.hpp" + +// hipcub API +#include "hipcub/util_device.hpp" + +template +__global__ +void alias_temporaries_kernel(T* data, size_t* temp_storage_bytes) +{ + T* allocations[10]; + size_t allocation_sizes[10] = {1, 2, 3, 5, 8, 13, 21, 34, 55, 89}; + (void)hipcub::AliasTemporaries(data, *temp_storage_bytes, allocations, allocation_sizes); +} + +TEST(HipcubUtilDevice, AliasTemporariesDevice) +{ + int device_id = test_common_utils::obtain_device_from_ctest(); + SCOPED_TRACE(testing::Message() << "with device_id= " << device_id); + HIP_CHECK(hipSetDevice(device_id)); + + void* data = nullptr; + size_t temp_storage_bytes_host = 0; // Temporary storage on the host + size_t* device_temp_storage_bytes; + HIP_CHECK(test_common_utils::hipMallocHelper(&device_temp_storage_bytes, sizeof(size_t))); + + // First kernel call to determine required temp storage size + alias_temporaries_kernel<<<1, 1, 0, 0>>>(data, device_temp_storage_bytes); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipGetLastError()); + + // Copy the device storage size to host + HIP_CHECK(hipMemcpy(&temp_storage_bytes_host, + device_temp_storage_bytes, + sizeof(size_t), + hipMemcpyDeviceToHost)); + ASSERT_GT(temp_storage_bytes_host, 0U); + + // Allocate the actual data buffer on the device + HIP_CHECK(test_common_utils::hipMallocHelper(&data, temp_storage_bytes_host)); + + // Second kernel call with allocated buffer + alias_temporaries_kernel<<<1, 1, 0, 0>>>(data, device_temp_storage_bytes); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipGetLastError()); + + // Free device memory + HIP_CHECK(hipFree(device_temp_storage_bytes)); + HIP_CHECK(hipFree(data)); +} + +TEST(HipcubUtilDevice, AliasTemporariesHost) +{ + int device_id = test_common_utils::obtain_device_from_ctest(); + SCOPED_TRACE(testing::Message() << "with device_id= " << device_id); + HIP_CHECK(hipSetDevice(device_id)); + + void* data = nullptr; + size_t temp_storage_bytes = 0; + void* allocations[10]; + size_t allocation_sizes[10] = {1, 789, 3, 5, 8, 13, 21, 257, 256, 890}; + + size_t min_size = 0; + for(unsigned int i = 0; i < 10; i++) + { + min_size += allocation_sizes[i]; + } + + // Determine storage size + HIP_CHECK(hipcub::AliasTemporaries(data, temp_storage_bytes, allocations, allocation_sizes)); + + // Should be larger or equal to the sum of all sizes. + ASSERT_GT(temp_storage_bytes, min_size - 1); + + // Allocate the actual data buffer on the device + HIP_CHECK(test_common_utils::hipMallocHelper(&data, temp_storage_bytes)); + + size_t zero_size = 0; + // Check for error if it does not fit + hipError_t error = hipcub::AliasTemporaries(data, zero_size, allocations, allocation_sizes); + test_utils::assert_eq(error, hipErrorInvalidValue); + + HIP_CHECK(hipcub::AliasTemporaries(data, temp_storage_bytes, allocations, allocation_sizes)); + + test_utils::assert_eq(data, allocations[0]); + + for(unsigned int i = 1; i < 10; i++) + { + // The allocations should be in increasing order. + ASSERT_GT(allocations[i], allocations[i - 1]); + size_t current_pointer = (size_t)allocations[i]; + size_t before_pointer = (size_t)allocations[i - 1]; + size_t distance = current_pointer - before_pointer; + + // Check if all pointer have enough space + ASSERT_GT(distance + 1, allocation_sizes[i - 1]); + } + + size_t last_pointer = (size_t)allocations[9]; + size_t start_pointer = (size_t)data; + size_t max_size = start_pointer + temp_storage_bytes; + size_t last_size = max_size - last_pointer; + + // Last size should be equal or larger then the last value in allocation_sizes + ASSERT_GT(last_size + 1, allocation_sizes[9]); + + HIP_CHECK(hipFree(data)); +}