/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_common.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_common.h Source File#

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_common.h Source File
rocrand_common.h
1 // Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
2 //
3 // Permission is hereby granted, free of charge, to any person obtaining a copy
4 // of this software and associated documentation files (the "Software"), to deal
5 // in the Software without restriction, including without limitation the rights
6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 // copies of the Software, and to permit persons to whom the Software is
8 // furnished to do so, subject to the following conditions:
9 //
10 // The above copyright notice and this permission notice shall be included in
11 // all copies or substantial portions of the Software.
12 //
13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 // THE SOFTWARE.
20 
21 #ifndef ROCRAND_COMMON_H_
22 #define ROCRAND_COMMON_H_
23 
24 #define ROCRAND_2POW16_INV (1.5258789e-05f)
25 #define ROCRAND_2POW16_INV_2PI (9.58738e-05f)
26 #define ROCRAND_2POW32_INV (2.3283064e-10f)
27 #define ROCRAND_2POW32_INV_DOUBLE (2.3283064365386963e-10)
28 #define ROCRAND_2POW64_INV (5.4210109e-20f)
29 #define ROCRAND_2POW64_INV_DOUBLE (5.4210108624275221700372640043497e-20)
30 #define ROCRAND_2POW32_INV_2PI (1.46291807e-09f)
31 #define ROCRAND_2POW53_INV_DOUBLE (1.1102230246251565e-16)
32 #define ROCRAND_PI (3.141592653f)
33 #define ROCRAND_PI_DOUBLE (3.1415926535897932)
34 #define ROCRAND_2PI (6.2831855f)
35 #define ROCRAND_SQRT2 (1.4142135f)
36 #define ROCRAND_SQRT2_DOUBLE (1.4142135623730951)
37 #define ROCRAND_NAN_FLOAT (0x7fc00000)
38 #define ROCRAND_NAN_DOUBLE (0x7ff8000000000000)
39 
40 #include <hip/hip_runtime.h>
41 #if !defined(__HIPCC_RTC__)
42  #include <utility>
43 #endif // !defined(__HIPCC_RTC__)
44 
45 #define ROCRAND_KERNEL __global__ static
46 
47 #if __HIP_DEVICE_COMPILE__ \
48  && (defined(__HIP_PLATFORM_AMD__) \
49  || (defined(__HIP_PLATFORM_NVCC__) && (__CUDA_ARCH__ >= 530)))
50  #define ROCRAND_HALF_MATH_SUPPORTED
51 #endif
52 
53 // Copyright 2001 John Maddock.
54 // Copyright 2017 Peter Dimov.
55 //
56 // Distributed under the Boost Software License, Version 1.0.
57 //
58 // See http://www.boost.org/LICENSE_1_0.txt
59 //
60 // BOOST_STRINGIZE(X)
61 #define ROCRAND_STRINGIZE(X) ROCRAND_DO_STRINGIZE(X)
62 #define ROCRAND_DO_STRINGIZE(X) #X
63 
64 // Copyright 2017 Peter Dimov.
65 //
66 // Distributed under the Boost Software License, Version 1.0.
67 //
68 // See http://www.boost.org/LICENSE_1_0.txt
69 //
70 // BOOST_PRAGMA_MESSAGE("message")
71 //
72 // Expands to the equivalent of #pragma message("message")
73 #if defined(__INTEL_COMPILER)
74  #define ROCRAND_PRAGMA_MESSAGE(x) \
75  __pragma(message(__FILE__ "(" ROCRAND_STRINGIZE(__LINE__) "): note: " x))
76 #elif defined(__GNUC__)
77  #define ROCRAND_PRAGMA_MESSAGE(x) _Pragma(ROCRAND_STRINGIZE(message(x)))
78 #elif defined(_MSC_VER)
79  #define ROCRAND_PRAGMA_MESSAGE(x) \
80  __pragma(message(__FILE__ "(" ROCRAND_STRINGIZE(__LINE__) "): note: " x))
81 #else
82  #define ROCRAND_PRAGMA_MESSAGE(x)
83 #endif
84 
85 #if __cplusplus >= 201402L
86  #define ROCRAND_DEPRECATED(msg) [[deprecated(msg)]]
87 #elif defined(_MSC_VER) && !defined(__clang__)
88  #define ROCRAND_DEPRECATED(msg) __declspec(deprecated(msg))
89 #elif defined(__clang__) || defined(__GNUC__)
90  #define ROCRAND_DEPRECATED(msg) __attribute__((deprecated(msg)))
91 #else
92  #define ROCRAND_DEPRECATED(msg)
93 #endif
94 
95 // This is an accessor macro for HIP vector types (eg. int2, float4, etc.).
96 // Prior to HIP 7.0, individual elements could be accessed through the
97 // data member:
98 //
99 // int2 vec;
100 // vec.data[0] = 1;
101 //
102 // Beginning with HIP 7.0, the data member is hidden, and individual
103 // elements must be accessed like this:
104 //
105 // int2 vec;
106 // vec[0] = 1;
107 //
108 // You can use the macro like this:
109 //
110 // int2 vec;
111 // ROCRAND_HIPVEC_ACCESS(vec)[0];
112 //
113 #if defined(__HIP_PLATFORM_AMD__)
114  #if HIP_VERSION_MAJOR < 7
115  #define ROCRAND_HIPVEC_ACCESS(x) x.data
116  #else
117  #define ROCRAND_HIPVEC_ACCESS(x) x
118  #endif
119 #endif
120 
121 namespace rocrand_device
122 {
123 namespace detail
124 {
125 
126 __forceinline__ __device__ __host__
127 unsigned long long
128  mad_u64_u32(const unsigned int x, const unsigned int y, const unsigned long long z)
129 {
130  return static_cast<unsigned long long>(x) * static_cast<unsigned long long>(y) + z;
131 }
132 
133 __forceinline__ __device__ __host__
134 unsigned long long mul_u64_u32(const unsigned int x, const unsigned int y)
135 {
136  return static_cast<unsigned long long>(x) * static_cast<unsigned long long>(y);
137 }
138 
139 // This helps access fields of engine's internal state which
140 // saves floats and doubles generated using the Box–Muller transform
141 template<typename Engine>
142 struct engine_boxmuller_helper
143 {
144  static __forceinline__ __device__ __host__
145  bool has_float(const Engine* engine)
146  {
147  return engine->m_state.boxmuller_float != ROCRAND_NAN_FLOAT;
148  }
149 
150  static __forceinline__ __device__ __host__
151  float get_float(Engine* engine)
152  {
153  const float ret = engine->m_state.boxmuller_float;
154  engine->m_state.boxmuller_float = ROCRAND_NAN_FLOAT;
155  return ret;
156  }
157 
158  static __forceinline__ __device__ __host__
159  void save_float(Engine* engine, float f)
160  {
161  engine->m_state.boxmuller_float = f;
162  }
163 
164  static __forceinline__ __device__ __host__
165  bool has_double(const Engine* engine)
166  {
167  return engine->m_state.boxmuller_double != ROCRAND_NAN_DOUBLE;
168  }
169 
170  static __forceinline__ __device__ __host__
171  double get_double(Engine* engine)
172  {
173  const double ret = engine->m_state.boxmuller_double;
174  engine->m_state.boxmuller_double = ROCRAND_NAN_DOUBLE;
175  return ret;
176  }
177 
178  static __forceinline__ __device__ __host__
179  void save_double(Engine* engine, double d)
180  {
181  engine->m_state.boxmuller_double = d;
182  }
183 };
184 
185 template<typename T>
186 __forceinline__ __device__ __host__
187 void split_ull(T& lo, T& hi, unsigned long long int val);
188 
189 template<>
190 __forceinline__ __device__ __host__
191 void split_ull(unsigned int& lo, unsigned int& hi, unsigned long long int val)
192 {
193  lo = val & 0xFFFFFFFF;
194  hi = (val >> 32) & 0xFFFFFFFF;
195 }
196 
197 template<>
198 __forceinline__ __device__ __host__
199 void split_ull(unsigned long long int& lo, unsigned long long int& hi, unsigned long long int val)
200 {
201  lo = val;
202  hi = 0;
203 }
204 
205 } // end namespace detail
206 } // end namespace rocrand_device
207 
208 #endif // ROCRAND_COMMON_H_