/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-6.0.0/hipcub/include/hipcub/config.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-6.0.0/hipcub/include/hipcub/config.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-6.0.0/hipcub/include/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-2023, 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/intrinsics/thread.hpp>
48  #define HIPCUB_WARP_THREADS ::rocprim::warp_size()
49  #define HIPCUB_DEVICE_WARP_THREADS ::rocprim::device_warp_size()
50  #define HIPCUB_HOST_WARP_THREADS ::rocprim::host_warp_size()
51  #define HIPCUB_ARCH 1 // ignored with rocPRIM backend
52 #elif defined(__HIP_PLATFORM_NVIDIA__)
53  #define HIPCUB_CUB_API 1
54  #define HIPCUB_RUNTIME_FUNCTION CUB_RUNTIME_FUNCTION
55 
56  #include <cub/util_arch.cuh>
57  #define HIPCUB_WARP_THREADS CUB_PTX_WARP_THREADS
58  #define HIPCUB_DEVICE_WARP_THREADS CUB_PTX_WARP_THREADS
59  #define HIPCUB_HOST_WARP_THREADS CUB_PTX_WARP_THREADS
60  #define HIPCUB_ARCH CUB_PTX_ARCH
61  BEGIN_HIPCUB_NAMESPACE
62  using namespace cub;
63  END_HIPCUB_NAMESPACE
64 #endif
65 
67 #define HIPCUB_WARP_SIZE_32 32u
68 #define HIPCUB_WARP_SIZE_64 64u
69 #define HIPCUB_MAX_WARP_SIZE HIPCUB_WARP_SIZE_64
70 
71 #define HIPCUB_HOST __host__
72 #define HIPCUB_DEVICE __device__
73 #define HIPCUB_HOST_DEVICE __host__ __device__
74 #define HIPCUB_SHARED_MEMORY __shared__
75 
76 // Helper macros to disable warnings in clang
77 #ifdef __clang__
78 #define HIPCUB_PRAGMA_TO_STR(x) _Pragma(#x)
79 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH _Pragma("clang diagnostic push")
80 #define HIPCUB_CLANG_SUPPRESS_WARNING(w) HIPCUB_PRAGMA_TO_STR(clang diagnostic ignored w)
81 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP _Pragma("clang diagnostic pop")
82 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w) \
83  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH HIPCUB_CLANG_SUPPRESS_WARNING(w)
84 #else // __clang__
85 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH
86 #define HIPCUB_CLANG_SUPPRESS_WARNING(w)
87 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP
88 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w)
89 #endif // __clang__
90 
91 BEGIN_HIPCUB_NAMESPACE
92 
94 #if (defined(DEBUG) || defined(_DEBUG)) && !defined(HIPCUB_STDERR)
95  #define HIPCUB_STDERR
96 #endif
97 
98 inline
99 hipError_t Debug(
100  hipError_t error,
101  const char* filename,
102  int line)
103 {
104  (void)filename;
105  (void)line;
106 #ifdef HIPCUB_STDERR
107  if (error)
108  {
109  fprintf(stderr, "HIP error %d [%s, %d]: %s\n", error, filename, line, hipGetErrorString(error));
110  fflush(stderr);
111  }
112 #endif
113  return error;
114 }
115 
116 #ifndef HipcubDebug
117  #define HipcubDebug(e) hipcub::Debug((hipError_t) (e), __FILE__, __LINE__)
118 #endif
119 
120 #if __cpp_if_constexpr
121  #define HIPCUB_IF_CONSTEXPR constexpr
122 #else
123  #if defined(_MSC_VER) && !defined(__clang__)
124  // MSVC (and not Clang pretending to be MSVC) unconditionally exposes if constexpr (even in C++14 mode),
125  // moreover it triggers warning C4127 (conditional expression is constant) when not using it. nvcc will
126  // be calling cl.exe for host-side codegen.
127  #define HIPCUB_IF_CONSTEXPR constexpr
128  #else
129  #define HIPCUB_IF_CONSTEXPR
130  #endif
131 #endif
132 
133 END_HIPCUB_NAMESPACE
134 
135 #endif // HIPCUB_CONFIG_HPP_