Ctrl+K

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-2024, 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 #define HIPCUB_NAMESPACE hipcub
36 
37 #define BEGIN_HIPCUB_NAMESPACE \
38  namespace hipcub {
39 
40 #define END_HIPCUB_NAMESPACE \
41  } /* hipcub */
42 
43 #ifdef __HIP_PLATFORM_AMD__
44  #define HIPCUB_ROCPRIM_API 1
45  #define HIPCUB_RUNTIME_FUNCTION __host__
46 
47  #include <rocprim/device/config_types.hpp>
48 
49 BEGIN_HIPCUB_NAMESPACE
50 namespace detail
51 {
52 inline unsigned int host_warp_size_wrapper()
53 {
54  int device_id = 0;
55  unsigned int host_warp_size = 0;
56  hipError_t error = hipGetDevice(&device_id);
57  if(error != hipSuccess)
58  {
59  fprintf(stderr, "HIP error: %d line: %d: %s\n", error, __LINE__, hipGetErrorString(error));
60  fflush(stderr);
61  }
62  if(::rocprim::host_warp_size(device_id, host_warp_size) != hipSuccess)
63  {
64  return 0u;
65  }
66  return host_warp_size;
67 }
68 } // namespace detail
69 END_HIPCUB_NAMESPACE
70 
71  #define HIPCUB_WARP_THREADS ::rocprim::warp_size()
72  #define HIPCUB_DEVICE_WARP_THREADS ::rocprim::device_warp_size()
73  #define HIPCUB_HOST_WARP_THREADS ::hipcub::detail::host_warp_size_wrapper()
74  #define HIPCUB_ARCH 1 // ignored with rocPRIM backend
75 #elif defined(__HIP_PLATFORM_NVIDIA__)
76  #define HIPCUB_CUB_API 1
77  #define HIPCUB_RUNTIME_FUNCTION CUB_RUNTIME_FUNCTION
78 
79  #include <cub/util_arch.cuh>
80  #define HIPCUB_WARP_THREADS CUB_PTX_WARP_THREADS
81  #define HIPCUB_DEVICE_WARP_THREADS CUB_PTX_WARP_THREADS
82  #define HIPCUB_HOST_WARP_THREADS CUB_PTX_WARP_THREADS
83  #define HIPCUB_ARCH CUB_PTX_ARCH
84  BEGIN_HIPCUB_NAMESPACE
85  using namespace cub;
86  END_HIPCUB_NAMESPACE
87 #endif
88 
90 #define HIPCUB_WARP_SIZE_32 32u
91 #define HIPCUB_WARP_SIZE_64 64u
92 #define HIPCUB_MAX_WARP_SIZE HIPCUB_WARP_SIZE_64
93 
94 #define HIPCUB_HOST __host__
95 #define HIPCUB_DEVICE __device__
96 #define HIPCUB_HOST_DEVICE __host__ __device__
97 #define HIPCUB_SHARED_MEMORY __shared__
98 
99 // Helper macros to disable warnings in clang
100 #ifdef __clang__
101 #define HIPCUB_PRAGMA_TO_STR(x) _Pragma(#x)
102 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH _Pragma("clang diagnostic push")
103 #define HIPCUB_CLANG_SUPPRESS_WARNING(w) HIPCUB_PRAGMA_TO_STR(clang diagnostic ignored w)
104 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP _Pragma("clang diagnostic pop")
105 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w) \
106  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH HIPCUB_CLANG_SUPPRESS_WARNING(w)
107 #else // __clang__
108 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH
109 #define HIPCUB_CLANG_SUPPRESS_WARNING(w)
110 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP
111 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w)
112 #endif // __clang__
113 
115 #if (defined(DEBUG) || defined(_DEBUG)) && !defined(HIPCUB_STDERR)
116  #define HIPCUB_STDERR
117 #endif
118 
119 BEGIN_HIPCUB_NAMESPACE
120 
125 inline
126 hipError_t Debug(
127  hipError_t error,
128  const char* filename,
129  int line)
130 {
131  (void)filename;
132  (void)line;
133 #ifdef HIPCUB_STDERR
134  if (error)
135  {
136  fprintf(stderr, "HIP error %d [%s, %d]: %s\n", error, filename, line, hipGetErrorString(error));
137  fflush(stderr);
138  }
139 #endif
140  return error;
141 }
142 
146 inline void Log(const char* message, const char* filename, int line)
147 {
148  printf("hipcub: %s [%s:%d]\n", message, filename, line);
149 }
150 
151 END_HIPCUB_NAMESPACE
152 
153 #ifndef HipcubDebug
154  #define HipcubDebug(e) ::hipcub::Debug((hipError_t)(e), __FILE__, __LINE__)
155 #endif
156 
157 #ifndef HipcubLog
158  #define HipcubLog(msg) ::hipcub::Log(msg, __FILE__, __LINE__)
159 #endif
160 
161 #if __cpp_if_constexpr
162  #define HIPCUB_IF_CONSTEXPR constexpr
163 #else
164  #if defined(_MSC_VER) && !defined(__clang__)
165  // MSVC (and not Clang pretending to be MSVC) unconditionally exposes if constexpr (even in C++14 mode),
166  // moreover it triggers warning C4127 (conditional expression is constant) when not using it. nvcc will
167  // be calling cl.exe for host-side codegen.
168  #define HIPCUB_IF_CONSTEXPR constexpr
169  #else
170  #define HIPCUB_IF_CONSTEXPR
171  #endif
172 #endif
173 
174 #ifdef DOXYGEN_SHOULD_SKIP_THIS // Documentation only
175 
180  #define HIPCUB_DEBUG_SYNC
181 
182 #endif // DOXYGEN_SHOULD_SKIP_THIS
183 
184 #if defined(HIPCUB_CUB_API) && defined(HIPCUB_DEBUG_SYNC) && !defined(CUB_DEBUG_SYNC)
185  #define CUB_DEBUG_SYNC
186 #endif
187 
188 #if !defined(HIPCUB_DEBUG_SYNC) \
189  && (defined(CUB_DEBUG_SYNC) || defined(CUB_DEBUG_HOST_ASSERTIONS) \
190  || defined(CUB_DEBUG_DEVICE_ASSERTIONS) || defined(CUB_DEBUG_ALL))
191  #define HIPCUB_DEBUG_SYNC
192 #endif
193 
194 #ifdef HIPCUB_ROCPRIM_API
195  // TODO C++17: use an inline constexpr variable
196  #ifdef HIPCUB_DEBUG_SYNC
197  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE true
198  #else
199  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE false
200  #endif
201 #endif // HIPCUB_ROCPRIM_API
202 
203 #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:126
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:146