hipCUB

https://rocm.docs.amd.com/projects/hipCUB/en/latest/
Admin
hipCUB is a thin, header-only wrapper library on top of rocPRIM or CUB that provides a unified HIP...
Tokens:9,028
Snippets:84
Trust Score:10
Update:1 month ago
Show doc for...
Tokens:
Raw
### Build and install hipCUB with rmake.py

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Runs the build and installation process, which also downloads and installs rocPRIM to C:\hipSDK.

```bash
python rmake.py -i
```

--------------------------------

### Install hipCUB

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Install the built hipCUB files to the system.

```bash
make install
```

--------------------------------

### Install ROCm CCL on RHEL-based distributions

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/install.html

Command to install the latest ROCm CCL development package on RHEL-based systems.

```bash
sudo dnf install amdrocm-ccl-devel
```

--------------------------------

### Install ROCm CCL on SLES

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/install.html

Command to install the latest ROCm CCL development package on SLES.

```bash
sudo zypper install amdrocm-ccl-devel
```

--------------------------------

### Install ROCm CCL on Debian-based distributions

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/install.html

Command to install the latest ROCm CCL development package on Debian-based systems like Ubuntu.

```bash
sudo apt install amdrocm-ccl-dev
```

--------------------------------

### Create and Enter Build Directory

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Create a build directory within the hipCUB root and navigate into it.

```bash
mkdir build
cd build
```

--------------------------------

### List rmake.py options

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Displays the complete list of available options for the rmake.py script.

```bash
python rmake.py --help
```

--------------------------------

### Build hipCUB clients with rmake.py

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Uses the -c option to build all clients, including unit tests.

```bash
python rmake.py -c
```

--------------------------------

### Generate Makefile with CMake

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Run the cmake command to generate the makefile, optionally passing build configuration options.

```bash
cmake ../. [-D<OPTION1=VALUE1> [-D<OPTION2=VALUE2>] ...]
```

--------------------------------

### Build hipCUB

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Compile the project using the generated makefile with the make command.

```bash
make -j4
```

--------------------------------

### Generate Packages

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Optionally generate tar, zip, and deb packages after the build process is complete.

```bash
make package
```

--------------------------------

### hipcub/grid/grid_queue.hpp header content

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/grid__queue_8hpp_source.html

The header file uses preprocessor directives to include either the rocprim-based grid queue for AMD platforms or the CUB-based grid queue for NVIDIA platforms.

```cpp
/******************************************************************************
* Copyright (c) 2011, Duane Merrill. All rights reserved.
* Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2021-2025, 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_GRID_GRID_QUEUE_HPP_
#define HIPCUB_GRID_GRID_QUEUE_HPP_

#ifdef __HIP_PLATFORM_AMD__
#include "../backend/rocprim/grid/grid_queue.hpp" // IWYU pragma: export
#elif defined(__HIP_PLATFORM_NVIDIA__)
#include "../config.hpp"
#include <cub/grid/grid_queue.cuh> // IWYU pragma: export
#endif

#endif // HIPCUB_GRID_GRID_QUEUE_HPP_
```

--------------------------------

### hipcub/block/block_load.hpp header guard and platform includes

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/block__load_8hpp_source.html

Header guard for hipcub/block/block_load.hpp. Includes the ROCm backend implementation for AMD platforms and the CUB implementation for NVIDIA platforms, with the necessary config header for NVIDIA.

```cpp
#ifndef HIPCUB_BLOCK_BLOCK_LOAD_HPP_
#define HIPCUB_BLOCK_BLOCK_LOAD_HPP_

#ifdef __HIP_PLATFORM_AMD__
#include "../backend/rocprim/block/block_load.hpp" // IWYU pragma: export
#elif defined(__HIP_PLATFORM_NVIDIA__)
#include "../config.hpp"
#include <cub/block/block_load.cuh> // IWYU pragma: export
#endif

#endif // HIPCUB_BLOCK_BLOCK_LOAD_HPP_
```

--------------------------------

### Clone hipCUB with sparse checkout

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Use sparse checkout to clone only the hipCUB project from the ROCm libraries repository. This avoids downloading the entire repository and is required before checking out a specific branch.

```bash
git clone --no-checkout --filter=blob:none https://github.com/ROCm/rocm-libraries.git
cd rocm-libraries
git sparse-checkout init --cone
git sparse-checkout set projects/hipcub
```

--------------------------------

### hipcub/block/block_shuffle.hpp Header Content

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/block__shuffle_8hpp_source.html

