source: sasmodels/sasmodels/kernel_header.c @ b0de252

core_shell_microgelsmagnetic_modelticket-1257-vesicle-productticket_1156ticket_1265_superballticket_822_more_unit_tests
Last change on this file since b0de252 was b0de252, checked in by pkienzle, 5 years ago

improve control over cuda context

  • Property mode set to 100644
File size: 9.4 KB
Line 
1#ifdef __OPENCL_VERSION__
2# define USE_OPENCL
3#elif defined(__CUDACC__)
4# define USE_CUDA
5#elif defined(_OPENMP)
6# define USE_OPENMP
7#elif defined(__CUDACC__)
8# define USE_CUDA
9#endif
10
11// Use SAS_DOUBLE to force the use of double even for float kernels
12#define SAS_DOUBLE dou ## ble
13
14// If opencl is not available, then we are compiling a C function
15// Note: if using a C++ compiler, then define kernel as extern "C"
16#ifdef USE_OPENCL
17
18   #define USE_GPU
19   typedef int int32_t;
20   #define global_par global
21   #define local_par local
22   #define constant_par constant
23   #define global_var global
24   #define local_var local
25   #define constant_var constant
26   #define __device__
27
28   #if defined(USE_SINCOS)
29   #  define SINCOS(angle,svar,cvar) svar=sincos(angle,&cvar)
30   #else
31   #  define SINCOS(angle,svar,cvar) do {const double _t_=angle; svar=sin(_t_);cvar=cos(_t_);} while (0)
32   #endif
33   // Intel CPU on Mac gives strange values for erf(); on the verified
34   // platforms (intel, nvidia, amd), the cephes erf() is significantly
35   // faster than that available in the native OpenCL.
36   #define NEED_ERF
37   // OpenCL only has type generic math
38   #define expf exp
39   #ifndef NEED_ERF
40   #  define erff erf
41   #  define erfcf erfc
42   #endif
43
44#elif defined(USE_CUDA)
45
46   #define USE_GPU
47   #define global_par
48   #define local_par
49   #define constant_par const
50   #define global_var
51   #define local_var __shared__
52   #define constant_var __constant__
53
54   #define kernel extern "C" __global__
55
56   // OpenCL powr(a,b) = C99 pow(a,b), b >= 0
57   // OpenCL pown(a,b) = C99 pow(a,b), b integer
58   #define powr(a,b) pow(a,b)
59   #define pown(a,b) pow(a,b)
60   //typedef int int32_t;
61   #if defined(USE_SINCOS)
62   #  define SINCOS(angle,svar,cvar) sincos(angle,&svar,&cvar)
63   #else
64   #  define SINCOS(angle,svar,cvar) do {const double _t_=angle; svar=sin(_t_);cvar=cos(_t_);} while (0)
65   #endif
66
67#else // !USE_OPENCL && !USE_CUDA
68
69   #define global_par
70   #define local_par
71   #define constant_par const
72   #define global_var
73   #define local_var
74   #define constant_var const
75   #define __device__
76
77   #ifdef __cplusplus
78      #include <cstdio>
79      #include <cmath>
80      using namespace std;
81      #if defined(_MSC_VER)
82         #include <limits>
83         #include <float.h>
84         #define kernel extern "C" __declspec( dllexport )
85         inline double trunc(double x) { return x>=0?floor(x):-floor(-x); }
86         inline double fmin(double x, double y) { return x>y ? y : x; }
87         inline double fmax(double x, double y) { return x<y ? y : x; }
88         #define isnan(x) _isnan(x)
89         #define isinf(x) (!_finite(x))
90         #define isfinite(x) _finite(x)
91         #define NAN (std::numeric_limits<double>::quiet_NaN()) // non-signalling NaN
92         #define INFINITY (std::numeric_limits<double>::infinity())
93         #define NEED_ERF
94         #define NEED_EXPM1
95         #define NEED_TGAMMA
96     #else
97         #define kernel extern "C"
98         #include <cstdint>
99     #endif
100     inline void SINCOS(double angle, double &svar, double &cvar) { svar=sin(angle); cvar=cos(angle); }
101   #else // !__cplusplus
102     #include <inttypes.h>  // C99 guarantees that int32_t types is here
103     #include <stdio.h>
104     #if defined(__TINYC__)
105         typedef int int32_t;
106         #include <math.h>
107         // TODO: check isnan is correct
108         inline double _isnan(double x) { return x != x; } // hope this doesn't optimize away!
109         #undef isnan
110         #define isnan(x) _isnan(x)
111         // Defeat the double->float conversion since we don't have tgmath
112         inline SAS_DOUBLE trunc(SAS_DOUBLE x) { return x>=0?floor(x):-floor(-x); }
113         inline SAS_DOUBLE fmin(SAS_DOUBLE x, SAS_DOUBLE y) { return x>y ? y : x; }
114         inline SAS_DOUBLE fmax(SAS_DOUBLE x, SAS_DOUBLE y) { return x<y ? y : x; }
115         #define NEED_ERF
116         #define NEED_EXPM1
117         #define NEED_TGAMMA
118         // expf missing from windows?
119         #define expf exp
120     #else
121         #include <tgmath.h> // C99 type-generic math, so sin(float) => sinf
122     #endif
123     // MSVC doesn't support C99, so no need for dllexport on C99 branch
124     #define kernel
125     #define SINCOS(angle,svar,cvar) do {const double _t_=angle; svar=sin(_t_);cvar=cos(_t_);} while (0)
126   #endif  // !__cplusplus
127   // OpenCL powr(a,b) = C99 pow(a,b), b >= 0
128   // OpenCL pown(a,b) = C99 pow(a,b), b integer
129   #define powr(a,b) pow(a,b)
130   #define pown(a,b) pow(a,b)
131
132#endif // !USE_OPENCL
133
134#if defined(NEED_EXPM1)
135   // TODO: precision is a half digit lower than numpy on mac in [1e-7, 0.5]
136   // Run "explore/precision.py sas_expm1" to see this (may have to fiddle
137   // the xrange for log to see the complete range).
138   static SAS_DOUBLE expm1(SAS_DOUBLE x_in) {
139      double x = (double)x_in;  // go back to float for single precision kernels
140      // Adapted from the cephes math library.
141      // Copyright 1984 - 1992 by Stephen L. Moshier
142      if (x != x || x == 0.0) {
143         return x; // NaN and +/- 0
144      } else if (x < -0.5 || x > 0.5) {
145         return exp(x) - 1.0;
146      } else {
147         const double xsq = x*x;
148         const double p = (((
149            +1.2617719307481059087798E-4)*xsq
150            +3.0299440770744196129956E-2)*xsq
151            +9.9999999999999999991025E-1);
152         const double q = ((((
153            +3.0019850513866445504159E-6)*xsq
154            +2.5244834034968410419224E-3)*xsq
155            +2.2726554820815502876593E-1)*xsq
156            +2.0000000000000000000897E0);
157         double r = x * p;
158         r =  r / (q - r);
159         return r+r;
160       }
161   }
162#endif
163
164// Standard mathematical constants:
165//   M_E, M_LOG2E, M_LOG10E, M_LN2, M_LN10, M_PI, M_PI_2=pi/2, M_PI_4=pi/4,
166//   M_1_PI=1/pi, M_2_PI=2/pi, M_2_SQRTPI=2/sqrt(pi), SQRT2, SQRT1_2=sqrt(1/2)
167// OpenCL defines M_constant_F for float constants, and nothing if double
168// is not enabled on the card, which is why these constants may be missing
169#ifndef M_PI
170#  define M_PI 3.141592653589793
171#endif
172#ifndef M_PI_2
173#  define M_PI_2 1.570796326794897
174#endif
175#ifndef M_PI_4
176#  define M_PI_4 0.7853981633974483
177#endif
178#ifndef M_E
179#  define M_E 2.718281828459045091
180#endif
181#ifndef M_SQRT1_2
182#  define M_SQRT1_2 0.70710678118654746
183#endif
184
185// Non-standard function library
186// pi/180, used for converting between degrees and radians
187// 4/3 pi for computing sphere volumes
188// square and cube for computing squares and cubes
189#ifndef M_PI_180
190#  define M_PI_180 0.017453292519943295
191#endif
192#ifndef M_4PI_3
193#  define M_4PI_3 4.18879020478639
194#endif
195__device__
196inline double square(double x) { return x*x; }
197__device__
198inline double cube(double x) { return x*x*x; }
199__device__
200inline double sas_sinx_x(double x) { return x==0 ? 1.0 : sin(x)/x; }
201
202// CRUFT: support old style models with orientation received qx, qy and angles
203
204// To rotate from the canonical position to theta, phi, psi, first rotate by
205// psi about the major axis, oriented along z, which is a rotation in the
206// detector plane xy. Next rotate by theta about the y axis, aligning the major
207// axis in the xz plane. Finally, rotate by phi in the detector plane xy.
208// To compute the scattering, undo these rotations in reverse order:
209//     rotate in xy by -phi, rotate in xz by -theta, rotate in xy by -psi
210// The returned q is the length of the q vector and (xhat, yhat, zhat) is a unit
211// vector in the q direction.
212// To change between counterclockwise and clockwise rotation, change the
213// sign of phi and psi.
214
215#if 1
216//think cos(theta) should be sin(theta) in new coords, RKH 11Jan2017
217#define ORIENT_SYMMETRIC(qx, qy, theta, phi, q, sn, cn) do { \
218    SINCOS(phi*M_PI_180, sn, cn); \
219    q = sqrt(qx*qx + qy*qy); \
220    cn  = (q==0. ? 1.0 : (cn*qx + sn*qy)/q * sin(theta*M_PI_180));  \
221    sn = sqrt(1 - cn*cn); \
222    } while (0)
223#else
224// SasView 3.x definition of orientation
225#define ORIENT_SYMMETRIC(qx, qy, theta, phi, q, sn, cn) do { \
226    SINCOS(theta*M_PI_180, sn, cn); \
227    q = sqrt(qx*qx + qy*qy);\
228    cn = (q==0. ? 1.0 : (cn*cos(phi*M_PI_180)*qx + sn*qy)/q); \
229    sn = sqrt(1 - cn*cn); \
230    } while (0)
231#endif
232
233#if 1
234#define ORIENT_ASYMMETRIC(qx, qy, theta, phi, psi, q, xhat, yhat, zhat) do { \
235    q = sqrt(qx*qx + qy*qy); \
236    const double qxhat = qx/q; \
237    const double qyhat = qy/q; \
238    double sin_theta, cos_theta; \
239    double sin_phi, cos_phi; \
240    double sin_psi, cos_psi; \
241    SINCOS(theta*M_PI_180, sin_theta, cos_theta); \
242    SINCOS(phi*M_PI_180, sin_phi, cos_phi); \
243    SINCOS(psi*M_PI_180, sin_psi, cos_psi); \
244    xhat = qxhat*(-sin_phi*sin_psi + cos_theta*cos_phi*cos_psi) \
245         + qyhat*( cos_phi*sin_psi + cos_theta*sin_phi*cos_psi); \
246    yhat = qxhat*(-sin_phi*cos_psi - cos_theta*cos_phi*sin_psi) \
247         + qyhat*( cos_phi*cos_psi - cos_theta*sin_phi*sin_psi); \
248    zhat = qxhat*(-sin_theta*cos_phi) \
249         + qyhat*(-sin_theta*sin_phi); \
250    } while (0)
251#else
252// SasView 3.x definition of orientation
253#define ORIENT_ASYMMETRIC(qx, qy, theta, phi, psi, q, cos_alpha, cos_mu, cos_nu) do { \
254    q = sqrt(qx*qx + qy*qy); \
255    const double qxhat = qx/q; \
256    const double qyhat = qy/q; \
257    double sin_theta, cos_theta; \
258    double sin_phi, cos_phi; \
259    double sin_psi, cos_psi; \
260    SINCOS(theta*M_PI_180, sin_theta, cos_theta); \
261    SINCOS(phi*M_PI_180, sin_phi, cos_phi); \
262    SINCOS(psi*M_PI_180, sin_psi, cos_psi); \
263    cos_alpha = cos_theta*cos_phi*qxhat + sin_theta*qyhat; \
264    cos_mu = (-sin_theta*cos_psi*cos_phi - sin_psi*sin_phi)*qxhat + cos_theta*cos_psi*qyhat; \
265    cos_nu = (-cos_phi*sin_psi*sin_theta + sin_phi*cos_psi)*qxhat + sin_psi*cos_theta*qyhat; \
266    } while (0)
267#endif
Note: See TracBrowser for help on using the repository browser.