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