The header file uses preprocessor directives to include the appropriate backend implementation based on the target HIP platform.

```cpp
/******************************************************************************
* Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
* Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2021-2025, 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_BLOCK_BLOCK_SHUFFLE_HPP_
#define HIPCUB_BLOCK_BLOCK_SHUFFLE_HPP_

#ifdef __HIP_PLATFORM_AMD__
#include "../backend/rocprim/block/block_shuffle.hpp" // IWYU pragma: export
#elif defined(__HIP_PLATFORM_NVIDIA__)
#include "../config.hpp"
#include <cub/block/block_shuffle.cuh> // IWYU pragma: export
#endif

#endif // HIPCUB_BLOCK_BLOCK_SHUFFLE_HPP_
```

--------------------------------

### Set Compiler Environment Variables

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/build.html

Set the CXX and CMAKE_CXX_COMPILER environment variables to the appropriate compiler path before running CMake.

```bash
CXX=amdclang++
CMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++
```

--------------------------------

### hipcub/config.hpp Source Code

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

The complete source code for hipcub/config.hpp, defining the library's namespace, versioning, and platform-specific configurations.

```cpp
/******************************************************************************
* Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
* Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
* Modifications Copyright (c) 2019-2025, 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_CONFIG_HPP_
#define HIPCUB_CONFIG_HPP_

#include <hip/hip_runtime.h>

// Version
#include "hipcub_version.hpp" // IWYU pragma: export

#define HIPCUB_NAMESPACE hipcub

// Inline namespace (e.g. HIPCUB_300400_NS where 300400 is the hipCUB version) is used to
// eliminate issues when shared libraries are built with different versions of hipCUB so they may
// have symbols with the same name but different content.
// HIPCUB_DISABLE_INLINE_NAMESPACE can be defined to disable inline namespaces (the old behavior).
// HIPCUB_INLINE_NAMESPACE can be defined to override the standard inline namespace name.
// Additionally, all kernels have hidden visibility.
// For rocPRIM backend, see rocprim/config.hpp for similar definitions.
#if defined(DOXYGEN_SHOULD_SKIP_THIS) || defined(HIPCUB_DISABLE_INLINE_NAMESPACE)
#define HIPCUB_INLINE_NAMESPACE
#define BEGIN_HIPCUB_INLINE_NAMESPACE
#define END_HIPCUB_INLINE_NAMESPACE
#else
#define HIPCUB_CONCAT_(SEP, A, B) A##SEP##B
#define HIPCUB_CONCAT(SEP, A, B) HIPCUB_CONCAT_(SEP, A, B)

#ifndef HIPCUB_INLINE_NAMESPACE
#define HIPCUB_INLINE_NAMESPACE \
HIPCUB_CONCAT(_, HIPCUB, HIPCUB_CONCAT(_, HIPCUB_VERSION, NS))
#endif
#define BEGIN_HIPCUB_INLINE_NAMESPACE \
inline namespace HIPCUB_INLINE_NAMESPACE \
{
#define END_HIPCUB_INLINE_NAMESPACE } /* inline namespace */
#endif

#define BEGIN_HIPCUB_NAMESPACE \
namespace HIPCUB_NAMESPACE \
{ \
BEGIN_HIPCUB_INLINE_NAMESPACE

#define END_HIPCUB_NAMESPACE \
END_HIPCUB_INLINE_NAMESPACE \
} /* namespace hipcub */

#ifdef __HIP_PLATFORM_AMD__
#define HIPCUB_ROCPRIM_API 1
#define HIPCUB_RUNTIME_FUNCTION __host__

#include <rocprim/device/config_types.hpp>
#include <rocprim/intrinsics/arch.hpp>
#include <rocprim/intrinsics/thread.hpp>

BEGIN_HIPCUB_NAMESPACE
namespace detail
{
inline unsigned int host_warp_size_wrapper()
{
int device_id = 0;
unsigned int host_warp_size = 0;
hipError_t error = hipGetDevice(&device_id);
if(error != hipSuccess)
{
fprintf(stderr, "HIP error: %d line: %d: %s\n", error, __LINE__, hipGetErrorString(error));
fflush(stderr);
}
if(::rocprim::host_warp_size(device_id, host_warp_size) != hipSuccess)
{
return 0u;
}
return host_warp_size;
}
} // namespace detail
END_HIPCUB_NAMESPACE
#include <rocprim/intrinsics/arch.hpp>

#define HIPCUB_WARP_THREADS ::rocprim::warp_size()
// HIPCUB (and CUB) don't have a method to express min and max warp size.
#define HIPCUB_DEVICE_WARP_THREADS ::rocprim::arch::wavefront::max_size()
```

