hipcub/config.hpp Source File

hipcub/config.hpp Source File#

hipCUB: hipcub/config.hpp Source File
config.hpp
1 /******************************************************************************
2  * Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4  * Modifications Copyright (c) 2019-2026, Advanced Micro Devices, Inc. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  *
28  ******************************************************************************/
29 
30 #ifndef HIPCUB_CONFIG_HPP_
31 #define HIPCUB_CONFIG_HPP_
32 
33 #include <hip/hip_runtime.h>
34 
35 // Version
36 #include "hipcub_version.hpp" // IWYU pragma: export
37 
38 // Manage std implementation
39 #include "libcxx.hpp" // IWYU pragma: export
40 
41 // For _CCCL_IMPLICIT_SYSTEM_HEADER
42 #if _HIPCUB_HAS_DEVICE_SYSTEM_STD
43  #include _HIPCUB_LIBCXX_INCLUDE(__cccl_config) // IWYU pragma: export
44 #endif
45 
46 #define HIPCUB_NAMESPACE hipcub
47 
48 // Inline namespace (e.g. HIPCUB_300400_NS where 300400 is the hipCUB version) is used to
49 // eliminate issues when shared libraries are built with different versions of hipCUB so they may
50 // have symbols with the same name but different content.
51 // HIPCUB_DISABLE_INLINE_NAMESPACE can be defined to disable inline namespaces (the old behavior).
52 // HIPCUB_INLINE_NAMESPACE can be defined to override the standard inline namespace name.
53 // Additionally, all kernels have hidden visibility.
54 // For rocPRIM backend, see rocprim/config.hpp for similar definitions.
55 #if defined(DOXYGEN_SHOULD_SKIP_THIS) || defined(HIPCUB_DISABLE_INLINE_NAMESPACE)
56  #define HIPCUB_INLINE_NAMESPACE
57  #define BEGIN_HIPCUB_INLINE_NAMESPACE
58  #define END_HIPCUB_INLINE_NAMESPACE
59 #else
60  #define HIPCUB_CONCAT_(SEP, A, B) A##SEP##B
61  #define HIPCUB_CONCAT(SEP, A, B) HIPCUB_CONCAT_(SEP, A, B)
62 
63  #ifndef HIPCUB_INLINE_NAMESPACE
64  #define HIPCUB_INLINE_NAMESPACE \
65  HIPCUB_CONCAT(_, HIPCUB, HIPCUB_CONCAT(_, HIPCUB_VERSION, NS))
66  #endif
67  #define BEGIN_HIPCUB_INLINE_NAMESPACE \
68  inline namespace HIPCUB_INLINE_NAMESPACE \
69  {
70  #define END_HIPCUB_INLINE_NAMESPACE } /* inline namespace */
71 #endif
72 
73 #define BEGIN_HIPCUB_NAMESPACE \
74  namespace HIPCUB_NAMESPACE \
75  { \
76  BEGIN_HIPCUB_INLINE_NAMESPACE
77 
78 #define END_HIPCUB_NAMESPACE \
79  END_HIPCUB_INLINE_NAMESPACE \
80  } /* namespace hipcub */
81 
82 #ifdef __HIP_PLATFORM_AMD__
83  #define HIPCUB_ROCPRIM_API 1
84  #define HIPCUB_RUNTIME_FUNCTION __host__
85 
86  #include <rocprim/device/config_types.hpp>
87  #include <rocprim/intrinsics/arch.hpp>
88  #include <rocprim/intrinsics/thread.hpp>
89 
90 BEGIN_HIPCUB_NAMESPACE
91 namespace detail
92 {
93 inline unsigned int host_warp_size_wrapper()
94 {
95  int device_id = 0;
96  unsigned int host_warp_size = 0;
97  hipError_t error = hipGetDevice(&device_id);
98  if(error != hipSuccess)
99  {
100  fprintf(stderr, "HIP error: %d line: %d: %s\n", error, __LINE__, hipGetErrorString(error));
101  fflush(stderr);
102  }
103  if(::rocprim::host_warp_size(device_id, host_warp_size) != hipSuccess)
104  {
105  return 0u;
106  }
107  return host_warp_size;
108 }
109 } // namespace detail
110 END_HIPCUB_NAMESPACE
111  #include <rocprim/intrinsics/arch.hpp>
112 
113  #define HIPCUB_WARP_THREADS ::rocprim::warp_size()
114  // HIPCUB (and CUB) don't have a method to express min and max warp size.
115  #define HIPCUB_DEVICE_WARP_THREADS ::rocprim::arch::wavefront::max_size()
116  #define HIPCUB_HOST_WARP_THREADS ::hipcub::detail::host_warp_size_wrapper()
117  #define HIPCUB_ARCH 1 // ignored with rocPRIM backend
118 #elif defined(__HIP_PLATFORM_NVIDIA__)
119  #define HIPCUB_CUB_API 1
120  #define HIPCUB_RUNTIME_FUNCTION CUB_RUNTIME_FUNCTION
121 
122  #include <cub/util_arch.cuh>
123  #include <cuda/std/limits>
124  #include <cuda/std/type_traits>
125  #define HIPCUB_WARP_THREADS CUB_PTX_WARP_THREADS
126  #define HIPCUB_DEVICE_WARP_THREADS CUB_PTX_WARP_THREADS
127  #define HIPCUB_HOST_WARP_THREADS CUB_PTX_WARP_THREADS
128  #define HIPCUB_ARCH CUB_PTX_ARCH
129  BEGIN_HIPCUB_NAMESPACE
130  using namespace cub;
131  END_HIPCUB_NAMESPACE
132 #endif
133 
135 #define HIPCUB_WARP_SIZE_32 32u
136 #define HIPCUB_WARP_SIZE_64 64u
137 #define HIPCUB_MAX_WARP_SIZE HIPCUB_WARP_SIZE_64
138 
139 #define HIPCUB_HOST __host__
140 #define HIPCUB_DEVICE __device__
141 #define HIPCUB_HOST_DEVICE __host__ __device__
142 #define HIPCUB_FORCEINLINE __forceinline__
143 #define HIPCUB_SHARED_MEMORY __shared__
144 
145 // Helper macros to disable warnings in clang
146 #ifdef __clang__
147 #define HIPCUB_PRAGMA_TO_STR(x) _Pragma(#x)
148 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH _Pragma("clang diagnostic push")
149 #define HIPCUB_CLANG_SUPPRESS_WARNING(w) HIPCUB_PRAGMA_TO_STR(clang diagnostic ignored w)
150 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP _Pragma("clang diagnostic pop")
151 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w) \
152  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH HIPCUB_CLANG_SUPPRESS_WARNING(w)
153 #else // __clang__
154 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH
155 #define HIPCUB_CLANG_SUPPRESS_WARNING(w)
156 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP
157 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w)
158 #endif // __clang__
159 
160 #define HIPCUB_CLANG_SUPPRESS_DEPRECATED_PUSH \
161  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH \
162  HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated") \
163  HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated-declarations")
164 #define HIPCUB_CLANG_SUPPRESS_DEPRECATED_POP HIPCUB_CLANG_SUPPRESS_WARNING_POP
165 
167 #if (defined(DEBUG) || defined(_DEBUG)) && !defined(HIPCUB_STDERR)
168  #define HIPCUB_STDERR
169 #endif
170 
171 BEGIN_HIPCUB_NAMESPACE
172 
177 inline
178 hipError_t Debug(
179  hipError_t error,
180  const char* filename,
181  int line)
182 {
183  (void)filename;
184  (void)line;
185 #ifdef HIPCUB_STDERR
186  if (error)
187  {
188  fprintf(stderr, "HIP error %d [%s, %d]: %s\n", error, filename, line, hipGetErrorString(error));
189  fflush(stderr);
190  }
191 #endif
192  return error;
193 }
194 
198 inline void Log(const char* message, const char* filename, int line)
199 {
200  printf("hipcub: %s [%s:%d]\n", message, filename, line);
201 }
202 
203 END_HIPCUB_NAMESPACE
204 
205 #ifndef HipcubDebug
206  #define HipcubDebug(e) ::hipcub::Debug((hipError_t)(e), __FILE__, __LINE__)
207 #endif
208 
209 #ifndef HipcubLog
210  #define HipcubLog(msg) ::hipcub::Log(msg, __FILE__, __LINE__)
211 #endif
212 
213 #ifdef DOXYGEN_SHOULD_SKIP_THIS // Documentation only
214 
219  #define HIPCUB_DEBUG_SYNC
220 
221 #endif // DOXYGEN_SHOULD_SKIP_THIS
222 
223 #if defined(HIPCUB_CUB_API) && defined(HIPCUB_DEBUG_SYNC) && !defined(CUB_DEBUG_SYNC)
224  #define CUB_DEBUG_SYNC
225 #endif
226 
227 #if !defined(HIPCUB_DEBUG_SYNC) \
228  && (defined(CUB_DEBUG_SYNC) || defined(CUB_DEBUG_HOST_ASSERTIONS) \
229  || defined(CUB_DEBUG_DEVICE_ASSERTIONS) || defined(CUB_DEBUG_ALL))
230  #define HIPCUB_DEBUG_SYNC
231 #endif
232 
233 #ifdef HIPCUB_ROCPRIM_API
234  // TODO C++17: use an inline constexpr variable
235  #ifdef HIPCUB_DEBUG_SYNC
236  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE true
237  #else
238  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE false
239  #endif
240 #endif // HIPCUB_ROCPRIM_API
241 
242 // This API needs to be deprecated once libhipcxx is available.
243 #if !defined(_CCCL_PRAGMA_UNROLL_FULL)
244  #define _CCCL_PRAGMA_UNROLL_FULL() _Pragma("unroll")
245 #endif // !defined(_CCCL_PRAGMA_UNROLL_FULL)
246 
247 #if !defined(_CCCL_PRAGMA_NOUNROLL)
248  #define _CCCL_PRAGMA_NOUNROLL() _Pragma("nounroll")
249 #endif // !defined(_CCCL_PRAGMA_NOUNROLL)
250 
251 #endif // HIPCUB_CONFIG_HPP_
hipcub::Debug
hipError_t Debug(hipError_t error, const char *filename, int line)
Don't use this function directly, but via the HipcubDebug macro instead. If error is not hipSuccess,...
Definition: config.hpp:178
hipcub::Log
void Log(const char *message, const char *filename, int line)
Don't use this function directly, but via the HipcubLog macro instead. Prints the provided message co...
Definition: config.hpp:198