/build/rocrand-sFFtDi/rocrand-10.0/library/include/rocrand/rocrand_common.h Source File

/build/rocrand-sFFtDi/rocrand-10.0/library/include/rocrand/rocrand_common.h Source File#

API library: /build/rocrand-sFFtDi/rocrand-10.0/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
121namespace rocrand_device
122{
123namespace detail
124{
125
126__forceinline__ __device__ __host__
127unsigned 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__
134unsigned 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
141template<typename Engine>
142struct 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
185template<typename T>
186__forceinline__ __device__ __host__
187void split_ull(T& lo, T& hi, unsigned long long int val);
188
189template<>
190__forceinline__ __device__ __host__
191void 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
197template<>
198__forceinline__ __device__ __host__
199void 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_