--------------------------------

### HIPCUB backend macro definitions (rocPRIM and CUB)

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB macros for the rocPRIM and CUB backends. The rocPRIM branch sets HIPCUB_HOST_WARP_THREADS to a runtime wrapper and ignores HIPCUB_ARCH; the NVIDIA branch includes cub/util_arch.cuh and maps HIPCUB macros to CUB equivalents.

```cpp
#define HIPCUB_HOST_WARP_THREADS ::hipcub::detail::host_warp_size_wrapper()
#define HIPCUB_ARCH 1 // ignored with rocPRIM backend
#elif defined(__HIP_PLATFORM_NVIDIA__)
#define HIPCUB_CUB_API 1
#define HIPCUB_RUNTIME_FUNCTION CUB_RUNTIME_FUNCTION

#include <cub/util_arch.cuh>
#define HIPCUB_WARP_THREADS CUB_PTX_WARP_THREADS
#define HIPCUB_DEVICE_WARP_THREADS CUB_PTX_WARP_THREADS
#define HIPCUB_HOST_WARP_THREADS CUB_PTX_WARP_THREADS
#define HIPCUB_ARCH CUB_PTX_ARCH
BEGIN_HIPCUB_NAMESPACE
using namespace cub;
END_HIPCUB_NAMESPACE
#endif
```

--------------------------------

### HipcubDebug and HipcubLog macros

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines the HipcubDebug and HipcubLog macros that wrap the Debug and Log functions, passing the current file and line.

```cpp
#ifndef HipcubDebug
#define HipcubDebug(e) ::hipcub::Debug((hipError_t)(e), __FILE__, __LINE__)
#endif

#ifndef HipcubLog
#define HipcubLog(msg) ::hipcub::Log(msg, __FILE__, __LINE__)
#endif
```

--------------------------------

### HIPCUB_DEBUG_SYNC documentation-only definition

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB_DEBUG_SYNC only for documentation purposes when DOXYGEN_SHOULD_SKIP_THIS is defined.

```cpp
#ifdef DOXYGEN_SHOULD_SKIP_THIS // Documentation only

#define HIPCUB_DEBUG_SYNC

#endif // DOXYGEN_SHOULD_SKIP_THIS
```

--------------------------------

### CUB_DEBUG_SYNC definition

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines CUB_DEBUG_SYNC when HIPCUB_CUB_API and HIPCUB_DEBUG_SYNC are defined and CUB_DEBUG_SYNC is not already defined.

```cpp
#if defined(HIPCUB_CUB_API) && defined(HIPCUB_DEBUG_SYNC) && !defined(CUB_DEBUG_SYNC)
#define CUB_DEBUG_SYNC
#endif
```

--------------------------------

### HIPCUB_DETAIL_DEBUG_SYNC_VALUE definition

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB_DETAIL_DEBUG_SYNC_VALUE as true or false based on HIPCUB_DEBUG_SYNC, only when HIPCUB_ROCPRIM_API is defined.

```cpp
#ifdef HIPCUB_ROCPRIM_API
// TODO C++17: use an inline constexpr variable
#ifdef HIPCUB_DEBUG_SYNC
#define HIPCUB_DETAIL_DEBUG_SYNC_VALUE true
#else
#define HIPCUB_DETAIL_DEBUG_SYNC_VALUE false
#endif
#endif // HIPCUB_ROCPRIM_API
```

--------------------------------

### HIPCUB_DEBUG_SYNC conditional definition

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB_DEBUG_SYNC when it is not already defined and any of CUB_DEBUG_SYNC, CUB_DEBUG_HOST_ASSERTIONS, CUB_DEBUG_DEVICE_ASSERTIONS, or CUB_DEBUG_ALL is defined.

```cpp
#if !defined(HIPCUB_DEBUG_SYNC) \
&& (defined(CUB_DEBUG_SYNC) || defined(CUB_DEBUG_HOST_ASSERTIONS) \
|| defined(CUB_DEBUG_DEVICE_ASSERTIONS) || defined(CUB_DEBUG_ALL))
#define HIPCUB_DEBUG_SYNC
#endif
```

--------------------------------

### hipcub::Log function

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines the Log function that prints a message with filename and line number to stdout. Use the HipcubLog macro instead of calling this directly.

```cpp
inline void Log(const char* message, const char* filename, int line)
{
printf("hipcub: %s [%s:%d]\n", message, filename, line);
}
```

