hipcub/config.hpp Source File

• hipcub

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_FORCEINLINE __forceinline__

98 #define HIPCUB_SHARED_MEMORY __shared__

99 

100 // Helper macros to disable warnings in clang

101 #ifdef __clang__

102 #define HIPCUB_PRAGMA_TO_STR(x) _Pragma(#x)

103 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH _Pragma("clang diagnostic push")

104 #define HIPCUB_CLANG_SUPPRESS_WARNING(w) HIPCUB_PRAGMA_TO_STR(clang diagnostic ignored w)

105 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP _Pragma("clang diagnostic pop")

106 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w) \

107  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH HIPCUB_CLANG_SUPPRESS_WARNING(w)

108 #else // __clang__

109 #define HIPCUB_CLANG_SUPPRESS_WARNING_PUSH

110 #define HIPCUB_CLANG_SUPPRESS_WARNING(w)

111 #define HIPCUB_CLANG_SUPPRESS_WARNING_POP

112 #define HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH(w)

113 #endif // __clang__

114 

115 #define HIPCUB_CLANG_SUPPRESS_DEPRECATED_PUSH \

116  HIPCUB_CLANG_SUPPRESS_WARNING_PUSH \

117  HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated") \

118  HIPCUB_CLANG_SUPPRESS_WARNING("-Wdeprecated-declarations")

119 #define HIPCUB_CLANG_SUPPRESS_DEPRECATED_POP HIPCUB_CLANG_SUPPRESS_WARNING_POP

120 

122 #if (defined(DEBUG) || defined(_DEBUG)) && !defined(HIPCUB_STDERR)

123  #define HIPCUB_STDERR

124 #endif

125 

126 BEGIN_HIPCUB_NAMESPACE

127 

132 inline

133 hipError_t Debug(

134  hipError_t error,

135  const char* filename,

136  int line)

137 {

138  (void)filename;

139  (void)line;

140 #ifdef HIPCUB_STDERR

141  if (error)

142  {

143  fprintf(stderr, "HIP error %d [%s, %d]: %s\n", error, filename, line, hipGetErrorString(error));

144  fflush(stderr);

145  }

146 #endif

147  return error;

148 }

149 

153 inline void Log(const char* message, const char* filename, int line)

154 {

155  printf("hipcub: %s [%s:%d]\n", message, filename, line);

156 }

157 

158 END_HIPCUB_NAMESPACE

159 

160 #ifndef HipcubDebug

161  #define HipcubDebug(e) ::hipcub::Debug((hipError_t)(e), __FILE__, __LINE__)

162 #endif

163 

164 #ifndef HipcubLog

165  #define HipcubLog(msg) ::hipcub::Log(msg, __FILE__, __LINE__)

166 #endif

167 

168 #if __cpp_if_constexpr

169  #define HIPCUB_IF_CONSTEXPR constexpr

170 #else

171  #if defined(_MSC_VER) && !defined(__clang__)

172  // MSVC (and not Clang pretending to be MSVC) unconditionally exposes if constexpr (even in C++14 mode),

173  // moreover it triggers warning C4127 (conditional expression is constant) when not using it. nvcc will

174  // be calling cl.exe for host-side codegen.

175  #define HIPCUB_IF_CONSTEXPR constexpr

176  #else

177  #define HIPCUB_IF_CONSTEXPR

178  #endif

179 #endif

180 

181 #ifdef DOXYGEN_SHOULD_SKIP_THIS // Documentation only

182 

187  #define HIPCUB_DEBUG_SYNC

188 

189 #endif // DOXYGEN_SHOULD_SKIP_THIS

190 

191 #if defined(HIPCUB_CUB_API) && defined(HIPCUB_DEBUG_SYNC) && !defined(CUB_DEBUG_SYNC)

192  #define CUB_DEBUG_SYNC

193 #endif

194 

195 #if !defined(HIPCUB_DEBUG_SYNC) \

196  && (defined(CUB_DEBUG_SYNC) || defined(CUB_DEBUG_HOST_ASSERTIONS) \

197  || defined(CUB_DEBUG_DEVICE_ASSERTIONS) || defined(CUB_DEBUG_ALL))

198  #define HIPCUB_DEBUG_SYNC

199 #endif

200 

201 #ifdef HIPCUB_ROCPRIM_API

202  // TODO C++17: use an inline constexpr variable

203  #ifdef HIPCUB_DEBUG_SYNC

204  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE true

205  #else

206  #define HIPCUB_DETAIL_DEBUG_SYNC_VALUE false

207  #endif

208 #endif // HIPCUB_ROCPRIM_API

209 

210 #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:133

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:153