--------------------------------

### HIPCUB warp size constants

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines warp size constants: 32, 64, and the maximum warp size (64).

```cpp
#define HIPCUB_WARP_SIZE_32 32u
#define HIPCUB_WARP_SIZE_64 64u
#define HIPCUB_MAX_WARP_SIZE HIPCUB_WARP_SIZE_64
```

--------------------------------

### HIPCUB_IF_CONSTEXPR macro

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB_IF_CONSTEXPR to 'constexpr' when if constexpr is available (C++17 or MSVC), otherwise empty. Includes a comment about MSVC behavior.

```cpp
#if __cpp_if_constexpr
#define HIPCUB_IF_CONSTEXPR constexpr
#else
#if defined(_MSC_VER) && !defined(__clang__)
// MSVC (and not Clang pretending to be MSVC) unconditionally exposes if constexpr (even in C++14 mode),
// moreover it triggers warning C4127 (conditional expression is constant) when not using it. nvcc will
// be calling cl.exe for host-side codegen.
#define HIPCUB_IF_CONSTEXPR constexpr
#else
#define HIPCUB_IF_CONSTEXPR
#endif
#endif
```

--------------------------------

### HIPCUB execution specifier macros

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB macros for CUDA execution specifiers and force-inline.

```cpp
#define HIPCUB_HOST __host__
#define HIPCUB_DEVICE __device__
#define HIPCUB_HOST_DEVICE __host__ __device__
#define HIPCUB_FORCEINLINE __forceinline__
#define HIPCUB_SHARED_MEMORY __shared__
```

--------------------------------

### Clang warning suppression macros

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines macros to suppress Clang warnings. When __clang__ is defined, the macros use _Pragma to push, ignore, and pop diagnostic states; otherwise they expand to nothing.

```cpp
#ifdef __clang__
#define HIPCUB_PRAGMA_TO_STR(x) _Pragma(#x)
#define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH _Pragma("clang diagnostic push")
#define HIPCUB_CLANG_SUPPRESS_WARNING(w) HIPCUB_PRAGMA_TO_STR(clang diagnostic ignored w)
#define HIPCUB_CLANG_SUPPRESS_WARNING_POP _Pragma("clang diagnostic pop")
#define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w) \
HIPCUB_CLANG_SUPPRESS_WARNING_PUSH HIPCUB_CLANG_SUPPRESS_WARNING(w)
#else // __clang__
#define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH
#define HIPCUB_CLANG_SUPPRESS_WARNING(w)
#define HIPCUB_CLANG_SUPPRESS_WARNING_POP
#define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w)
#endif // __clang__
```

--------------------------------

### HIPCUB_STDERR definition

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines HIPCUB_STDERR when DEBUG or _DEBUG is defined and HIPCUB_STDERR is not already defined.

```cpp
#if (defined(DEBUG) || defined(_DEBUG)) && !defined(HIPCUB_STDERR)
#define HIPCUB_STDERR
#endif
```

--------------------------------

### hipcub::Debug function

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines the Debug function that prints an error message to stderr if HIPCUB_STDERR is defined and the error is non-zero. It returns the error unchanged. Use the HipcubDebug macro instead of calling this directly.

```cpp
inline
hipError_t Debug(
hipError_t error,
const char* filename,
int line)
{
(void)filename;
(void)line;
#ifdef HIPCUB_STDERR
if (error)
{
fprintf(stderr, "HIP error %d [%s, %d]: %s\n", error, filename, line, hipGetErrorString(error));
fflush(stderr);
}
#endif
return error;
}
```

--------------------------------

### ROCm CCL package naming format

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/install.html

The naming convention for the ROCm CCL package, which includes hipCUB, rocPRIM, and rocThrust.

```text
amdrocm-ccl<dev/devel><rocm_version>-<llvm_target>
```

--------------------------------

### HIPCUB deprecated warning suppression macros

Source: https://rocm.docs.amd.com/projects/hipCUB/en/latest/doxygen/html/config_8hpp_source.html

Defines macros to suppress deprecated warnings in Clang, using the previously defined warning suppression macros.

```cpp
#define HIPCUB_CLANG_SUPPRESS_DEPRECATED_PUSH \
HIPCUB_CLANG_SUPPRESS_WARNING_PUSH \
HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated") \
HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated-declarations")
#define HIPCUB_CLANG_SUPPRESS_DEPRECATED_POP HIPCUB_CLANG_SUPPRESS_WARNING_POP
```
Tokens:
Raw