wrfv2_fire/phys/wsm5.cu.c

   1 /*
   2
   3    This WSM5 microphysics accelerated for the NVIDIA GPU.  It is experimental and
   4    is not supported as part of WRF.  There is additional information available
   5    at http://www.mmm.ucar.edu/people/michalakes.  Requests for assistance will be
   6    considered only on a case by case basis, favoring active collaborators.
   7
   8    Required: a Linux x86 or x86_64 system with a CUDA-enabled NVIDIA GPU installed
   9    as a co-processor as well as the CUDA libraries on a directory in your system,
  10    for example:
  11
  12       /usr/local/cuda/lib/libcublas.so
  13
  14    included in the CUDA SDK 1.1 from NVIDIA (see nvidia.com).
  15
  16    To use with WRF:
  17
  18    1)  Compile this file and companion file as:
  19
  20          gcc -c wsm5.cu.c
  21          gcc -c wsm5_gpu.cu.c
  22
  23        producing wsm5.cu.o and wsm5_gpu.cu.o
  24
  25    2)  configure WRF, generating a configure.wrf file for your system
  26        Note that serial and dmpar work with the GPU, but smpar
  27        and dm+sm may not.
  28
  29    3)  Modify configure.wrf:
  30
  31       a) add -DTEST_ON_GPU_RK -DRUN_ON_GPU to ARCH_LOCAL
  32       b) add ../phys/wsm5.cu.o and ../phys/wsm5_gpu.cu.o to LIB_LOCAL
  33          (define LIB_LOCAL it does not already exist)
  34       c) add -L/usr/local/cuda/lib -lcuda -lcudart to LIB_LOCAL
  35          (or wherever the cuda lib is on your system)
  36
  37    3)  Compile wrf as usual.
  38
  39    Note: The GPU code is compiled for a maximum number of 41 vertical levels
  40    If you need a larger number, contact below.
  41
  42    20080721, JM  (michalak@ucar.edu)
  43
  44 */
  45
  46 # 1 "/tmp/tmpxft_00001ecc_00000000-0.c"
  47 # 1 "<built-in>"
  48 # 1 "<command line>"
  49 # 1 "/tmp/tmpxft_00001ecc_00000000-0.c"
  50 # 1 "y.cu"
  51 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  52 struct _Z7textureIcLi1EL19cudaTextureReadMode0EE;
  53 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  54 struct _Z7textureIaLi1EL19cudaTextureReadMode0EE;
  55 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  56 struct _Z7textureIhLi1EL19cudaTextureReadMode0EE;
  57 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  58 struct _Z7textureI5char1Li1EL19cudaTextureReadMode0EE;
  59 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  60 struct _Z7textureI6uchar1Li1EL19cudaTextureReadMode0EE;
  61 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  62 struct _Z7textureI5char2Li1EL19cudaTextureReadMode0EE;
  63 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  64 struct _Z7textureI6uchar2Li1EL19cudaTextureReadMode0EE;
  65 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  66 struct _Z7textureI5char3Li1EL19cudaTextureReadMode0EE;
  67 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  68 struct _Z7textureI6uchar3Li1EL19cudaTextureReadMode0EE;
  69 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  70 struct _Z7textureI5char4Li1EL19cudaTextureReadMode0EE;
  71 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  72 struct _Z7textureI6uchar4Li1EL19cudaTextureReadMode0EE;
  73 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  74 struct _Z7textureIsLi1EL19cudaTextureReadMode0EE;
  75 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  76 struct _Z7textureItLi1EL19cudaTextureReadMode0EE;
  77 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  78 struct _Z7textureI6short1Li1EL19cudaTextureReadMode0EE;
  79 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  80 struct _Z7textureI7ushort1Li1EL19cudaTextureReadMode0EE;
  81 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  82 struct _Z7textureI6short2Li1EL19cudaTextureReadMode0EE;
  83 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  84 struct _Z7textureI7ushort2Li1EL19cudaTextureReadMode0EE;
  85 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  86 struct _Z7textureI6short3Li1EL19cudaTextureReadMode0EE;
  87 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  88 struct _Z7textureI7ushort3Li1EL19cudaTextureReadMode0EE;
  89 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  90 struct _Z7textureI6short4Li1EL19cudaTextureReadMode0EE;
  91 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  92 struct _Z7textureI7ushort4Li1EL19cudaTextureReadMode0EE;
  93 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  94 struct _Z7textureIiLi1EL19cudaTextureReadMode0EE;
  95 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  96 struct _Z7textureIjLi1EL19cudaTextureReadMode0EE;
  97 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
  98 struct _Z7textureI4int1Li1EL19cudaTextureReadMode0EE;
  99 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 100 struct _Z7textureI5uint1Li1EL19cudaTextureReadMode0EE;
 101 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 102 struct _Z7textureI4int2Li1EL19cudaTextureReadMode0EE;
 103 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 104 struct _Z7textureI5uint2Li1EL19cudaTextureReadMode0EE;
 105 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 106 struct _Z7textureI4int3Li1EL19cudaTextureReadMode0EE;
 107 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 108 struct _Z7textureI5uint3Li1EL19cudaTextureReadMode0EE;
 109 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 110 struct _Z7textureI4int4Li1EL19cudaTextureReadMode0EE;
 111 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 112 struct _Z7textureI5uint4Li1EL19cudaTextureReadMode0EE;
 113 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 114 struct _Z7textureIcLi1EL19cudaTextureReadMode1EE;
 115 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 116 struct _Z7textureIaLi1EL19cudaTextureReadMode1EE;
 117 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 118 struct _Z7textureIhLi1EL19cudaTextureReadMode1EE;
 119 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 120 struct _Z7textureI5char1Li1EL19cudaTextureReadMode1EE;
 121 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 122 struct _Z7textureI6uchar1Li1EL19cudaTextureReadMode1EE;
 123 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 124 struct _Z7textureI5char2Li1EL19cudaTextureReadMode1EE;
 125 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 126 struct _Z7textureI6uchar2Li1EL19cudaTextureReadMode1EE;
 127 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 128 struct _Z7textureI5char3Li1EL19cudaTextureReadMode1EE;
 129 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 130 struct _Z7textureI6uchar3Li1EL19cudaTextureReadMode1EE;
 131 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 132 struct _Z7textureI5char4Li1EL19cudaTextureReadMode1EE;
 133 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 134 struct _Z7textureI6uchar4Li1EL19cudaTextureReadMode1EE;
 135 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 136 struct _Z7textureIsLi1EL19cudaTextureReadMode1EE;
 137 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 138 struct _Z7textureItLi1EL19cudaTextureReadMode1EE;
 139 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 140 struct _Z7textureI6short1Li1EL19cudaTextureReadMode1EE;
 141 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 142 struct _Z7textureI7ushort1Li1EL19cudaTextureReadMode1EE;
 143 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 144 struct _Z7textureI6short2Li1EL19cudaTextureReadMode1EE;
 145 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 146 struct _Z7textureI7ushort2Li1EL19cudaTextureReadMode1EE;
 147 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 148 struct _Z7textureI6short3Li1EL19cudaTextureReadMode1EE;
 149 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 150 struct _Z7textureI7ushort3Li1EL19cudaTextureReadMode1EE;
 151 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 152 struct _Z7textureI6short4Li1EL19cudaTextureReadMode1EE;
 153 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 154 struct _Z7textureI7ushort4Li1EL19cudaTextureReadMode1EE;
 155 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 156 struct _Z7textureIcLi2EL19cudaTextureReadMode0EE;
 157 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 158 struct _Z7textureIaLi2EL19cudaTextureReadMode0EE;
 159 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 160 struct _Z7textureIhLi2EL19cudaTextureReadMode0EE;
 161 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 162 struct _Z7textureI5char1Li2EL19cudaTextureReadMode0EE;
 163 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 164 struct _Z7textureI6uchar1Li2EL19cudaTextureReadMode0EE;
 165 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 166 struct _Z7textureI5char2Li2EL19cudaTextureReadMode0EE;
 167 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 168 struct _Z7textureI6uchar2Li2EL19cudaTextureReadMode0EE;
 169 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 170 struct _Z7textureI5char3Li2EL19cudaTextureReadMode0EE;
 171 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 172 struct _Z7textureI6uchar3Li2EL19cudaTextureReadMode0EE;
 173 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 174 struct _Z7textureI5char4Li2EL19cudaTextureReadMode0EE;
 175 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 176 struct _Z7textureI6uchar4Li2EL19cudaTextureReadMode0EE;
 177 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 178 struct _Z7textureIsLi2EL19cudaTextureReadMode0EE;
 179 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 180 struct _Z7textureItLi2EL19cudaTextureReadMode0EE;
 181 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 182 struct _Z7textureI6short1Li2EL19cudaTextureReadMode0EE;
 183 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 184 struct _Z7textureI7ushort1Li2EL19cudaTextureReadMode0EE;
 185 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 186 struct _Z7textureI6short2Li2EL19cudaTextureReadMode0EE;
 187 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 188 struct _Z7textureI7ushort2Li2EL19cudaTextureReadMode0EE;
 189 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 190 struct _Z7textureI6short3Li2EL19cudaTextureReadMode0EE;
 191 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 192 struct _Z7textureI7ushort3Li2EL19cudaTextureReadMode0EE;
 193 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 194 struct _Z7textureI6short4Li2EL19cudaTextureReadMode0EE;
 195 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 196 struct _Z7textureI7ushort4Li2EL19cudaTextureReadMode0EE;
 197 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 198 struct _Z7textureIiLi2EL19cudaTextureReadMode0EE;
 199 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 200 struct _Z7textureIjLi2EL19cudaTextureReadMode0EE;
 201 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 202 struct _Z7textureI4int1Li2EL19cudaTextureReadMode0EE;
 203 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 204 struct _Z7textureI5uint1Li2EL19cudaTextureReadMode0EE;
 205 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 206 struct _Z7textureI4int2Li2EL19cudaTextureReadMode0EE;
 207 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 208 struct _Z7textureI5uint2Li2EL19cudaTextureReadMode0EE;
 209 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 210 struct _Z7textureI4int3Li2EL19cudaTextureReadMode0EE;
 211 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 212 struct _Z7textureI5uint3Li2EL19cudaTextureReadMode0EE;
 213 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 214 struct _Z7textureI4int4Li2EL19cudaTextureReadMode0EE;
 215 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 216 struct _Z7textureI5uint4Li2EL19cudaTextureReadMode0EE;
 217 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 218 struct _Z7textureIiLi1EL19cudaTextureReadMode1EE;
 219 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 220 struct _Z7textureIjLi1EL19cudaTextureReadMode1EE;
 221 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 222 struct _Z7textureI4int1Li1EL19cudaTextureReadMode1EE;
 223 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 224 struct _Z7textureI5uint1Li1EL19cudaTextureReadMode1EE;
 225 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 226 struct _Z7textureI4int2Li1EL19cudaTextureReadMode1EE;
 227 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 228 struct _Z7textureI5uint2Li1EL19cudaTextureReadMode1EE;
 229 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 230 struct _Z7textureI4int3Li1EL19cudaTextureReadMode1EE;
 231 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 232 struct _Z7textureI5uint3Li1EL19cudaTextureReadMode1EE;
 233 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 234 struct _Z7textureI4int4Li1EL19cudaTextureReadMode1EE;
 235 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 236 struct _Z7textureI5uint4Li1EL19cudaTextureReadMode1EE;
 237 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 238 struct _Z7textureIcLi2EL19cudaTextureReadMode1EE;
 239 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 240 struct _Z7textureIaLi2EL19cudaTextureReadMode1EE;
 241 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 242 struct _Z7textureIhLi2EL19cudaTextureReadMode1EE;
 243 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 244 struct _Z7textureI5char1Li2EL19cudaTextureReadMode1EE;
 245 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 246 struct _Z7textureI6uchar1Li2EL19cudaTextureReadMode1EE;
 247 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 248 struct _Z7textureI5char2Li2EL19cudaTextureReadMode1EE;
 249 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 250 struct _Z7textureI6uchar2Li2EL19cudaTextureReadMode1EE;
 251 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 252 struct _Z7textureI5char3Li2EL19cudaTextureReadMode1EE;
 253 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 254 struct _Z7textureI6uchar3Li2EL19cudaTextureReadMode1EE;
 255 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 256 struct _Z7textureI5char4Li2EL19cudaTextureReadMode1EE;
 257 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 258 struct _Z7textureI6uchar4Li2EL19cudaTextureReadMode1EE;
 259 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 260 struct _Z7textureIsLi2EL19cudaTextureReadMode1EE;
 261 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 262 struct _Z7textureItLi2EL19cudaTextureReadMode1EE;
 263 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 264 struct _Z7textureI6short1Li2EL19cudaTextureReadMode1EE;
 265 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 266 struct _Z7textureI7ushort1Li2EL19cudaTextureReadMode1EE;
 267 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 268 struct _Z7textureI6short2Li2EL19cudaTextureReadMode1EE;
 269 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 270 struct _Z7textureI7ushort2Li2EL19cudaTextureReadMode1EE;
 271 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 272 struct _Z7textureI6short3Li2EL19cudaTextureReadMode1EE;
 273 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 274 struct _Z7textureI7ushort3Li2EL19cudaTextureReadMode1EE;
 275 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 276 struct _Z7textureI6short4Li2EL19cudaTextureReadMode1EE;
 277 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 278 struct _Z7textureI7ushort4Li2EL19cudaTextureReadMode1EE;
 279 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 280 struct _Z7textureIiLi2EL19cudaTextureReadMode1EE;
 281 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 282 struct _Z7textureIjLi2EL19cudaTextureReadMode1EE;
 283 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 284 struct _Z7textureI4int1Li2EL19cudaTextureReadMode1EE;
 285 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 286 struct _Z7textureI5uint1Li2EL19cudaTextureReadMode1EE;
 287 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 288 struct _Z7textureI4int2Li2EL19cudaTextureReadMode1EE;
 289 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 290 struct _Z7textureI5uint2Li2EL19cudaTextureReadMode1EE;
 291 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 292 struct _Z7textureI4int3Li2EL19cudaTextureReadMode1EE;
 293 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 294 struct _Z7textureI5uint3Li2EL19cudaTextureReadMode1EE;
 295 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 296 struct _Z7textureI4int4Li2EL19cudaTextureReadMode1EE;
 297 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 298 struct _Z7textureI5uint4Li2EL19cudaTextureReadMode1EE;
 299 # 268 "/usr/include/libio.h" 3
 300 struct _IO_FILE;
 301 # 214 "/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/stddef.h" 3
 302 typedef unsigned long size_t;
 303 # 1 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 1 3
 304 # 56 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 3
 305 typedef char bool;
 306
 307
 308
 309 # 1 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 1 3
 310 # 48 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 3
 311 # 1 "/usr/local/cuda/bin/../include/host_defines.h" 1 3
 312 # 49 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 2 3
 313 # 1 "/usr/local/cuda/bin/../include/builtin_types.h" 1 3
 314 # 42 "/usr/local/cuda/bin/../include/builtin_types.h" 3
 315 # 1 "/usr/local/cuda/bin/../include/device_types.h" 1 3
 316 # 46 "/usr/local/cuda/bin/../include/device_types.h" 3
 317 enum cudaRoundMode
 318 {
 319   cudaRoundNearest,
 320   cudaRoundZero,
 321   cudaRoundPosInf,
 322   cudaRoundMinInf
 323 };
 324 # 43 "/usr/local/cuda/bin/../include/builtin_types.h" 2 3
 325 # 1 "/usr/local/cuda/bin/../include/driver_types.h" 1 3
 326 # 60 "/usr/local/cuda/bin/../include/driver_types.h" 3
 327 enum cudaError
 328 {
 329   cudaSuccess = 0,
 330   cudaErrorMissingConfiguration,
 331   cudaErrorMemoryAllocation,
 332   cudaErrorInitializationError,
 333   cudaErrorLaunchFailure,
 334   cudaErrorPriorLaunchFailure,
 335   cudaErrorLaunchTimeout,
 336   cudaErrorLaunchOutOfResources,
 337   cudaErrorInvalidDeviceFunction,
 338   cudaErrorInvalidConfiguration,
 339   cudaErrorInvalidDevice,
 340   cudaErrorInvalidValue,
 341   cudaErrorInvalidPitchValue,
 342   cudaErrorInvalidSymbol,
 343   cudaErrorMapBufferObjectFailed,
 344   cudaErrorUnmapBufferObjectFailed,
 345   cudaErrorInvalidHostPointer,
 346   cudaErrorInvalidDevicePointer,
 347   cudaErrorInvalidTexture,
 348   cudaErrorInvalidTextureBinding,
 349   cudaErrorInvalidChannelDescriptor,
 350   cudaErrorInvalidMemcpyDirection,
 351   cudaErrorAddressOfConstant,
 352   cudaErrorTextureFetchFailed,
 353   cudaErrorTextureNotBound,
 354   cudaErrorSynchronizationError,
 355   cudaErrorInvalidFilterSetting,
 356   cudaErrorInvalidNormSetting,
 357   cudaErrorMixedDeviceExecution,
 358   cudaErrorCudartUnloading,
 359   cudaErrorUnknown,
 360   cudaErrorNotYetImplemented,
 361   cudaErrorMemoryValueTooLarge,
 362   cudaErrorInvalidResourceHandle,
 363   cudaErrorNotReady,
 364   cudaErrorStartupFailure = 0x7f,
 365   cudaErrorApiFailureBase = 10000
 366 };
 367
 368
 369 enum cudaMemcpyKind
 370 {
 371   cudaMemcpyHostToHost = 0,
 372   cudaMemcpyHostToDevice,
 373   cudaMemcpyDeviceToHost,
 374   cudaMemcpyDeviceToDevice
 375 };
 376
 377
 378 struct cudaDeviceProp
 379 {
 380   char name[256];
 381   size_t totalGlobalMem;
 382   size_t sharedMemPerBlock;
 383   int regsPerBlock;
 384   int warpSize;
 385   size_t memPitch;
 386   int maxThreadsPerBlock;
 387   int maxThreadsDim[3];
 388   int maxGridSize[3];
 389   size_t totalConstMem;
 390   int major;
 391   int minor;
 392   int clockRate;
 393   size_t textureAlignment;
 394 };
 395 # 154 "/usr/local/cuda/bin/../include/driver_types.h" 3
 396 typedef enum cudaError cudaError_t;
 397
 398
 399 typedef int cudaStream_t;
 400
 401
 402 typedef int cudaEvent_t;
 403 # 44 "/usr/local/cuda/bin/../include/builtin_types.h" 2 3
 404 # 1 "/usr/local/cuda/bin/../include/texture_types.h" 1 3
 405 # 46 "/usr/local/cuda/bin/../include/texture_types.h" 3
 406 struct cudaArray;
 407
 408
 409 enum cudaChannelFormatKind
 410 {
 411   cudaChannelFormatKindSigned,
 412   cudaChannelFormatKindUnsigned,
 413   cudaChannelFormatKindFloat
 414 };
 415
 416
 417 struct cudaChannelFormatDesc
 418 {
 419   int x;
 420   int y;
 421   int z;
 422   int w;
 423   enum cudaChannelFormatKind f;
 424 };
 425
 426
 427 enum cudaTextureAddressMode
 428 {
 429   cudaAddressModeWrap,
 430   cudaAddressModeClamp
 431 };
 432
 433
 434 enum cudaTextureFilterMode
 435 {
 436   cudaFilterModePoint,
 437   cudaFilterModeLinear
 438 };
 439
 440
 441 enum cudaTextureReadMode
 442 {
 443   cudaReadModeElementType,
 444   cudaReadModeNormalizedFloat
 445 };
 446
 447
 448 struct textureReference
 449 {
 450   int normalized;
 451   enum cudaTextureFilterMode filterMode;
 452   enum cudaTextureAddressMode addressMode[2];
 453   struct cudaChannelFormatDesc channelDesc;
 454 };
 455 # 45 "/usr/local/cuda/bin/../include/builtin_types.h" 2 3
 456 # 1 "/usr/local/cuda/bin/../include/vector_types.h" 1 3
 457 # 54 "/usr/local/cuda/bin/../include/vector_types.h" 3
 458 struct char1
 459 {
 460   signed char x;
 461 };
 462
 463
 464 struct uchar1
 465 {
 466   unsigned char x;
 467 };
 468
 469
 470 struct char2
 471 {
 472   signed char x, y;
 473 };
 474
 475
 476 struct uchar2
 477 {
 478   unsigned char x, y;
 479 };
 480
 481
 482 struct char3
 483 {
 484   signed char x, y, z;
 485 };
 486
 487
 488 struct uchar3
 489 {
 490   unsigned char x, y, z;
 491 };
 492
 493
 494 struct char4
 495 {
 496   signed char x, y, z, w;
 497 };
 498
 499
 500 struct uchar4
 501 {
 502   unsigned char x, y, z, w;
 503 };
 504
 505
 506 struct short1
 507 {
 508   short x;
 509 };
 510
 511
 512 struct ushort1
 513 {
 514   unsigned short x;
 515 };
 516
 517
 518 struct short2
 519 {
 520   short x, y;
 521 };
 522
 523
 524 struct ushort2
 525 {
 526   unsigned short x, y;
 527 };
 528
 529
 530 struct short3
 531 {
 532   short x, y, z;
 533 };
 534
 535
 536 struct ushort3
 537 {
 538   unsigned short x, y, z;
 539 };
 540
 541
 542 struct short4
 543 {
 544   short x, y, z, w;
 545 };
 546
 547
 548 struct ushort4
 549 {
 550   unsigned short x, y, z, w;
 551 };
 552
 553
 554 struct int1
 555 {
 556   int x;
 557 };
 558
 559
 560 struct uint1
 561 {
 562   unsigned int x;
 563 };
 564
 565
 566 struct int2
 567 {
 568   int x, y;
 569 };
 570
 571
 572 struct uint2
 573 {
 574   unsigned int x, y;
 575 };
 576
 577
 578 struct int3
 579 {
 580   int x, y, z;
 581 };
 582
 583
 584 struct uint3
 585 {
 586   unsigned int x, y, z;
 587 };
 588
 589
 590 struct int4
 591 {
 592   int x, y, z, w;
 593 };
 594
 595
 596 struct uint4
 597 {
 598   unsigned int x, y, z, w;
 599 };
 600
 601
 602 struct long1
 603 {
 604   long x;
 605 };
 606
 607
 608 struct ulong1
 609 {
 610   unsigned long x;
 611 };
 612
 613
 614 struct long2
 615 {
 616   long x, y;
 617 };
 618
 619
 620 struct ulong2
 621 {
 622   unsigned long x, y;
 623 };
 624
 625
 626 struct long3
 627 {
 628   long x, y, z;
 629 };
 630
 631
 632 struct ulong3
 633 {
 634   unsigned long x, y, z;
 635 };
 636
 637
 638 struct long4
 639 {
 640   long x, y, z, w;
 641 };
 642
 643
 644 struct ulong4
 645 {
 646   unsigned long x, y, z, w;
 647 };
 648
 649
 650 struct float1
 651 {
 652   float x;
 653 };
 654
 655
 656 struct float2
 657 {
 658   float x, y;
 659 };
 660
 661
 662 struct float3
 663 {
 664   float x, y, z;
 665 };
 666
 667
 668 struct float4
 669 {
 670   float x, y, z, w;
 671 };
 672
 673
 674 struct double2
 675 {
 676   double x, y;
 677 };
 678 # 282 "/usr/local/cuda/bin/../include/vector_types.h" 3
 679 typedef struct char1 char1;
 680
 681 typedef struct uchar1 uchar1;
 682
 683 typedef struct char2 char2;
 684
 685 typedef struct uchar2 uchar2;
 686
 687 typedef struct char3 char3;
 688
 689 typedef struct uchar3 uchar3;
 690
 691 typedef struct char4 char4;
 692
 693 typedef struct uchar4 uchar4;
 694
 695 typedef struct short1 short1;
 696
 697 typedef struct ushort1 ushort1;
 698
 699 typedef struct short2 short2;
 700
 701 typedef struct ushort2 ushort2;
 702
 703 typedef struct short3 short3;
 704
 705 typedef struct ushort3 ushort3;
 706
 707 typedef struct short4 short4;
 708
 709 typedef struct ushort4 ushort4;
 710
 711 typedef struct int1 int1;
 712
 713 typedef struct uint1 uint1;
 714
 715 typedef struct int2 int2;
 716
 717 typedef struct uint2 uint2;
 718
 719 typedef struct int3 int3;
 720
 721 typedef struct uint3 uint3;
 722
 723 typedef struct int4 int4;
 724
 725 typedef struct uint4 uint4;
 726
 727 typedef struct long1 long1;
 728
 729 typedef struct ulong1 ulong1;
 730
 731 typedef struct long2 long2;
 732
 733 typedef struct ulong2 ulong2;
 734
 735 typedef struct long3 long3;
 736
 737 typedef struct ulong3 ulong3;
 738
 739 typedef struct long4 long4;
 740
 741 typedef struct ulong4 ulong4;
 742
 743 typedef struct float1 float1;
 744
 745 typedef struct float2 float2;
 746
 747 typedef struct float3 float3;
 748
 749 typedef struct float4 float4;
 750
 751 typedef struct double2 double2;
 752 # 363 "/usr/local/cuda/bin/../include/vector_types.h" 3
 753 typedef struct dim3 dim3;
 754
 755
 756 struct dim3
 757 {
 758     unsigned int x, y, z;
 759
 760
 761
 762
 763
 764 };
 765 # 45 "/usr/local/cuda/bin/../include/builtin_types.h" 2 3
 766 # 50 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 2 3
 767 # 82 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 3
 768 extern cudaError_t cudaMalloc(void **devPtr, size_t size);
 769 extern cudaError_t cudaMallocHost(void **ptr, size_t size);
 770 extern cudaError_t cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height);
 771 extern cudaError_t cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height );
 772 extern cudaError_t cudaFree(void *devPtr);
 773 extern cudaError_t cudaFreeHost(void *ptr);
 774 extern cudaError_t cudaFreeArray(struct cudaArray *array);
 775 # 97 "/usr/local/cuda/bin/../include/cuda_runtime_api.h" 3
 776 extern cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
 777 extern cudaError_t cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind);
 778 extern cudaError_t cudaMemcpyFromArray(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind);
 779 extern cudaError_t cudaMemcpyArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t count, enum cudaMemcpyKind kind );
 780 extern cudaError_t cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);
 781 extern cudaError_t cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);
 782 extern cudaError_t cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind);
 783 extern cudaError_t cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind );
 784 extern cudaError_t cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset , enum cudaMemcpyKind kind );
 785 extern cudaError_t cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset , enum cudaMemcpyKind kind );
 786
 787
 788
 789
 790
 791
 792
 793 extern cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
 794 extern cudaError_t cudaMemcpyToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
 795 extern cudaError_t cudaMemcpyFromArrayAsync(void *dst, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
 796 extern cudaError_t cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
 797 extern cudaError_t cudaMemcpy2DToArrayAsync(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
 798 extern cudaError_t cudaMemcpy2DFromArrayAsync(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
 799
 800
 801
 802
 803
 804
 805
 806 extern cudaError_t cudaMemset(void *mem, int c, size_t count);
 807 extern cudaError_t cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height);
 808
 809
 810
 811
 812
 813
 814
 815 extern cudaError_t cudaGetSymbolAddress(void **devPtr, const char *symbol);
 816 extern cudaError_t cudaGetSymbolSize(size_t *size, const char *symbol);
 817
 818
 819
 820
 821
 822
 823
 824 extern cudaError_t cudaGetDeviceCount(int *count);
 825 extern cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);
 826 extern cudaError_t cudaChooseDevice(int *device, const struct cudaDeviceProp *prop);
 827 extern cudaError_t cudaSetDevice(int device);
 828 extern cudaError_t cudaGetDevice(int *device);
 829
 830
 831
 832
 833
 834
 835
 836 extern cudaError_t cudaBindTexture(size_t *offset, const struct textureReference *texref, const void *devPtr, const struct cudaChannelFormatDesc *desc, size_t size );
 837 extern cudaError_t cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc);
 838 extern cudaError_t cudaUnbindTexture(const struct textureReference *texref);
 839 extern cudaError_t cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref);
 840 extern cudaError_t cudaGetTextureReference(const struct textureReference **texref, const char *symbol);
 841
 842
 843
 844
 845
 846
 847
 848 extern cudaError_t cudaGetChannelDesc(struct cudaChannelFormatDesc *desc, const struct cudaArray *array);
 849 extern struct cudaChannelFormatDesc cudaCreateChannelDesc(int x, int y, int z, int w, enum cudaChannelFormatKind f);
 850
 851
 852
 853
 854
 855
 856
 857 extern cudaError_t cudaGetLastError(void);
 858 extern const char* cudaGetErrorString(cudaError_t error);
 859
 860
 861
 862
 863
 864
 865
 866 extern cudaError_t cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem , cudaStream_t stream );
 867 extern cudaError_t cudaSetupArgument(const void *arg, size_t size, size_t offset);
 868 extern cudaError_t cudaLaunch(const char *symbol);
 869
 870
 871
 872
 873
 874
 875
 876 extern cudaError_t cudaStreamCreate(cudaStream_t *stream);
 877 extern cudaError_t cudaStreamDestroy(cudaStream_t stream);
 878 extern cudaError_t cudaStreamSynchronize(cudaStream_t stream);
 879 extern cudaError_t cudaStreamQuery(cudaStream_t stream);
 880
 881
 882
 883
 884
 885
 886
 887 extern cudaError_t cudaEventCreate(cudaEvent_t *event);
 888 extern cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream);
 889 extern cudaError_t cudaEventQuery(cudaEvent_t event);
 890 extern cudaError_t cudaEventSynchronize(cudaEvent_t event);
 891 extern cudaError_t cudaEventDestroy(cudaEvent_t event);
 892 extern cudaError_t cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end);
 893
 894
 895
 896
 897
 898
 899
 900 extern cudaError_t cudaThreadExit(void);
 901 extern cudaError_t cudaThreadSynchronize(void);
 902 # 61 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 2 3
 903 # 1 "/usr/local/cuda/bin/../include/crt/storage_class.h" 1 3
 904 # 62 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 2 3
 905 # 216 "/usr/lib/gcc/x86_64-redhat-linux/4.1.2/include/stddef.h" 2 3
 906 # 148 "/usr/include/bits/types.h" 3
 907 typedef long __clock_t;
 908 # 61 "/usr/include/time.h" 3
 909 typedef __clock_t clock_t;
 910 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 911 struct _Z7textureIcLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 912 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 913 struct _Z7textureIaLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 914 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 915 struct _Z7textureIhLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 916 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 917 struct _Z7textureI5char1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 918 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 919 struct _Z7textureI6uchar1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 920 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 921 struct _Z7textureI5char2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 922 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 923 struct _Z7textureI6uchar2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 924 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 925 struct _Z7textureI5char3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 926 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 927 struct _Z7textureI6uchar3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 928 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 929 struct _Z7textureI5char4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 930 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 931 struct _Z7textureI6uchar4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 932 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 933 struct _Z7textureIsLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 934 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 935 struct _Z7textureItLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 936 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 937 struct _Z7textureI6short1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 938 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 939 struct _Z7textureI7ushort1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 940 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 941 struct _Z7textureI6short2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 942 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 943 struct _Z7textureI7ushort2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 944 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 945 struct _Z7textureI6short3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 946 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 947 struct _Z7textureI7ushort3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 948 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 949 struct _Z7textureI6short4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 950 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 951 struct _Z7textureI7ushort4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 952 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 953 struct _Z7textureIiLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 954 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 955 struct _Z7textureIjLi1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 956 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 957 struct _Z7textureI4int1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 958 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 959 struct _Z7textureI5uint1Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 960 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 961 struct _Z7textureI4int2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 962 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 963 struct _Z7textureI5uint2Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 964 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 965 struct _Z7textureI4int3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 966 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 967 struct _Z7textureI5uint3Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 968 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 969 struct _Z7textureI4int4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 970 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 971 struct _Z7textureI5uint4Li1EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
 972 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 973 struct _Z7textureIcLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 974 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 975 struct _Z7textureIaLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 976 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 977 struct _Z7textureIhLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 978 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 979 struct _Z7textureI5char1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 980 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 981 struct _Z7textureI6uchar1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 982 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 983 struct _Z7textureI5char2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 984 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 985 struct _Z7textureI6uchar2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 986 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 987 struct _Z7textureI5char3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 988 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 989 struct _Z7textureI6uchar3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 990 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 991 struct _Z7textureI5char4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 992 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 993 struct _Z7textureI6uchar4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 994 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 995 struct _Z7textureIsLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 996 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 997 struct _Z7textureItLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
 998 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
 999 struct _Z7textureI6short1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1000 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1001 struct _Z7textureI7ushort1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1002 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1003 struct _Z7textureI6short2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1004 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1005 struct _Z7textureI7ushort2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1006 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1007 struct _Z7textureI6short3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1008 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1009 struct _Z7textureI7ushort3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1010 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1011 struct _Z7textureI6short4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1012 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1013 struct _Z7textureI7ushort4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1014 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1015 struct _Z7textureIcLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1016 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1017 struct _Z7textureIaLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1018 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1019 struct _Z7textureIhLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1020 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1021 struct _Z7textureI5char1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1022 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1023 struct _Z7textureI6uchar1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1024 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1025 struct _Z7textureI5char2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1026 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1027 struct _Z7textureI6uchar2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1028 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1029 struct _Z7textureI5char3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1030 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1031 struct _Z7textureI6uchar3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1032 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1033 struct _Z7textureI5char4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1034 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1035 struct _Z7textureI6uchar4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1036 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1037 struct _Z7textureIsLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1038 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1039 struct _Z7textureItLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1040 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1041 struct _Z7textureI6short1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1042 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1043 struct _Z7textureI7ushort1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1044 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1045 struct _Z7textureI6short2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1046 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1047 struct _Z7textureI7ushort2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1048 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1049 struct _Z7textureI6short3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1050 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1051 struct _Z7textureI7ushort3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1052 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1053 struct _Z7textureI6short4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1054 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1055 struct _Z7textureI7ushort4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1056 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1057 struct _Z7textureIiLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1058 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1059 struct _Z7textureIjLi2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1060 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1061 struct _Z7textureI4int1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1062 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1063 struct _Z7textureI5uint1Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1064 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1065 struct _Z7textureI4int2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1066 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1067 struct _Z7textureI5uint2Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1068 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1069 struct _Z7textureI4int3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1070 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1071 struct _Z7textureI5uint3Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1072 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1073 struct _Z7textureI4int4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1074 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1075 struct _Z7textureI5uint4Li2EL19cudaTextureReadMode0EE { struct textureReference __b_16textureReference;};
1076 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1077 struct _Z7textureIiLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1078 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1079 struct _Z7textureIjLi1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1080 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1081 struct _Z7textureI4int1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1082 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1083 struct _Z7textureI5uint1Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1084 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1085 struct _Z7textureI4int2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1086 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1087 struct _Z7textureI5uint2Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1088 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1089 struct _Z7textureI4int3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1090 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1091 struct _Z7textureI5uint3Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1092 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1093 struct _Z7textureI4int4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1094 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1095 struct _Z7textureI5uint4Li1EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1096 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1097 struct _Z7textureIcLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1098 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1099 struct _Z7textureIaLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1100 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1101 struct _Z7textureIhLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1102 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1103 struct _Z7textureI5char1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1104 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1105 struct _Z7textureI6uchar1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1106 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1107 struct _Z7textureI5char2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1108 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1109 struct _Z7textureI6uchar2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1110 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1111 struct _Z7textureI5char3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1112 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1113 struct _Z7textureI6uchar3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1114 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1115 struct _Z7textureI5char4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1116 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1117 struct _Z7textureI6uchar4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1118 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1119 struct _Z7textureIsLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1120 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1121 struct _Z7textureItLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1122 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1123 struct _Z7textureI6short1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1124 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1125 struct _Z7textureI7ushort1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1126 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1127 struct _Z7textureI6short2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1128 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1129 struct _Z7textureI7ushort2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1130 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1131 struct _Z7textureI6short3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1132 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1133 struct _Z7textureI7ushort3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1134 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1135 struct _Z7textureI6short4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1136 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1137 struct _Z7textureI7ushort4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1138 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1139 struct _Z7textureIiLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1140 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1141 struct _Z7textureIjLi2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1142 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1143 struct _Z7textureI4int1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1144 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1145 struct _Z7textureI5uint1Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1146 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1147 struct _Z7textureI4int2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1148 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1149 struct _Z7textureI5uint2Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1150 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1151 struct _Z7textureI4int3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1152 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1153 struct _Z7textureI5uint3Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1154 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1155 struct _Z7textureI4int4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1156 # 115 "/usr/local/cuda/bin/../include/texture_types.h"
1157 struct _Z7textureI5uint4Li2EL19cudaTextureReadMode1EE { struct textureReference __b_16textureReference;};
1158 # 46 "/usr/include/stdio.h" 3
1159 typedef struct _IO_FILE FILE;
1160 void *memcpy(void*, const void*, size_t); void *memset(void*, int, size_t);
1161 # 82 "/usr/local/cuda/bin/../include/cuda_runtime_api.h"
1162 extern cudaError_t cudaMalloc(void **, size_t);
1163
1164
1165
1166 extern cudaError_t cudaFree(void *);
1167 # 97 "/usr/local/cuda/bin/../include/cuda_runtime_api.h"
1168 extern cudaError_t cudaMemcpy(void *, const void *, size_t, enum cudaMemcpyKind);
1169 # 145 "/usr/local/cuda/bin/../include/cuda_runtime_api.h"
1170 extern cudaError_t cudaGetDeviceCount(int *);
1171 extern cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *, int);
1172
1173 extern cudaError_t cudaSetDevice(int);
1174 # 187 "/usr/local/cuda/bin/../include/cuda_runtime_api.h"
1175 extern cudaError_t cudaConfigureCall(dim3, dim3, size_t, cudaStream_t);
1176 # 222 "/usr/local/cuda/bin/../include/cuda_runtime_api.h"
1177 extern cudaError_t cudaThreadSynchronize(void);
1178 # 184 "/usr/include/time.h" 3
1179 extern __attribute__((__weak__)) clock_t clock(void);
1180 # 58 "/usr/local/cuda/bin/../include/common_functions.h"
1181 extern __attribute__((__weak__)) void *memset(void *, int, size_t);
1182 # 65 "/usr/local/cuda/bin/../include/math_functions.h"
1183 extern __attribute__((__weak__)) int abs(int) __attribute__((__const__));
1184
1185 extern __attribute__((__weak__)) long labs(long) __attribute__((__const__));
1186
1187 extern __attribute__((__weak__)) long long llabs(long long) __attribute__((__const__));
1188
1189 extern __attribute__((__weak__)) double fabs(double) __attribute__((__const__));
1190
1191 extern __attribute__((__weak__)) float fabsf(float) __attribute__((__const__));
1192
1193
1194 extern __attribute__((__weak__)) int min(int, int);
1195
1196 extern __attribute__((__weak__)) unsigned umin(unsigned, unsigned);
1197
1198 extern __attribute__((__weak__)) float fminf(float, float);
1199
1200 extern __attribute__((__weak__)) double fmin(double, double);
1201
1202
1203 extern __attribute__((__weak__)) int max(int, int);
1204
1205 extern __attribute__((__weak__)) unsigned umax(unsigned, unsigned);
1206
1207 extern __attribute__((__weak__)) float fmaxf(float, float);
1208
1209 extern __attribute__((__weak__)) double fmax(double, double);
1210
1211
1212 extern __attribute__((__weak__)) double sin(double);
1213
1214 extern __attribute__((__weak__)) float sinf(float);
1215
1216
1217 extern __attribute__((__weak__)) double cos(double);
1218
1219 extern __attribute__((__weak__)) float cosf(float);
1220
1221
1222 extern __attribute__((__weak__)) void sincos(double, double *, double *);
1223
1224 extern __attribute__((__weak__)) void sincosf(float, float *, float *);
1225
1226
1227 extern __attribute__((__weak__)) double tan(double);
1228
1229 extern __attribute__((__weak__)) float tanf(float);
1230
1231
1232 extern __attribute__((__weak__)) double sqrt(double);
1233
1234 extern __attribute__((__weak__)) float sqrtf(float);
1235
1236
1237 extern __attribute__((__weak__)) double rsqrt(double);
1238
1239 extern __attribute__((__weak__)) float rsqrtf(float);
1240
1241
1242 extern __attribute__((__weak__)) double exp2(double);
1243
1244 extern __attribute__((__weak__)) float exp2f(float);
1245
1246
1247 extern __attribute__((__weak__)) double exp10(double);
1248
1249 extern __attribute__((__weak__)) float exp10f(float);
1250
1251
1252 extern __attribute__((__weak__)) double expm1(double);
1253
1254 extern __attribute__((__weak__)) float expm1f(float);
1255
1256
1257 extern __attribute__((__weak__)) double log2(double);
1258
1259 extern __attribute__((__weak__)) float log2f(float);
1260
1261
1262 extern __attribute__((__weak__)) double log10(double);
1263
1264 extern __attribute__((__weak__)) float log10f(float);
1265
1266
1267 extern __attribute__((__weak__)) double log(double);
1268
1269 extern __attribute__((__weak__)) float logf(float);
1270
1271
1272 extern __attribute__((__weak__)) double log1p(double);
1273
1274 extern __attribute__((__weak__)) float log1pf(float);
1275
1276
1277 extern __attribute__((__weak__)) double floor(double) __attribute__((__const__));
1278
1279 extern __attribute__((__weak__)) float floorf(float) __attribute__((__const__));
1280
1281
1282 extern __attribute__((__weak__)) double exp(double);
1283
1284 extern __attribute__((__weak__)) float expf(float);
1285
1286
1287 extern __attribute__((__weak__)) double cosh(double);
1288
1289 extern __attribute__((__weak__)) float coshf(float);
1290
1291
1292 extern __attribute__((__weak__)) double sinh(double);
1293
1294 extern __attribute__((__weak__)) float sinhf(float);
1295
1296
1297 extern __attribute__((__weak__)) double tanh(double);
1298
1299 extern __attribute__((__weak__)) float tanhf(float);
1300
1301
1302 extern __attribute__((__weak__)) double acosh(double);
1303
1304 extern __attribute__((__weak__)) float acoshf(float);
1305
1306
1307 extern __attribute__((__weak__)) double asinh(double);
1308
1309 extern __attribute__((__weak__)) float asinhf(float);
1310
1311
1312 extern __attribute__((__weak__)) double atanh(double);
1313
1314 extern __attribute__((__weak__)) float atanhf(float);
1315
1316
1317 extern __attribute__((__weak__)) double ldexp(double, int);
1318
1319 extern __attribute__((__weak__)) float ldexpf(float, int);
1320
1321
1322 extern __attribute__((__weak__)) double logb(double);
1323
1324 extern __attribute__((__weak__)) float logbf(float);
1325
1326
1327 extern __attribute__((__weak__)) int ilogb(double);
1328
1329 extern __attribute__((__weak__)) int ilogbf(float);
1330
1331
1332 extern __attribute__((__weak__)) double scalbn(double, int);
1333
1334 extern __attribute__((__weak__)) float scalbnf(float, int);
1335
1336
1337 extern __attribute__((__weak__)) double scalbln(double, long);
1338
1339 extern __attribute__((__weak__)) float scalblnf(float, long);
1340
1341
1342 extern __attribute__((__weak__)) double frexp(double, int *);
1343
1344 extern __attribute__((__weak__)) float frexpf(float, int *);
1345
1346
1347 extern __attribute__((__weak__)) double round(double) __attribute__((__const__));
1348
1349 extern __attribute__((__weak__)) float roundf(float) __attribute__((__const__));
1350
1351
1352 extern __attribute__((__weak__)) long lround(double);
1353
1354 extern __attribute__((__weak__)) long lroundf(float);
1355
1356
1357 extern __attribute__((__weak__)) long long llround(double);
1358
1359 extern __attribute__((__weak__)) long long llroundf(float);
1360
1361
1362 extern __attribute__((__weak__)) double rint(double);
1363
1364 extern __attribute__((__weak__)) float rintf(float);
1365
1366
1367 extern __attribute__((__weak__)) long lrint(double);
1368
1369 extern __attribute__((__weak__)) long lrintf(float);
1370
1371
1372 extern __attribute__((__weak__)) long long llrint(double);
1373
1374 extern __attribute__((__weak__)) long long llrintf(float);
1375
1376
1377 extern __attribute__((__weak__)) double nearbyint(double);
1378
1379 extern __attribute__((__weak__)) float nearbyintf(float);
1380
1381
1382 extern __attribute__((__weak__)) double ceil(double) __attribute__((__const__));
1383
1384 extern __attribute__((__weak__)) float ceilf(float) __attribute__((__const__));
1385
1386
1387 extern __attribute__((__weak__)) double trunc(double) __attribute__((__const__));
1388
1389 extern __attribute__((__weak__)) float truncf(float) __attribute__((__const__));
1390
1391
1392 extern __attribute__((__weak__)) double fdim(double, double);
1393
1394 extern __attribute__((__weak__)) float fdimf(float, float);
1395
1396
1397 extern __attribute__((__weak__)) double atan2(double, double);
1398
1399 extern __attribute__((__weak__)) float atan2f(float, float);
1400
1401
1402 extern __attribute__((__weak__)) double atan(double);
1403
1404 extern __attribute__((__weak__)) float atanf(float);
1405
1406
1407 extern __attribute__((__weak__)) double asin(double);
1408
1409 extern __attribute__((__weak__)) float asinf(float);
1410
1411
1412 extern __attribute__((__weak__)) double acos(double);
1413
1414 extern __attribute__((__weak__)) float acosf(float);
1415
1416
1417 extern __attribute__((__weak__)) double hypot(double, double);
1418
1419 extern __attribute__((__weak__)) float hypotf(float, float);
1420
1421
1422 extern __attribute__((__weak__)) double cbrt(double);
1423
1424 extern __attribute__((__weak__)) float cbrtf(float);
1425
1426
1427 extern __attribute__((__weak__)) double pow(double, double);
1428
1429 extern __attribute__((__weak__)) float powf(float, float);
1430
1431
1432 extern __attribute__((__weak__)) double modf(double, double *);
1433
1434 extern __attribute__((__weak__)) float modff(float, float *);
1435
1436
1437 extern __attribute__((__weak__)) double fmod(double, double);
1438
1439 extern __attribute__((__weak__)) float fmodf(float, float);
1440
1441
1442 extern __attribute__((__weak__)) double remainder(double, double);
1443
1444 extern __attribute__((__weak__)) float remainderf(float, float);
1445
1446
1447 extern __attribute__((__weak__)) double remquo(double, double, int *);
1448
1449 extern __attribute__((__weak__)) float remquof(float, float, int *);
1450
1451
1452 extern __attribute__((__weak__)) double erf(double);
1453
1454 extern __attribute__((__weak__)) float erff(float);
1455
1456
1457 extern __attribute__((__weak__)) double erfc(double);
1458
1459 extern __attribute__((__weak__)) float erfcf(float);
1460
1461
1462 extern __attribute__((__weak__)) double lgamma(double);
1463
1464 extern __attribute__((__weak__)) float lgammaf(float);
1465
1466
1467 extern __attribute__((__weak__)) double tgamma(double);
1468
1469 extern __attribute__((__weak__)) float tgammaf(float);
1470
1471
1472 extern __attribute__((__weak__)) double copysign(double, double) __attribute__((__const__));
1473
1474 extern __attribute__((__weak__)) float copysignf(float, float) __attribute__((__const__));
1475
1476
1477 extern __attribute__((__weak__)) double nextafter(double, double) __attribute__((__const__));
1478
1479 extern __attribute__((__weak__)) float nextafterf(float, float) __attribute__((__const__));
1480
1481
1482 extern __attribute__((__weak__)) double nan(const char *) __attribute__((__const__));
1483
1484 extern __attribute__((__weak__)) float nanf(const char *) __attribute__((__const__));
1485
1486
1487 extern __attribute__((__weak__)) int __signbit(double) __attribute__((__const__));
1488
1489 extern __attribute__((__weak__)) int __signbitf(float) __attribute__((__const__));
1490
1491
1492 extern __attribute__((__weak__)) int __isinf(double) __attribute__((__const__));
1493
1494 extern __attribute__((__weak__)) int __isinff(float) __attribute__((__const__));
1495
1496
1497 extern __attribute__((__weak__)) int __isnan(double) __attribute__((__const__));
1498
1499 extern __attribute__((__weak__)) int __isnanf(float) __attribute__((__const__));
1500
1501
1502 extern __attribute__((__weak__)) int __finite(double) __attribute__((__const__));
1503
1504 extern __attribute__((__weak__)) int __finitef(float) __attribute__((__const__));
1505
1506
1507 extern __attribute__((__weak__)) double fma(double, double, double);
1508
1509 extern __attribute__((__weak__)) float fmaf(float, float, float);
1510 # 193 "/usr/include/bits/mathcalls.h" 3
1511 extern __attribute__((__weak__)) int __isinfl(long double) __attribute__((__const__));
1512
1513
1514 extern __attribute__((__weak__)) int __finitel(long double) __attribute__((__const__));
1515 # 231 "/usr/include/bits/mathcalls.h" 3
1516 extern __attribute__((__weak__)) int __isnanl(long double) __attribute__((__const__));
1517 # 350 "/usr/include/bits/mathcalls.h" 3
1518 extern __attribute__((__weak__)) int __signbitl(long double) __attribute__((__const__));
1519 # 589 "/usr/include/stdlib.h" 3
1520 extern void *malloc(size_t) __attribute__((__malloc__));
1521 # 327 "/usr/include/stdio.h" 3
1522 extern int fprintf(FILE *, const char *, ...);
1523 # 113 "y.cu"
1524 extern int rsl_internal_microclock_(void);
1525 # 135 "y.cu"
1526 extern int gethostname(char *, size_t);
1527 # 142 "y.cu"
1528 extern int wsm5_gpu_init_(int *, int *, int *);
1529 # 199 "y.cu"
1530 extern int wsm5_host_(float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *, int *);
1531 # 470 "y.cu"
1532 extern int get_wsm5_gpu_levels_(int *);
1533 extern void __sti___29_tmpxft_00001ecc_00000000_2_ii_91788a12(void) __attribute__((__constructor__));
1534 # 144 "/usr/include/stdio.h" 3
1535 extern struct _IO_FILE *stderr;
1536 # 1 "/tmp/tmpxft_00001ecc_00000000-0.stub.h" 1 3
1537
1538
1539
1540
1541 extern void __device_stub__Z8wsm5_gpuPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_fS_iiiiiiiiiiiiiiiiii(float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float *, float, float *, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int);
1542 # 146 "/usr/include/stdio.h" 2 3
1543 # 142 "y.cu"
1544 int wsm5_gpu_init_( int *myproc, int *nproc, int *mydevice)
1545 {
1546 auto float x;
1547 # 144 "y.cu"
1548 auto float *x_d;
1549 auto int s;
1550 # 145 "y.cu"
1551 auto int e;
1552 auto int i;
1553 # 146 "y.cu"
1554 auto int dc;
1555 auto cudaError_t cerr;
1556 auto char hostname[64];
1557 auto struct cudaDeviceProp dp;
1558
1559 cudaGetDeviceCount((&dc));
1560 if (dc > 4)
1561 { fprintf(stderr, "warning: more than %d devices on node (%d)\n", 4, dc); dc = 4; }
1562 fprintf(stderr, "Number of devices on this node: %d\n", dc);
1563
1564
1565
1566 i = ((*mydevice));
1567 if (dc > 0)
1568 {
1569 if ((int)(cerr = (cudaSetDevice(i)))) {
1570 fprintf(stderr, "    non-zero cerr %d\n", ((int)cerr));
1571 }
1572 }
1573 gethostname(((char *)hostname), 64UL);
1574 fprintf(stderr, "Setting device %02d for task %03d on host %s\n", i, ((*myproc)), ((char *)hostname));
1575
1576 if ((int)(cerr = (cudaGetDeviceProperties((&dp), i)))) {
1577 fprintf(stderr, "Device %02d: cerr = %d\n", ((int)cerr));
1578 } else {
1579 fprintf(stderr, "Device %02d: name %s\n", i, ((char *)(&dp.name)));
1580 fprintf(stderr, "Device %02d: mem       %d\n", i, ((dp.totalGlobalMem)));
1581 fprintf(stderr, "Device %02d: smem      %d\n", i, ((dp.sharedMemPerBlock)));
1582 fprintf(stderr, "Device %02d: nreg      %d\n", i, ((dp.regsPerBlock)));
1583 fprintf(stderr, "Device %02d: warp      %d\n", i, ((dp.warpSize)));
1584 fprintf(stderr, "Device %02d: pitch     %d\n", i, ((dp.memPitch)));
1585 fprintf(stderr, "Device %02d: maxthrds  %d\n", i, ((dp.maxThreadsPerBlock)));
1586 fprintf(stderr, "Device %02d: maxtdim   %d %d %d\n", i, (((int *)(&dp.maxThreadsDim))[0]), (((int *)(&dp.maxThreadsDim))[1]), (((int *)(&dp.maxThreadsDim))[2]));
1587
1588
1589 fprintf(stderr, "Device %02d: maxgdim   %d %d %d\n", i, (((int *)(&dp.maxGridSize))[0]), (((int *)(&dp.maxGridSize))[1]), (((int *)(&dp.maxGridSize))[2]));
1590
1591
1592 fprintf(stderr, "Device %02d: clock     %d\n", i, ((dp.clockRate)));
1593 fprintf(stderr, "Device %02d: talign    %d\n", i, ((dp.textureAlignment)));
1594 }
1595
1596
1597 s = (rsl_internal_microclock_());
1598 cudaMalloc(((void **)(&x_d)), 4UL);
1599 cudaMemcpy(((void *)x_d), ((const void *)(&x)), 4UL, cudaMemcpyHostToDevice);
1600 cudaFree(((void *)x_d));
1601 e = (rsl_internal_microclock_());
1602 fprintf(stderr, "wsm5_init: %d\n", (e - s));
1603 return 0;
1604 }
1605
1606
1607 int wsm5_host_(
1608 float *th, float *pii,
1609 float *q,
1610 float *qc, float *qi, float *qr, float *qs,
1611 float *den, float *p, float *delz,
1612
1613
1614
1615 float *delt,
1616 float *rain, float *rainncv,
1617 float *sr,
1618 float *snow, float *snowncv,
1619 int *ids, int *ide, int *jds, int *jde, int *kds, int *kde,
1620 int *ims, int *ime, int *jms, int *jme, int *kms, int *kme,
1621 int *ips, int *ipe, int *jps, int *jpe, int *kps, int *kpe)
1622
1623 { auto unsigned __T20;
1624 auto unsigned __T21;
1625 auto float *bigbuf;
1626 auto int s;
1627 # 218 "y.cu"
1628 auto int e;
1629 # 218 "y.cu"
1630 auto int s2;
1631 # 218 "y.cu"
1632 auto int e2;
1633 auto int d3;
1634 auto int d2;
1635 # 229 "y.cu"
1636 auto int dips;
1637 # 229 "y.cu"
1638 auto int dipe;
1639
1640 auto int djps;
1641 # 231 "y.cu"
1642 auto int djpe;
1643 auto int dkps;
1644 # 232 "y.cu"
1645 auto int dkpe;
1646 # 242 "y.cu"
1647 auto float *th_d;
1648 auto float *pii_d;
1649 auto float *q_d;
1650 auto float *qc_d;
1651 auto float *qi_d;
1652 auto float *qr_d;
1653 auto float *qs_d;
1654 auto float *den_d;
1655 auto float *p_d;
1656 auto float *delz_d;
1657
1658
1659
1660 auto float *rain_d;
1661 auto float *rainncv_d;
1662 auto float *sr_d;
1663 auto float *snow_d;
1664 auto float *snowncv_d;
1665 auto float retvals[100];
1666
1667
1668
1669 auto float *retvals_d;
1670
1671 auto int remx;
1672 # 266 "y.cu"
1673 auto int remy;
1674
1675
1676
1677
1678 auto dim3 dimBlock;
1679
1680 auto dim3 dimGrid;
1681 # 219 "y.cu"
1682 d3 = ((((((*ime)) - ((*ims))) + 1) * ((((*jme)) - ((*jms))) + 1)) * ((((*kme)) - ((*kms))) + 1));
1683 d2 = (((((*ime)) - ((*ims))) + 1) * ((((*jme)) - ((*jms))) + 1));
1684 # 229 "y.cu"
1685 dips = 0; dipe = ((((*ipe)) - ((*ips))) + 1);
1686
1687 djps = 0; djpe = ((((*jpe)) - ((*jps))) + 1);
1688 dkps = 0; dkpe = ((((*kpe)) - ((*kps))) + 1);
1689
1690 bigbuf = ((float *)(malloc((((unsigned long)((dipe * djpe) * dkpe)) * 4UL))));
1691 # 241 "y.cu"
1692 s = (rsl_internal_microclock_());
1693 cudaMalloc(((void **)(&th_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)th_d), ((const void *)th), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1694 cudaMalloc(((void **)(&pii_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)pii_d), ((const void *)pii), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1695 cudaMalloc(((void **)(&q_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)q_d), ((const void *)q), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1696 cudaMalloc(((void **)(&qc_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)qc_d), ((const void *)qc), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1697 cudaMalloc(((void **)(&qi_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)qi_d), ((const void *)qi), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1698 cudaMalloc(((void **)(&qr_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)qr_d), ((const void *)qr), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1699 cudaMalloc(((void **)(&qs_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)qs_d), ((const void *)qs), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1700 cudaMalloc(((void **)(&den_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)den_d), ((const void *)den), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1701 cudaMalloc(((void **)(&p_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)p_d), ((const void *)p), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1702 cudaMalloc(((void **)(&delz_d)), (((unsigned long)d3) * 4UL)); cudaMemcpy(((void *)delz_d), ((const void *)delz), (((unsigned long)d3) * 4UL), cudaMemcpyHostToDevice);
1703
1704
1705
1706 cudaMalloc(((void **)(&rain_d)), (((unsigned long)d2) * 4UL)); cudaMemcpy(((void *)rain_d), ((const void *)rain), (((unsigned long)d2) * 4UL), cudaMemcpyHostToDevice);
1707 cudaMalloc(((void **)(&rainncv_d)), (((unsigned long)d2) * 4UL)); cudaMemcpy(((void *)rainncv_d), ((const void *)rainncv), (((unsigned long)d2) * 4UL), cudaMemcpyHostToDevice);
1708 cudaMalloc(((void **)(&sr_d)), (((unsigned long)d2) * 4UL)); cudaMemcpy(((void *)sr_d), ((const void *)sr), (((unsigned long)d2) * 4UL), cudaMemcpyHostToDevice);
1709 cudaMalloc(((void **)(&snow_d)), (((unsigned long)d2) * 4UL)); cudaMemcpy(((void *)snow_d), ((const void *)snow), (((unsigned long)d2) * 4UL), cudaMemcpyHostToDevice);
1710 cudaMalloc(((void **)(&snowncv_d)), (((unsigned long)d2) * 4UL)); cudaMemcpy(((void *)snowncv_d), ((const void *)snowncv), (((unsigned long)d2) * 4UL), cudaMemcpyHostToDevice);
1711
1712 { auto int k;
1713 for (k = 0; (k < ((((*kme)) - ((*kms))) + 1)); k++) { (((float *)retvals)[k]) = (0.0F); }
1714 }
1715 cudaMalloc(((void **)(&retvals_d)), (((unsigned long)((((*kme)) - ((*kms))) + 1)) * 4UL)); cudaMemcpy(((void *)retvals_d), ((const void *)((float *)retvals)), (((unsigned long)((((*kme)) - ((*kms))) + 1)) * 4UL), cudaMemcpyHostToDevice);
1716
1717
1718
1719 remx = ((((((*ipe)) - ((*ips))) + 1) % 16) ? 1 : 0);
1720 remy = ((((((*jpe)) - ((*jps))) + 1) % 8) ? 1 : 0);
1721
1722 { (dimBlock.x) = 16U; (dimBlock.y) = 8U; (dimBlock.z) = 1U; }
1723
1724 { __T20 = ((unsigned)((((((*ipe)) - ((*ips))) + 1) / 16) + remx)); __T21 = ((unsigned)((((((*jpe)) - ((*jps))) + 1) / 8) + remy)); { (dimGrid.x) = __T20; (dimGrid.y) = __T21; (dimGrid.z) = 1U; } }
1725
1726 fprintf(stderr, "Call to wsm5_gpu: block dims %d %d\n", ((dimBlock.x)), ((dimBlock.y)));
1727 fprintf(stderr, "Call to wsm5_gpu: grid  dims %d %d\n", ((dimGrid.x)), ((dimGrid.y)));
1728 # 284 "y.cu"
1729 s2 = (rsl_internal_microclock_());
1730 ((int)(cudaConfigureCall(dimGrid, dimBlock, 0UL, 0))) ? ((void)0) : (__device_stub__Z8wsm5_gpuPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_fS_iiiiiiiiiiiiiiiiii(th_d, pii_d, q_d, qc_d, qi_d, qr_d, qs_d, den_d, p_d, delz_d, rain_d, rainncv_d, sr_d, snow_d, snowncv_d, ((*delt)), retvals_d, (dips + 1), ((((*ipe)) - ((*ips))) + 1), (djps + 1), ((((*jpe)) - ((*jps))) + 1), (dkps + 1), ((((*kpe)) - ((*kps))) + 1), (dips + 1), dipe, (djps + 1), djpe, (dkps + 1), dkpe, (dips + 1), dipe, (djps + 1), djpe, (dkps + 1), dkpe));
1731 # 299 "y.cu"
1732 cudaThreadSynchronize();
1733 e2 = (rsl_internal_microclock_());
1734 fprintf(stderr, "Call to wsm5_gpu (not including data xfer): %d microseconds\n", (e2 - s2));
1735
1736
1737 cudaMemcpy(((void *)th), ((const void *)th_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1738 cudaMemcpy(((void *)pii), ((const void *)pii_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1739 cudaMemcpy(((void *)q), ((const void *)q_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1740 cudaMemcpy(((void *)qc), ((const void *)qc_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1741 cudaMemcpy(((void *)qi), ((const void *)qi_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1742 cudaMemcpy(((void *)qr), ((const void *)qr_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1743 cudaMemcpy(((void *)qs), ((const void *)qs_d), (((unsigned long)d3) * 4UL), cudaMemcpyDeviceToHost);
1744
1745
1746
1747 cudaMemcpy(((void *)rain), ((const void *)rain_d), (((unsigned long)d2) * 4UL), cudaMemcpyDeviceToHost);
1748 cudaMemcpy(((void *)rainncv), ((const void *)rainncv_d), (((unsigned long)d2) * 4UL), cudaMemcpyDeviceToHost);
1749 cudaMemcpy(((void *)sr), ((const void *)sr_d), (((unsigned long)d2) * 4UL), cudaMemcpyDeviceToHost);
1750 cudaMemcpy(((void *)snow), ((const void *)snow_d), (((unsigned long)d2) * 4UL), cudaMemcpyDeviceToHost);
1751 cudaMemcpy(((void *)snowncv), ((const void *)snowncv_d), (((unsigned long)d2) * 4UL), cudaMemcpyDeviceToHost);
1752 e = (rsl_internal_microclock_());
1753
1754 cudaMemcpy(((void *)((float *)retvals)), ((const void *)retvals_d), (((unsigned long)((((*kme)) - ((*kms))) + 1)) * 4UL), cudaMemcpyDeviceToHost);
1755 fprintf(stderr, "Call to wsm5_gpu (including data xfer): %d microseconds\n", (e - s));
1756
1757 {
1758
1759
1760 }
1761
1762 cudaFree(((void *)th_d));
1763 cudaFree(((void *)pii_d));
1764 cudaFree(((void *)q_d));
1765 cudaFree(((void *)qc_d));
1766 cudaFree(((void *)qi_d));
1767 cudaFree(((void *)qr_d));
1768 cudaFree(((void *)qs_d));
1769 cudaFree(((void *)den_d));
1770 cudaFree(((void *)p_d));
1771 cudaFree(((void *)delz_d));
1772
1773
1774
1775 cudaFree(((void *)rain_d));
1776 cudaFree(((void *)rainncv_d));
1777 cudaFree(((void *)sr_d));
1778 cudaFree(((void *)snow_d));
1779 cudaFree(((void *)snowncv_d));
1780 cudaFree(((void *)retvals_d));
1781
1782 return 0;
1783 }
1784 # 470 "y.cu"
1785 int get_wsm5_gpu_levels_( int *retval)
1786 {
1787 (*retval) = 41;
1788 }
1789 void __sti___29_tmpxft_00001ecc_00000000_2_ii_91788a12(void) { }
1790 # 1 "/tmp/tmpxft_00001ecc_00000000-0.stub.c" 1
1791
1792
1793
1794 # 1 "/tmp/tmpxft_00001ecc_00000000-1.c" 1
1795 # 1 "/usr/local/cuda/bin/../include/__cudaFatFormat.h" 1
1796 # 97 "/usr/local/cuda/bin/../include/__cudaFatFormat.h"
1797 typedef struct {
1798     char* gpuProfileName;
1799     char* cubin;
1800 } __cudaFatCubinEntry;
1801 # 113 "/usr/local/cuda/bin/../include/__cudaFatFormat.h"
1802 typedef struct {
1803     char* gpuProfileName;
1804     char* ptx;
1805 } __cudaFatPtxEntry;
1806 # 125 "/usr/local/cuda/bin/../include/__cudaFatFormat.h"
1807 typedef struct {
1808     char* gpuProfileName;
1809     char* debug;
1810 } __cudaFatDebugEntry;
1811
1812
1813 typedef enum {
1814       __cudaFatDontSearchFlag = (1 << 0),
1815       __cudaFatDontCacheFlag = (1 << 1)
1816 } __cudaFatCudaBinaryFlag;
1817 # 145 "/usr/local/cuda/bin/../include/__cudaFatFormat.h"
1818 typedef struct {
1819     unsigned long magic;
1820     unsigned long version;
1821     unsigned long gpuInfoVersion;
1822     char* key;
1823     char* ident;
1824     char* usageMode;
1825     __cudaFatPtxEntry *ptx;
1826     __cudaFatCubinEntry *cubin;
1827     __cudaFatDebugEntry *debug;
1828     void* debugInfo;
1829     unsigned int flags;
1830 } __cudaFatCudaBinary;
1831 # 189 "/usr/local/cuda/bin/../include/__cudaFatFormat.h"
1832 void fatGetCubinForGpu( __cudaFatCudaBinary *binary, char* gpuName, char* *cubin, char* *dbgInfoFile );
1833 # 2 "/tmp/tmpxft_00001ecc_00000000-1.c" 2
1834
1835
1836
1837
1838
1839
1840
1841 static const unsigned char __deviceText[] = {
1842 0x61,0x72,0x63,0x68,0x69,0x74,0x65,0x63,0x74,0x75,0x72,0x65,0x20,0x7b,0x73,0x6d,
1843 0x5f,0x31,0x30,0x7d,0x0a,0x61,0x62,0x69,0x76,0x65,0x72,0x73,0x69,0x6f,0x6e,0x20,
1844 0x7b,0x30,0x7d,0x0a,0x63,0x6f,0x64,0x65,0x20,0x20,0x7b,0x0a,0x09,0x6e,0x61,0x6d,
1845 0x65,0x20,0x3d,0x20,0x5f,0x5f,0x64,0x75,0x6d,0x6d,0x79,0x5f,0x65,0x6e,0x74,0x72,
1846 0x79,0x5f,0x5f,0x0a,0x09,0x6c,0x6d,0x65,0x6d,0x20,0x3d,0x20,0x30,0x0a,0x09,0x73,
1847 0x6d,0x65,0x6d,0x20,0x3d,0x20,0x30,0x0a,0x09,0x72,0x65,0x67,0x20,0x3d,0x20,0x30,
1848 0x0a,0x09,0x62,0x61,0x72,0x20,0x3d,0x20,0x30,0x0a,0x09,0x62,0x69,0x6e,0x63,0x6f,
1849 0x64,0x65,0x20,0x20,0x7b,0x0a,0x09,0x09,0x30,0x78,0x66,0x30,0x30,0x30,0x30,0x30,
1850 0x30,0x31,0x20,0x30,0x78,0x65,0x30,0x30,0x30,0x30,0x30,0x30,0x31,0x20,0x0a,0x09,
1851 0x7d,0x0a,0x7d,0x0a,0x00
1852 };
1853
1854
1855
1856
1857
1858 static __cudaFatPtxEntry __ptxEntries [] = {{0,0}};
1859 static __cudaFatCubinEntry __cubinEntries[] = {{"sm_10",(char*)__deviceText},{0,0}};
1860 static __cudaFatDebugEntry __debugEntries[] = {{0,0}};
1861
1862 static __cudaFatCudaBinary __fatDeviceText __attribute__ ((section (".nvFatBinSegment")))= {0x1ee55a01,0x00000002,0x840b5bca,"81bb892378501d16","y.cu"," ",__ptxEntries,__cubinEntries,__debugEntries,0,0};
1863 # 5 "/tmp/tmpxft_00001ecc_00000000-0.stub.c" 2
1864 # 1 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 1
1865 # 65 "/usr/local/cuda/bin/../include/crt/host_runtime.h"
1866 # 1 "/usr/local/cuda/bin/../include/host_defines.h" 1
1867 # 66 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 2
1868 # 88 "/usr/local/cuda/bin/../include/crt/host_runtime.h"
1869 extern void** __cudaRegisterFatBinary(
1870   void *fatCubin
1871 );
1872
1873 extern void __cudaUnregisterFatBinary(
1874   void **fatCubinHandle
1875 );
1876
1877 extern void __cudaRegisterVar(
1878         void **fatCubinHandle,
1879         char *hostVar,
1880         char *deviceAddress,
1881   const char *deviceName,
1882         int ext,
1883         int size,
1884         int constant,
1885         int global
1886 );
1887
1888 extern void __cudaRegisterTexture(
1889         void **fatCubinHandle,
1890   const struct textureReference *hostVar,
1891   const void **deviceAddress,
1892   const char *deviceName,
1893         int dim,
1894         int norm,
1895         int ext
1896 );
1897
1898 extern void __cudaRegisterShared(
1899   void **fatCubinHandle,
1900   void **devicePtr
1901 );
1902
1903 extern void __cudaRegisterFunction(
1904         void **fatCubinHandle,
1905   const char *hostFun,
1906         char *deviceFun,
1907   const char *deviceName,
1908         int thread_limit,
1909         uint3 *tid,
1910         uint3 *bid,
1911         dim3 *bDim,
1912         dim3 *gDim
1913 );
1914
1915
1916
1917
1918
1919 static void **__cudaFatCubinHandle;
1920
1921 static void __cudaUnregisterBinaryUtil(void)
1922 {
1923   __cudaUnregisterFatBinary(__cudaFatCubinHandle);
1924 }
1925
1926
1927
1928 __attribute__((destructor)) static void __cudaUnregisterBinary(void)
1929 {
1930   __cudaUnregisterBinaryUtil();
1931 }
1932 # 195 "/usr/local/cuda/bin/../include/crt/host_runtime.h"
1933 # 1 "/usr/local/cuda/bin/../include/common_functions.h" 1
1934 # 64 "/usr/local/cuda/bin/../include/common_functions.h"
1935 # 1 "/usr/local/cuda/bin/../include/crt/func_macro.h" 1 3
1936 # 65 "/usr/local/cuda/bin/../include/common_functions.h" 2
1937
1938 extern __attribute__((weak)) long __cuda_clock(void); long __cuda_clock(void)
1939 {
1940   return clock();
1941 }
1942
1943 extern __attribute__((weak)) void *__cuda_memset(void *s, int c, size_t n); void *__cuda_memset(void *s, int c, size_t n)
1944 {
1945   char *p = (char*)s;
1946
1947   while (n--) *p++ = (char)c;
1948
1949   return s;
1950 }
1951 # 88 "/usr/local/cuda/bin/../include/common_functions.h"
1952 # 1 "/usr/local/cuda/bin/../include/math_functions.h" 1 3
1953 # 794 "/usr/local/cuda/bin/../include/math_functions.h" 3
1954 extern __attribute__((weak)) int __cuda___signbitl(long double a); int __cuda___signbitl(long double a);
1955 extern __attribute__((weak)) int __cuda___isinfl(long double a); int __cuda___isinfl(long double a);
1956 extern __attribute__((weak)) int __cuda___isnanl(long double a); int __cuda___isnanl(long double a);
1957 extern __attribute__((weak)) int __cuda___finitel(long double a); int __cuda___finitel(long double a);
1958 # 834 "/usr/local/cuda/bin/../include/math_functions.h" 3
1959 extern __attribute__((weak)) int __cuda_abs(int a); int __cuda_abs(int a)
1960 {
1961   return abs(a);
1962 }
1963
1964 extern __attribute__((weak)) float __cuda_fabsf(float a); float __cuda_fabsf(float a)
1965 {
1966   return fabsf(a);
1967 }
1968
1969 extern __attribute__((weak)) long long int __cuda_llabs(long long int a); long long int __cuda_llabs(long long int a)
1970 {
1971
1972
1973
1974   return llabs(a);
1975
1976 }
1977
1978 extern __attribute__((weak)) float __cuda_exp2f(float a); float __cuda_exp2f(float a)
1979 {
1980   return exp2f(a);
1981 }
1982
1983 # 1 "/usr/local/cuda/bin/../include/device_functions.h" 1 3
1984 # 322 "/usr/local/cuda/bin/../include/device_functions.h" 3
1985 # 1 "/usr/local/cuda/bin/../include/math_constants.h" 1 3
1986 # 323 "/usr/local/cuda/bin/../include/device_functions.h" 2 3
1987
1988
1989
1990        extern __attribute__((weak)) int __cuda___isnan(double a); int __cuda___isnan(double a);
1991        extern __attribute__((weak)) int __cuda___isnanf(float a); int __cuda___isnanf(float a);
1992 static int __double2int_rz(double);
1993 static unsigned int __double2uint_rz(double);
1994 static long long int __double2ll_rz(double);
1995 static unsigned long long int __double2ull_rz(double);
1996 # 345 "/usr/local/cuda/bin/../include/device_functions.h" 3
1997 static int __mulhi(int a, int b)
1998 {
1999   long long int c = (long long int)a * (long long int)b;
2000
2001   return (int)(c >> 32);
2002 }
2003
2004 static unsigned int __umulhi(unsigned int a, unsigned int b)
2005 {
2006   unsigned long long int c = (unsigned long long int)a * (unsigned long long int)b;
2007
2008   return (unsigned int)(c >> 32);
2009 }
2010
2011 static unsigned long long int __umul64hi(unsigned long long int a, unsigned long long int b)
2012 {
2013   unsigned int a_lo = (unsigned int)a;
2014   unsigned long long int a_hi = a >> 32;
2015   unsigned int b_lo = (unsigned int)b;
2016   unsigned long long int b_hi = b >> 32;
2017   unsigned long long int m1 = a_lo * b_hi;
2018   unsigned long long int m2 = a_hi * b_lo;
2019   unsigned int carry;
2020
2021   carry = (0ULL + __umulhi(a_lo, b_lo) + (unsigned int)m1 + (unsigned int)m2) >> 32;
2022
2023   return a_hi * b_hi + (m1 >> 32) + (m2 >> 32) + carry;
2024 }
2025
2026 static long long int __mul64hi(long long int a, long long int b)
2027 {
2028   return __umul64hi(a, b) - (a < 0LL ? b : 0LL) - (b < 0LL ? a : 0LL);
2029 }
2030
2031 static float __saturatef(float a)
2032 {
2033   return a >= 1.0f ? 1.0f : a <= 0.0f ? 0.0f : a;
2034 }
2035
2036 static unsigned int __sad(int a, int b, unsigned int c)
2037 {
2038   long long int diff = (long long int)a - (long long int)b;
2039
2040   return (unsigned int)(__cuda_llabs(diff) + (long long int)c);
2041 }
2042
2043 static unsigned int __usad(unsigned int a, unsigned int b, unsigned int c)
2044 {
2045   long long int diff = (long long int)a - (long long int)b;
2046
2047   return (unsigned int)(__cuda_llabs(diff) + (long long int)c);
2048 }
2049
2050 static int __mul24(int a, int b)
2051 {
2052   a &= 0xffffff;
2053   a = (a & 0x800000) != 0 ? a | ~0xffffff : a;
2054   b &= 0xffffff;
2055   b = (b & 0x800000) != 0 ? b | ~0xffffff : b;
2056
2057   return a * b;
2058 }
2059
2060 static unsigned int __umul24(unsigned int a, unsigned int b)
2061 {
2062   a &= 0xffffff;
2063   b &= 0xffffff;
2064
2065   return a * b;
2066 }
2067
2068 static float __int_as_float(int a)
2069 {
2070   union {int a; float b;} u;
2071
2072   u.a = a;
2073
2074   return u.b;
2075 }
2076
2077 static int __float_as_int(float a)
2078 {
2079   union {float a; int b;} u;
2080
2081   u.a = a;
2082
2083   return u.b;
2084 }
2085
2086 static long long int __internal_float2ll_kernel(float a, long long int max, long long int min, long long int nan, enum cudaRoundMode rndMode)
2087 {
2088   unsigned long long int res, t = 0ULL;
2089   int shift;
2090   unsigned int ia;
2091
2092   if (sizeof(a) == sizeof(double) && __cuda___isnan((double)a)) return nan; if (sizeof(a) == sizeof(float) && __cuda___isnanf((float)a)) return nan; if (a >= max) return max; if (a <= min) return min;
2093   ia = __float_as_int(a);
2094   shift = 189 - ((ia >> 23) & 0xff);
2095   res = (unsigned long long int)(((ia << 8) | 0x80000000) >> 1) << 32;
2096   if (shift >= 64) {
2097     t = res;
2098     res = 0;
2099   } else if (shift) {
2100     t = res << (64 - shift);
2101     res = res >> shift;
2102   }
2103   if (rndMode == cudaRoundNearest && (long long int)t < 0LL) {
2104     res += t == 0x8000000000000000ULL ? res & 1ULL : 1ULL;
2105   }
2106   else if (rndMode == cudaRoundMinInf && t != 0ULL && ia > 0x80000000) {
2107     res++;
2108   }
2109   else if (rndMode == cudaRoundPosInf && t != 0ULL && (int)ia > 0) {
2110     res++;
2111   }
2112   if ((int)ia < 0) res = (unsigned long long int)-(long long int)res;
2113   return (long long int)res;
2114 }
2115
2116 static int __internal_float2int(float a, enum cudaRoundMode rndMode)
2117 {
2118   return (int)__internal_float2ll_kernel(a, 2147483647LL, -2147483648LL, 0LL, rndMode);
2119 }
2120
2121 static int __float2int_rz(float a)
2122 {
2123   return __internal_float2int(a, cudaRoundZero);
2124 }
2125
2126 static int __float2int_ru(float a)
2127 {
2128   return __internal_float2int(a, cudaRoundPosInf);
2129 }
2130
2131 static int __float2int_rd(float a)
2132 {
2133   return __internal_float2int(a, cudaRoundMinInf);
2134 }
2135
2136 static int __float2int_rn(float a)
2137 {
2138   return __internal_float2int(a, cudaRoundNearest);
2139 }
2140
2141 static long long int __internal_float2ll(float a, enum cudaRoundMode rndMode)
2142 {
2143   return __internal_float2ll_kernel(a, 9223372036854775807LL, -9223372036854775807LL -1LL, -9223372036854775807LL -1LL, rndMode);
2144 }
2145
2146 static long long int __float2ll_rz(float a)
2147 {
2148   return __internal_float2ll(a, cudaRoundZero);
2149 }
2150
2151 static long long int __float2ll_ru(float a)
2152 {
2153   return __internal_float2ll(a, cudaRoundPosInf);
2154 }
2155
2156 static long long int __float2ll_rd(float a)
2157 {
2158   return __internal_float2ll(a, cudaRoundMinInf);
2159 }
2160
2161 static long long int __float2ll_rn(float a)
2162 {
2163   return __internal_float2ll(a, cudaRoundNearest);
2164 }
2165
2166 static unsigned long long int __internal_float2ull_kernel(float a, unsigned long long int max, unsigned long long int nan, enum cudaRoundMode rndMode)
2167 {
2168   unsigned long long int res, t = 0ULL;
2169   int shift;
2170   unsigned int ia;
2171
2172   if (sizeof(a) == sizeof(double) && __cuda___isnan((double)a)) return nan; if (sizeof(a) == sizeof(float) && __cuda___isnanf((float)a)) return nan; if (a >= max) return max; if (a <= 0LL) return 0LL;
2173   ia = __float_as_int(a);
2174   shift = 190 - ((ia >> 23) & 0xff);
2175   res = (unsigned long long int)((ia << 8) | 0x80000000) << 32;
2176   if (shift >= 64) {
2177     t = res >> (int)(shift > 64);
2178     res = 0;
2179   } else if (shift) {
2180     t = res << (64 - shift);
2181     res = res >> shift;
2182   }
2183   if (rndMode == cudaRoundNearest && (long long int)t < 0LL) {
2184     res += t == 0x8000000000000000ULL ? res & 1ULL : 1ULL;
2185   }
2186   else if (rndMode == cudaRoundPosInf && t != 0ULL) {
2187     res++;
2188   }
2189   return res;
2190 }
2191
2192 static unsigned int __internal_float2uint(float a, enum cudaRoundMode rndMode)
2193 {
2194   return (unsigned int)__internal_float2ull_kernel(a, 4294967295U, 0U, rndMode);
2195 }
2196
2197 static unsigned int __float2uint_rz(float a)
2198 {
2199   return __internal_float2uint(a, cudaRoundZero);
2200 }
2201
2202 static unsigned int __float2uint_ru(float a)
2203 {
2204   return __internal_float2uint(a, cudaRoundPosInf);
2205 }
2206
2207 static unsigned int __float2uint_rd(float a)
2208 {
2209   return __internal_float2uint(a, cudaRoundMinInf);
2210 }
2211
2212 static unsigned int __float2uint_rn(float a)
2213 {
2214   return __internal_float2uint(a, cudaRoundNearest);
2215 }
2216
2217 static unsigned long long int __internal_float2ull(float a, enum cudaRoundMode rndMode)
2218 {
2219   return __internal_float2ull_kernel(a, 18446744073709551615ULL, 9223372036854775808ULL, rndMode);
2220 }
2221
2222 static unsigned long long int __float2ull_rz(float a)
2223 {
2224   return __internal_float2ull(a, cudaRoundZero);
2225 }
2226
2227 static unsigned long long int __float2ull_ru(float a)
2228 {
2229   return __internal_float2ull(a, cudaRoundPosInf);
2230 }
2231
2232 static unsigned long long int __float2ull_rd(float a)
2233 {
2234   return __internal_float2ull(a, cudaRoundMinInf);
2235 }
2236
2237 static unsigned long long int __float2ull_rn(float a)
2238 {
2239   return __internal_float2ull(a, cudaRoundNearest);
2240 }
2241
2242 static int __internal_normalize64(unsigned long long int *a)
2243 {
2244   int lz = 0;
2245
2246   if ((*a & 0xffffffff00000000ULL) == 0ULL) {
2247     *a <<= 32;
2248     lz += 32;
2249   }
2250   if ((*a & 0xffff000000000000ULL) == 0ULL) {
2251     *a <<= 16;
2252     lz += 16;
2253   }
2254   if ((*a & 0xff00000000000000ULL) == 0ULL) {
2255     *a <<= 8;
2256     lz += 8;
2257   }
2258   if ((*a & 0xf000000000000000ULL) == 0ULL) {
2259     *a <<= 4;
2260     lz += 4;
2261   }
2262   if ((*a & 0xC000000000000000ULL) == 0ULL) {
2263     *a <<= 2;
2264     lz += 2;
2265   }
2266   if ((*a & 0x8000000000000000ULL) == 0ULL) {
2267     *a <<= 1;
2268     lz += 1;
2269   }
2270   return lz;
2271 }
2272
2273 static int __internal_normalize(unsigned int *a)
2274 {
2275   unsigned long long int t = (unsigned long long int)*a;
2276   int lz = __internal_normalize64(&t);
2277
2278   *a = (unsigned int)(t >> 32);
2279
2280   return lz - 32;
2281 }
2282
2283 static float __internal_int2float_kernel(int a, enum cudaRoundMode rndMode)
2284 {
2285   volatile union {
2286     float f;
2287     unsigned int i;
2288   } res;
2289   int shift;
2290   unsigned int t;
2291   res.i = a;
2292   if (a == 0) return res.f;
2293   if (a < 0) res.i = (unsigned int)-a;
2294   shift = __internal_normalize((unsigned int*)&res.i);
2295   t = res.i << 24;
2296   res.i = (res.i >> 8);
2297   res.i += (127 + 30 - shift) << 23;
2298   if (a < 0) res.i |= 0x80000000;
2299   if ((rndMode == cudaRoundNearest) && (t >= 0x80000000)) {
2300     res.i += (t == 0x80000000) ? (res.i & 1) : (t >> 31);
2301   }
2302   else if ((rndMode == cudaRoundMinInf) && t && (a < 0)) {
2303     res.i++;
2304   }
2305   else if ((rndMode == cudaRoundPosInf) && t && (a > 0)) {
2306     res.i++;
2307   }
2308   return res.f;
2309 }
2310
2311 static float __int2float_rz(int a)
2312 {
2313   return __internal_int2float_kernel(a, cudaRoundZero);
2314 }
2315
2316 static float __int2float_ru(int a)
2317 {
2318   return __internal_int2float_kernel(a, cudaRoundPosInf);
2319 }
2320
2321 static float __int2float_rd(int a)
2322 {
2323   return __internal_int2float_kernel(a, cudaRoundMinInf);
2324 }
2325
2326 static float __int2float_rn(int a)
2327 {
2328   return __internal_int2float_kernel(a, cudaRoundNearest);
2329 }
2330
2331 static float __internal_uint2float_kernel(unsigned int a, enum cudaRoundMode rndMode)
2332 {
2333   volatile union {
2334     float f;
2335     unsigned int i;
2336   } res;
2337   int shift;
2338   unsigned int t;
2339   res.i = a;
2340   if (a == 0) return res.f;
2341   shift = __internal_normalize((unsigned int*)&res.i);
2342   t = res.i << 24;
2343   res.i = (res.i >> 8);
2344   res.i += (127 + 30 - shift) << 23;
2345   if ((rndMode == cudaRoundNearest) && (t >= 0x80000000)) {
2346     res.i += (t == 0x80000000) ? (res.i & 1) : (t >> 31);
2347   }
2348   else if ((rndMode == cudaRoundPosInf) && t) {
2349     res.i++;
2350   }
2351   return res.f;
2352 }
2353
2354 static float __uint2float_rz(unsigned int a)
2355 {
2356   return __internal_uint2float_kernel(a, cudaRoundZero);
2357 }
2358
2359 static float __uint2float_ru(unsigned int a)
2360 {
2361   return __internal_uint2float_kernel(a, cudaRoundPosInf);
2362 }
2363
2364 static float __uint2float_rd(unsigned int a)
2365 {
2366   return __internal_uint2float_kernel(a, cudaRoundMinInf);
2367 }
2368
2369 static float __uint2float_rn(unsigned int a)
2370 {
2371   return __internal_uint2float_kernel(a, cudaRoundNearest);
2372 }
2373
2374 static float __ll2float_rn(long long int a)
2375 {
2376   return (float)a;
2377 }
2378
2379 static float __ull2float_rn(unsigned long long int a)
2380 {
2381   unsigned long long int temp;
2382   unsigned int res, t;
2383   int shift;
2384   if (a == 0ULL) return 0.0f;
2385   temp = a;
2386   shift = __internal_normalize64(&temp);
2387   temp = (temp >> 8) | ((temp & 0xffULL) ? 1ULL : 0ULL);
2388   res = (unsigned int)(temp >> 32);
2389   t = (unsigned int)temp;
2390   res += (127 + 62 - shift) << 23;
2391   res += t == 0x80000000 ? res & 1 : t >> 31;
2392   return __int_as_float(res);
2393 }
2394
2395 static float __internal_fmul_kernel(float a, float b, int rndNearest)
2396 {
2397   unsigned long long product;
2398   volatile union {
2399     float f;
2400     unsigned int i;
2401   } xx, yy;
2402   unsigned expo_x, expo_y;
2403
2404   xx.f = a;
2405   yy.f = b;
2406
2407   expo_y = 0xFF;
2408   expo_x = expo_y & (xx.i >> 23);
2409   expo_x = expo_x - 1;
2410   expo_y = expo_y & (yy.i >> 23);
2411   expo_y = expo_y - 1;
2412
2413   if ((expo_x <= 0xFD) &&
2414       (expo_y <= 0xFD)) {
2415 multiply:
2416     expo_x = expo_x + expo_y;
2417     expo_y = xx.i ^ yy.i;
2418     xx.i = xx.i & 0x00ffffff;
2419     yy.i = yy.i << 8;
2420     xx.i = xx.i | 0x00800000;
2421     yy.i = yy.i | 0x80000000;
2422
2423     product = ((unsigned long long)xx.i) * yy.i;
2424     expo_x = expo_x - 127 + 2;
2425     expo_y = expo_y & 0x80000000;
2426     xx.i = (unsigned int)(product >> 32);
2427     yy.i = (unsigned int)(product & 0xffffffff);
2428
2429     if (xx.i < 0x00800000) {
2430       xx.i = (xx.i << 1) | (yy.i >> 31);
2431       yy.i = (yy.i << 1);
2432       expo_x--;
2433     }
2434     if (expo_x <= 0xFD) {
2435       xx.i = xx.i | expo_y;
2436       xx.i = xx.i + (expo_x << 23);
2437
2438       if (yy.i < 0x80000000) return xx.f;
2439       xx.i += (((yy.i == 0x80000000) ? (xx.i & 1) : (yy.i >> 31))
2440                && rndNearest);
2441       return xx.f;
2442     } else if ((int)expo_x >= 254) {
2443
2444       xx.i = (expo_y | 0x7F800000) - (!rndNearest);
2445       return xx.f;
2446     } else {
2447
2448       expo_x = ((unsigned int)-((int)expo_x));
2449       if (expo_x > 25) {
2450
2451         xx.i = expo_y;
2452         return xx.f;
2453       } else {
2454         yy.i = (xx.i << (32 - expo_x)) | ((yy.i) ? 1 : 0);
2455         xx.i = expo_y + (xx.i >> expo_x);
2456         xx.i += (((yy.i == 0x80000000) ? (xx.i & 1) : (yy.i >> 31))
2457                  && rndNearest);
2458         return xx.f;
2459       }
2460     }
2461   } else {
2462     product = xx.i ^ yy.i;
2463     product = product & 0x80000000;
2464     if (!(xx.i & 0x7fffffff)) {
2465       if (expo_y != 254) {
2466         xx.i = (unsigned int)product;
2467         return xx.f;
2468       }
2469       expo_y = yy.i << 1;
2470       if (expo_y == 0xFF000000) {
2471         xx.i = expo_y | 0x00C00000;
2472       } else {
2473         xx.i = yy.i | 0x00400000;
2474       }
2475       return xx.f;
2476     }
2477     if (!(yy.i & 0x7fffffff)) {
2478       if (expo_x != 254) {
2479         xx.i = (unsigned int)product;
2480         return xx.f;
2481       }
2482       expo_x = xx.i << 1;
2483       if (expo_x == 0xFF000000) {
2484         xx.i = expo_x | 0x00C00000;
2485       } else {
2486         xx.i = xx.i | 0x00400000;
2487       }
2488       return xx.f;
2489     }
2490     if ((expo_y != 254) && (expo_x != 254)) {
2491       expo_y++;
2492       expo_x++;
2493       if (expo_x == 0) {
2494         expo_y |= xx.i & 0x80000000;
2495
2496
2497
2498
2499         xx.i = xx.i << 8;
2500         while (!(xx.i & 0x80000000)) {
2501           xx.i <<= 1;
2502           expo_x--;
2503         }
2504         xx.i = (xx.i >> 8) | (expo_y & 0x80000000);
2505         expo_y &= ~0x80000000;
2506         expo_y--;
2507         goto multiply;
2508       }
2509       if (expo_y == 0) {
2510         expo_x |= yy.i & 0x80000000;
2511         yy.i = yy.i << 8;
2512         while (!(yy.i & 0x80000000)) {
2513           yy.i <<= 1;
2514           expo_y--;
2515         }
2516         yy.i = (yy.i >> 8) | (expo_x & 0x80000000);
2517         expo_x &= ~0x80000000;
2518         expo_x--;
2519         goto multiply;
2520       }
2521     }
2522     expo_x = xx.i << 1;
2523     expo_y = yy.i << 1;
2524
2525     if (expo_x > 0xFF000000) {
2526
2527       xx.i = xx.i | 0x00400000;
2528       return xx.f;
2529     }
2530
2531     if (expo_y > 0xFF000000) {
2532
2533       xx.i = yy.i | 0x00400000;
2534       return xx.f;
2535     }
2536     xx.i = (unsigned int)product | 0x7f800000;
2537     return xx.f;
2538   }
2539 }
2540
2541 static float __internal_fadd_kernel(float a, float b, int rndNearest)
2542 {
2543   volatile union {
2544     float f;
2545     unsigned int i;
2546   } xx, yy;
2547   unsigned int expo_x;
2548   unsigned int expo_y;
2549   unsigned int temp;
2550
2551   xx.f = a;
2552   yy.f = b;
2553
2554
2555   expo_y = yy.i << 1;
2556   if (expo_y > (xx.i << 1)) {
2557     expo_y = xx.i;
2558     xx.i = yy.i;
2559     yy.i = expo_y;
2560   }
2561
2562   temp = 0xff;
2563   expo_x = temp & (xx.i >> 23);
2564   expo_x = expo_x - 1;
2565   expo_y = temp & (yy.i >> 23);
2566   expo_y = expo_y - 1;
2567
2568   if ((expo_x <= 0xFD) &&
2569       (expo_y <= 0xFD)) {
2570
2571 add:
2572     expo_y = expo_x - expo_y;
2573     if (expo_y > 25) {
2574       expo_y = 31;
2575     }
2576     temp = xx.i ^ yy.i;
2577     xx.i = xx.i & ~0x7f000000;
2578     xx.i = xx.i | 0x00800000;
2579     yy.i = yy.i & ~0xff000000;
2580     yy.i = yy.i | 0x00800000;
2581
2582     if ((int)temp < 0) {
2583
2584       temp = 32 - expo_y;
2585       temp = (expo_y) ? (yy.i << temp) : 0;
2586       temp = (unsigned int)(-((int)temp));
2587       xx.i = xx.i - (yy.i >> expo_y) - (temp ? 1 : 0);
2588       if (xx.i & 0x00800000) {
2589         if (expo_x <= 0xFD) {
2590           xx.i = xx.i & ~0x00800000;
2591           xx.i = (xx.i + (expo_x << 23)) + 0x00800000;
2592           if (temp < 0x80000000) return xx.f;
2593           xx.i += (((temp == 0x80000000) ? (xx.i & 1) : (temp >> 31))
2594                    && rndNearest);
2595           return xx.f;
2596         }
2597       } else {
2598         if ((temp | (xx.i << 1)) == 0) {
2599
2600           xx.i = 0;
2601           return xx.f;
2602         }
2603
2604         yy.i = xx.i & 0x80000000;
2605         do {
2606           xx.i = (xx.i << 1) | (temp >> 31);
2607           temp <<= 1;
2608           expo_x--;
2609         } while (!(xx.i & 0x00800000));
2610         xx.i = xx.i | yy.i;
2611       }
2612     } else {
2613
2614       temp = 32 - expo_y;
2615       temp = (expo_y) ? (yy.i << temp) : 0;
2616       xx.i = xx.i + (yy.i >> expo_y);
2617       if (!(xx.i & 0x01000000)) {
2618         if (expo_x <= 0xFD) {
2619           expo_y = xx.i & 1;
2620           xx.i = xx.i + (expo_x << 23);
2621           if (temp < 0x80000000) return xx.f;
2622           xx.i += (((temp == 0x80000000) ? expo_y : (temp >> 31))
2623                    && rndNearest);
2624           return xx.f;
2625         }
2626       } else {
2627
2628         temp = (xx.i << 31) | (temp >> 1);
2629
2630         xx.i = ((xx.i & 0x80000000) | (xx.i >> 1)) & ~0x40000000;
2631         expo_x++;
2632       }
2633     }
2634     if (expo_x <= 0xFD) {
2635       expo_y = xx.i & 1;
2636       xx.i += (((temp == 0x80000000) ? expo_y : (temp >> 31))
2637                && rndNearest);
2638       xx.i = xx.i + (expo_x << 23);
2639       return xx.f;
2640     }
2641     if ((int)expo_x >= 254) {
2642
2643         xx.i = ((xx.i & 0x80000000) | 0x7f800000) - (!rndNearest);
2644         return xx.f;
2645     }
2646
2647     expo_y = expo_x + 32;
2648     yy.i = xx.i & 0x80000000;
2649     xx.i = xx.i & ~0xff000000;
2650
2651     expo_x = (unsigned int)(-((int)expo_x));
2652     temp = xx.i << expo_y | ((temp) ? 1 : 0);
2653     xx.i = yy.i | (xx.i >> expo_x);
2654     xx.i += (((temp == 0x80000000) ? (xx.i & 1) : (temp >> 31))
2655              && rndNearest);
2656     return xx.f;
2657   } else {
2658
2659     if (!(yy.i << 1)) {
2660       if (xx.i == 0x80000000) {
2661         xx.i = yy.i;
2662       }
2663       return xx.f;
2664     }
2665     if ((expo_y != 254) && (expo_x != 254)) {
2666
2667       if (expo_x == (unsigned int) -1) {
2668         temp = xx.i & 0x80000000;
2669         xx.i = xx.i << 8;
2670         while (!(xx.i & 0x80000000)) {
2671           xx.i <<= 1;
2672           expo_x--;
2673         }
2674         expo_x++;
2675         xx.i = (xx.i >> 8) | temp;
2676       }
2677       if (expo_y == (unsigned int) -1) {
2678         temp = yy.i & 0x80000000;
2679         yy.i = yy.i << 8;
2680         while (!(yy.i & 0x80000000)) {
2681           yy.i <<= 1;
2682           expo_y--;
2683         }
2684         expo_y++;
2685         yy.i = (yy.i >> 8) | temp;
2686       }
2687       goto add;
2688     }
2689     expo_x = xx.i << 1;
2690     expo_y = yy.i << 1;
2691
2692     if (expo_x > 0xff000000) {
2693
2694       xx.i = xx.i | 0x00400000;
2695       return xx.f;
2696     }
2697
2698     if (expo_y > 0xff000000) {
2699
2700       xx.i = yy.i | 0x00400000;
2701       return xx.f;
2702     }
2703     if ((expo_x == 0xff000000) && (expo_y == 0xff000000)) {
2704
2705
2706
2707
2708       expo_x = xx.i ^ yy.i;
2709       xx.i = xx.i | ((expo_x) ? 0xffc00000 : 0);
2710       return xx.f;
2711     }
2712
2713     if (expo_y == 0xff000000) {
2714       xx.i = yy.i;
2715     }
2716     return xx.f;
2717   }
2718 }
2719
2720 static float __fadd_rz(float a, float b)
2721 {
2722   return __internal_fadd_kernel(a, b, 0);
2723 }
2724
2725 static float __fmul_rz(float a, float b)
2726 {
2727   return __internal_fmul_kernel(a, b, 0);
2728 }
2729
2730 static float __fdividef(float a, float b)
2731 {
2732
2733   if (__cuda_fabsf(b) > 8.507059173e37f) {
2734     if (__cuda_fabsf(a) <= 3.402823466e38f) {
2735       return ((a / b) / 3.402823466e38f) / 3.402823466e38f;
2736     } else {
2737       return __int_as_float(0x7fffffff);
2738     }
2739   } else {
2740     return a / b;
2741   }
2742 }
2743
2744 static void __brkpt(int c)
2745 {
2746
2747 }
2748
2749 extern int __cudaSynchronizeThreads(void**, void*);
2750
2751
2752
2753 static inline __attribute__((always_inline)) void __syncthreads(void)
2754 {
2755   volatile int _ = 0;
2756   L: if (__cudaSynchronizeThreads((void**)&&L, (void*)&_)) goto L;
2757 }
2758
2759 static void __trap(void)
2760 {
2761   __builtin_trap();
2762 }
2763 # 1139 "/usr/local/cuda/bin/../include/device_functions.h" 3
2764 static float __sinf(float a)
2765 {
2766   return sinf(a);
2767 }
2768
2769 static float __cosf(float a)
2770 {
2771   return cosf(a);
2772 }
2773
2774 static float __log2f(float a)
2775 {
2776   return log2f(a);
2777 }
2778
2779
2780
2781
2782
2783
2784
2785 static float __internal_accurate_fdividef(float a, float b)
2786 {
2787   if (__cuda_fabsf(b) > 8.507059173e37f) {
2788     a *= .25f;
2789     b *= .25f;
2790   }
2791   return __fdividef(a, b);
2792 }
2793
2794 static float __tanf(float a)
2795 {
2796   return __sinf(a) / __cosf(a);
2797 }
2798
2799 static void __sincosf(float a, float *sptr, float *cptr)
2800 {
2801   *sptr = __sinf(a);
2802   *cptr = __cosf(a);
2803 }
2804
2805 static float __expf(float a)
2806 {
2807   return __cuda_exp2f(a * 1.442695041f);
2808 }
2809
2810 static float __exp10f(float a)
2811 {
2812   return __cuda_exp2f(a * 3.321928094f);
2813 }
2814
2815 static float __log10f(float a)
2816 {
2817   return 0.301029996f * __log2f(a);
2818 }
2819
2820 static float __logf(float a)
2821 {
2822   return 0.693147181f * __log2f(a);
2823 }
2824
2825 static float __powf(float a, float b)
2826 {
2827   return __cuda_exp2f(b * __log2f(a));
2828 }
2829
2830 static float fdividef(float a, float b)
2831 {
2832
2833
2834
2835   return __internal_accurate_fdividef(a, b);
2836
2837 }
2838
2839 static int __clz(int a)
2840 {
2841   return (a)?(158-(__float_as_int(__uint2float_rz((unsigned int)a))>>23)):32;
2842 }
2843
2844 static int __ffs(int a)
2845 {
2846   return 32 - __clz (a & -a);
2847 }
2848
2849 static int __clzll(long long int a)
2850 {
2851   int ahi = ((int)(a >> 32));
2852   int alo = ((int)(a & 0xffffffffULL));
2853   int res;
2854   if (ahi) {
2855       res = 0;
2856   } else {
2857       res = 32;
2858       ahi = alo;
2859   }
2860   res = res + __clz(ahi);
2861   return res;
2862 }
2863
2864 static int __ffsll(long long int a)
2865 {
2866   return 64 - __clzll (a & -a);
2867 }
2868 # 1252 "/usr/local/cuda/bin/../include/device_functions.h" 3
2869 static double fdivide(double a, double b)
2870 {
2871   return (double)fdividef((float)a, (float)b);
2872 }
2873
2874
2875
2876 static int __double2int_rz(double a)
2877 {
2878   return __float2int_rz((float)a);
2879 }
2880
2881 static unsigned int __double2uint_rz(double a)
2882 {
2883   return __float2uint_rz((float)a);
2884 }
2885
2886 static long long int __double2ll_rz(double a)
2887 {
2888   return __float2ll_rz((float)a);
2889 }
2890
2891 static unsigned long long int __double2ull_rz(double a)
2892 {
2893   return __float2ull_rz((float)a);
2894 }
2895 # 1291 "/usr/local/cuda/bin/../include/device_functions.h" 3
2896 # 1 "/usr/local/cuda/bin/../include/sm_11_atomic_functions.h" 1 3
2897 # 214 "/usr/local/cuda/bin/../include/sm_11_atomic_functions.h" 3
2898 static int __iAtomicAdd(int *address, int val)
2899 {
2900   int old = *address;
2901
2902   *address = old + val;
2903
2904   return old;
2905 }
2906
2907 static unsigned int __uAtomicAdd(unsigned int *address, unsigned int val)
2908 {
2909   unsigned int old = *address;
2910
2911   *address = old + val;
2912
2913   return old;
2914 }
2915
2916 static int __iAtomicExch(int *address, int val)
2917 {
2918   int old = *address;
2919
2920   *address = val;
2921
2922   return old;
2923 }
2924
2925 static unsigned int __uAtomicExch(unsigned int *address, unsigned int val)
2926 {
2927   unsigned int old = *address;
2928
2929   *address = val;
2930
2931   return old;
2932 }
2933
2934 static float __fAtomicExch(float *address, float val)
2935 {
2936   float old = *address;
2937
2938   *address = val;
2939
2940   return old;
2941 }
2942
2943 static int __iAtomicMin(int *address, int val)
2944 {
2945   int old = *address;
2946
2947   *address = old < val ? old : val;
2948
2949   return old;
2950 }
2951
2952 static unsigned int __uAtomicMin(unsigned int *address, unsigned int val)
2953 {
2954   unsigned int old = *address;
2955
2956   *address = old < val ? old : val;
2957
2958   return old;
2959 }
2960
2961 static int __iAtomicMax(int *address, int val)
2962 {
2963   int old = *address;
2964
2965   *address = old > val ? old : val;
2966
2967   return old;
2968 }
2969
2970 static unsigned int __uAtomicMax(unsigned int *address, unsigned int val)
2971 {
2972   unsigned int old = *address;
2973
2974   *address = old > val ? old : val;
2975
2976   return old;
2977 }
2978
2979 static unsigned int __uAtomicInc(unsigned int *address, unsigned int val)
2980 {
2981   unsigned int old = *address;
2982
2983   *address = (old >= val) ? 0 : old + 1;
2984
2985   return old;
2986 }
2987
2988 static unsigned int __uAtomicDec(unsigned int *address, unsigned int val)
2989 {
2990   unsigned int old = *address;
2991
2992   *address = ((old == 0) | (old > val)) ? val : (old - 1);
2993
2994   return old;
2995 }
2996
2997 static int __iAtomicAnd(int *address, int val)
2998 {
2999   int old = *address;
3000
3001   *address = old & val;
3002
3003   return old;
3004 }
3005
3006 static unsigned int __uAtomicAnd(unsigned int *address, unsigned int val)
3007 {
3008   unsigned int old = *address;
3009
3010   *address = old & val;
3011
3012   return old;
3013 }
3014
3015 static int __iAtomicOr(int *address, int val)
3016 {
3017   int old = *address;
3018
3019   *address = old | val;
3020
3021   return old;
3022 }
3023
3024 static unsigned int __uAtomicOr(unsigned int *address, unsigned int val)
3025 {
3026   unsigned int old = *address;
3027
3028   *address = old | val;
3029
3030   return old;
3031 }
3032
3033 static int __iAtomicXor(int *address, int val)
3034 {
3035   int old = *address;
3036
3037   *address = old ^ val;
3038
3039   return old;
3040 }
3041
3042 static unsigned int __uAtomicXor(unsigned int *address, unsigned int val)
3043 {
3044   unsigned int old = *address;
3045
3046   *address = old ^ val;
3047
3048   return old;
3049 }
3050
3051 static int __iAtomicCAS(int *address, int compare, int val)
3052 {
3053   int old = *address;
3054
3055   *address = old == compare ? val : old;
3056
3057   return old;
3058 }
3059
3060 static unsigned int __uAtomicCAS(unsigned int *address, unsigned int compare, unsigned int val)
3061 {
3062   unsigned int old = *address;
3063
3064   *address = old == compare ? val : old;
3065
3066   return old;
3067 }
3068 # 1292 "/usr/local/cuda/bin/../include/device_functions.h" 2 3
3069 # 1 "/usr/local/cuda/bin/../include/texture_fetch_functions.h" 1 3
3070 # 2007 "/usr/local/cuda/bin/../include/texture_fetch_functions.h" 3
3071 extern void __cudaTextureFetch(const void *tex, void *index, int integer, void *val);
3072
3073 static int4 __itexfetchi(const void *tex, int4 index)
3074 {
3075   int4 val;
3076
3077   __cudaTextureFetch(tex, &index, 1, &val);
3078
3079   return val;
3080 }
3081
3082 static uint4 __utexfetchi(const void *tex, int4 index)
3083 {
3084   uint4 val;
3085
3086   __cudaTextureFetch(tex, &index, 1, &val);
3087
3088   return val;
3089 }
3090
3091 static float4 __ftexfetchi(const void *tex, int4 index)
3092 {
3093   float4 val;
3094
3095   __cudaTextureFetch(tex, &index, 1, &val);
3096
3097   return val;
3098 }
3099
3100 static int4 __itexfetch(const void *tex, float4 index, int dim)
3101 {
3102   int4 val;
3103
3104   __cudaTextureFetch(tex, &index, 0, &val);
3105
3106   return val;
3107 }
3108
3109 static uint4 __utexfetch(const void *tex, float4 index, int dim)
3110 {
3111   uint4 val;
3112
3113   __cudaTextureFetch(tex, &index, 0, &val);
3114
3115   return val;
3116 }
3117
3118 static float4 __ftexfetch(const void *tex, float4 index, int dim)
3119 {
3120   float4 val;
3121
3122   __cudaTextureFetch(tex, &index, 0, &val);
3123
3124   return val;
3125 }
3126 # 1293 "/usr/local/cuda/bin/../include/device_functions.h" 2 3
3127 # 859 "/usr/local/cuda/bin/../include/math_functions.h" 2 3
3128
3129
3130 extern __attribute__((weak)) int __cuda___signbitf(float a); int __cuda___signbitf(float a)
3131 {
3132   return (int)((unsigned int)__float_as_int(a) >> 31);
3133 }
3134
3135
3136
3137
3138 extern __attribute__((weak)) float __cuda_copysignf(float a, float b); float __cuda_copysignf(float a, float b)
3139 {
3140   return __int_as_float((__float_as_int(b) & 0x80000000) |
3141                         (__float_as_int(a) & ~0x80000000));
3142 }
3143 # 883 "/usr/local/cuda/bin/../include/math_functions.h" 3
3144 extern __attribute__((weak)) int min(int a, int b); int min(int a, int b)
3145 {
3146   return a < b ? a : b;
3147 }
3148
3149 extern __attribute__((weak)) unsigned int umin(unsigned int a, unsigned int b); unsigned int umin(unsigned int a, unsigned int b)
3150 {
3151   return a < b ? a : b;
3152 }
3153
3154 extern __attribute__((weak)) int max(int a, int b); int max(int a, int b)
3155 {
3156   return a > b ? a : b;
3157 }
3158
3159 extern __attribute__((weak)) unsigned int umax(unsigned int a, unsigned int b); unsigned int umax(unsigned int a, unsigned int b)
3160 {
3161   return a > b ? a : b;
3162 }
3163 # 967 "/usr/local/cuda/bin/../include/math_functions.h" 3
3164 extern __attribute__((weak)) float __internal_nearbyintf(float a); float __internal_nearbyintf(float a)
3165 {
3166   float fa = fabsf(a);
3167
3168   if (fa >= 8388608.0f) {
3169     return a;
3170   } else {
3171     volatile float u = 8388608.0f + fa;
3172
3173     u = u - 8388608.0f;
3174     return copysignf(u, a);
3175   }
3176 }
3177
3178 extern __attribute__((weak)) float __internal_fminf(float a, float b); float __internal_fminf(float a, float b)
3179 {
3180   volatile union {
3181     float f;
3182     unsigned int i;
3183   } cvta, cvtb;
3184
3185   cvta.f = a;
3186   cvtb.f = b;
3187   if ((cvta.i << 1) > 0xff000000) return b;
3188   if ((cvtb.i << 1) > 0xff000000) return a;
3189   if ((cvta.i | cvtb.i) == 0x80000000) {
3190     return __int_as_float(0x80000000);
3191   }
3192   return a < b ? a : b;
3193 }
3194
3195 extern __attribute__((weak)) float __internal_fmaxf(float a, float b); float __internal_fmaxf(float a, float b)
3196 {
3197   volatile union {
3198     float f;
3199     unsigned int i;
3200   } cvta, cvtb;
3201
3202   cvta.f = a;
3203   cvtb.f = b;
3204   if ((cvta.i << 1) > 0xff000000) return b;
3205   if ((cvtb.i << 1) > 0xff000000) return a;
3206   if ((cvta.f == 0.0f) && (cvtb.f == 0.0f)) {
3207     cvta.i &= cvtb.i;
3208     return cvta.f;
3209   }
3210   return a > b ? a : b;
3211 }
3212 # 1055 "/usr/local/cuda/bin/../include/math_functions.h" 3
3213 extern __attribute__((weak)) long int __cuda_labs(long int a); long int __cuda_labs(long int a)
3214 {
3215   return labs(a);
3216 }
3217
3218 extern __attribute__((weak)) float __cuda_ceilf(float a); float __cuda_ceilf(float a)
3219 {
3220   return ceilf(a);
3221 }
3222
3223 extern __attribute__((weak)) float __cuda_floorf(float a); float __cuda_floorf(float a)
3224 {
3225   return floorf(a);
3226 }
3227
3228 extern __attribute__((weak)) float __cuda_sqrtf(float a); float __cuda_sqrtf(float a)
3229 {
3230    return sqrtf(a);
3231 }
3232
3233 extern __attribute__((weak)) float __cuda_rsqrtf(float a); float __cuda_rsqrtf(float a)
3234 {
3235    return 1.0f / sqrtf(a);
3236 }
3237
3238 extern __attribute__((weak)) float __cuda_truncf(float a); float __cuda_truncf(float a)
3239 {
3240   return truncf(a);
3241 }
3242
3243 extern __attribute__((weak)) int __cuda_max(int a, int b); int __cuda_max(int a, int b)
3244 {
3245   return max(a, b);
3246 }
3247
3248 extern __attribute__((weak)) int __cuda_min(int a, int b); int __cuda_min(int a, int b)
3249 {
3250   return min(a, b);
3251 }
3252
3253 extern __attribute__((weak)) unsigned int __cuda_umax(unsigned int a, unsigned int b); unsigned int __cuda_umax(unsigned int a, unsigned int b)
3254 {
3255   return umax(a, b);
3256 }
3257
3258 extern __attribute__((weak)) unsigned int __cuda_umin(unsigned int a, unsigned int b); unsigned int __cuda_umin(unsigned int a, unsigned int b)
3259 {
3260   return umin(a, b);
3261 }
3262
3263 extern __attribute__((weak)) long long int __cuda_llrintf(float a); long long int __cuda_llrintf(float a)
3264 {
3265   return __float2ll_rn(a);
3266 }
3267
3268 extern __attribute__((weak)) long int __cuda_lrintf(float a); long int __cuda_lrintf(float a)
3269 {
3270
3271   return (long int)__cuda_llrintf(a);
3272
3273
3274
3275 }
3276
3277 extern __attribute__((weak)) float __cuda_nearbyintf(float a); float __cuda_nearbyintf(float a)
3278 {
3279
3280
3281
3282   return __internal_nearbyintf(a);
3283
3284 }
3285
3286 extern __attribute__((weak)) float __cuda_fmaxf(float a, float b); float __cuda_fmaxf(float a, float b)
3287 {
3288
3289
3290
3291   return __internal_fmaxf(a, b);
3292
3293 }
3294
3295 extern __attribute__((weak)) float __cuda_fminf(float a, float b); float __cuda_fminf(float a, float b)
3296 {
3297
3298
3299
3300   return __internal_fminf(a, b);
3301
3302 }
3303 # 1162 "/usr/local/cuda/bin/../include/math_functions.h" 3
3304 extern __attribute__((weak)) int __cuda___finitef(float a); int __cuda___finitef(float a)
3305 {
3306   return __cuda_fabsf(a) < __int_as_float(0x7f800000);
3307 }
3308
3309 extern __attribute__((weak)) int __cuda___isinff(float a); int __cuda___isinff(float a)
3310 {
3311   return __cuda_fabsf(a) == __int_as_float(0x7f800000);
3312 }
3313
3314 extern __attribute__((weak)) int __cuda___isnanf(float a); int __cuda___isnanf(float a)
3315 {
3316   return !(__cuda_fabsf(a) <= __int_as_float(0x7f800000));
3317 }
3318
3319 extern __attribute__((weak)) float __cuda_nextafterf(float a, float b); float __cuda_nextafterf(float a, float b)
3320 {
3321   unsigned int ia;
3322   unsigned int ib;
3323   ia = __float_as_int(a);
3324   ib = __float_as_int(b);
3325
3326
3327
3328
3329   if (__cuda___isnanf(a) || __cuda___isnanf(b)) return a + b;
3330   if (__int_as_float (ia | ib) == 0.0f) return b;
3331
3332
3333
3334
3335
3336   if (__int_as_float(ia) == 0.0f) {
3337     return __cuda_copysignf(__int_as_float(0x00000001), b);
3338   }
3339
3340   if ((a < b) && (a < 0.0f)) ia--;
3341   if ((a < b) && (a > 0.0f)) ia++;
3342   if ((a > b) && (a < 0.0f)) ia++;
3343   if ((a > b) && (a > 0.0f)) ia--;
3344   a = __int_as_float(ia);
3345
3346
3347
3348
3349
3350   return a;
3351 }
3352
3353 extern __attribute__((weak)) float __cuda_nanf(const char *tagp); float __cuda_nanf(const char *tagp)
3354 {
3355
3356   return __int_as_float(0x7fffffff);
3357 }
3358
3359
3360 extern __attribute__((weak)) float __internal_atanhf_kernel(float a_1, float a_2); float __internal_atanhf_kernel(float a_1, float a_2)
3361 {
3362   float a, a2, t;
3363
3364   a = a_1 + a_2;
3365   a2 = a * a;
3366   t = 1.566305595598990E-001f/64.0f;
3367   t = t * a2 + 1.995081856004762E-001f/16.0f;
3368   t = t * a2 + 3.333382699617026E-001f/4.0f;
3369   t = t * a2;
3370   t = t * a + a_2;
3371   t = t + a_1;
3372   return t;
3373 }
3374
3375
3376
3377
3378 extern __attribute__((weak)) float __internal_atanf_kernel(float a); float __internal_atanf_kernel(float a)
3379 {
3380   float t4, t0, t1;
3381
3382   t4 = a * a;
3383   t0 = - 5.674867153f;
3384   t0 = t4 * - 0.823362947f + t0;
3385   t0 = t0 * t4 - 6.565555096f;
3386   t0 = t0 * t4;
3387   t0 = t0 * a;
3388   t1 = t4 + 11.33538818f;
3389   t1 = t1 * t4 + 28.84246826f;
3390   t1 = t1 * t4 + 19.69667053f;
3391   t1 = 1.0f / t1;
3392   a = t0 * t1 + a;
3393   return a;
3394 }
3395
3396
3397 extern __attribute__((weak)) float __internal_tan_kernel(float a); float __internal_tan_kernel(float a)
3398 {
3399   float a2, s, t;
3400
3401   a2 = a * a;
3402   t = 4.114678393115178E-003f * a2 - 8.231194034909670E-001f;
3403   s = a2 - 2.469348886157666E+000f;
3404   s = 1.0f / s;
3405   t = t * s;
3406   t = t * a2;
3407   t = t * a + a;
3408   return t;
3409 }
3410
3411 extern __attribute__((weak)) float __internal_accurate_logf(float a); float __internal_accurate_logf(float a)
3412 {
3413   float t;
3414   float z;
3415   float m;
3416   int ia, e;
3417   ia = __float_as_int(a);
3418
3419   if ((ia < 0x00800000) || (ia > 0x7f7fffff)) {
3420     return __logf(a);
3421   }
3422
3423   m = __int_as_float((ia & 0x807fffff) | 0x3f800000);
3424   e = ((unsigned)ia >> 23) - 127;
3425   if (m > 1.414213562f) {
3426     m = m * 0.5f;
3427     e = e + 1;
3428   }
3429   t = m - 1.0f;
3430   z = m + 1.0f;
3431   z = t / z;
3432   z = -t * z;
3433   z = __internal_atanhf_kernel(t, z);
3434   z = (float)e * 0.693147181f + z;
3435   return z;
3436 }
3437
3438 extern __attribute__((weak)) float __internal_accurate_log2f(float a); float __internal_accurate_log2f(float a)
3439 {
3440   return 1.442695041f * __internal_accurate_logf(a);
3441 }
3442
3443
3444 static unsigned int __cudart_i2opi_f [] = {
3445   0x3c439041,
3446   0xdb629599,
3447   0xf534ddc0,
3448   0xfc2757d1,
3449   0x4e441529,
3450   0xa2f9836e,
3451 };
3452
3453
3454 extern __attribute__((weak)) float __internal_trig_reduction_kernel(float a, int *quadrant); float __internal_trig_reduction_kernel(float a, int *quadrant)
3455 {
3456   float j;
3457   int q;
3458   if (__cuda_fabsf(a) > 48039.0f) {
3459
3460     unsigned int ia = __float_as_int(a);
3461     unsigned int s = ia & 0x80000000;
3462     unsigned int result[7];
3463     unsigned int phi, plo;
3464     unsigned int hi, lo;
3465     unsigned int e;
3466     int idx;
3467     e = ((ia >> 23) & 0xff) - 128;
3468     ia = (ia << 8) | 0x80000000;
3469
3470     idx = 4 - (e >> 5);
3471     hi = 0;
3472
3473
3474
3475     for (q = 0; q < 6; q++) {
3476       plo = __cudart_i2opi_f[q] * ia;
3477       phi = __umulhi (__cudart_i2opi_f[q], ia);
3478       lo = hi + plo;
3479       hi = phi + (lo < plo);
3480       result[q] = lo;
3481     }
3482     result[q] = hi;
3483     e = e & 31;
3484
3485
3486
3487     hi = result[idx+2];
3488     lo = result[idx+1];
3489     if (e) {
3490       q = 32 - e;
3491       hi = (hi << e) | (lo >> q);
3492       lo = (lo << e) | (result[idx] >> q);
3493     }
3494     q = hi >> 30;
3495
3496     hi = (hi << 2) | (lo >> 30);
3497     lo = (lo << 2);
3498     e = (hi + (lo > 0)) > 0x80000000;
3499     q += e;
3500     if (s) q = -q;
3501     if (e) {
3502       unsigned int t;
3503       hi = ~hi;
3504       lo = -(int)lo;
3505       t = (lo == 0);
3506       hi += t;
3507       s = s ^ 0x80000000;
3508     }
3509     *quadrant = q;
3510
3511     e = 0;
3512     while ((int)hi > 0) {
3513       hi = (hi << 1) | (lo >> 31);
3514       lo = (lo << 1);
3515       e--;
3516     }
3517     lo = hi * 0xc90fdaa2;
3518     hi = __umulhi(hi, 0xc90fdaa2);
3519     if ((int)hi > 0) {
3520       hi = (hi << 1) | (lo >> 31);
3521       lo = (lo << 1);
3522       e--;
3523     }
3524     hi = hi + (lo > 0);
3525     ia = s | (((e + 126) << 23) + (hi >> 8) + ((hi << 24) >= 0x80000000));
3526     return __int_as_float(ia);
3527   }
3528   q = __float2int_rn(a * 0.636619772f);
3529   j = (float)q;
3530   a = a - j * 1.5703125000000000e+000f;
3531   a = a - j * 4.8351287841796875e-004f;
3532   a = a - j * 3.1385570764541626e-007f;
3533   a = a - j * 6.0771005065061922e-011f;
3534   *quadrant = q;
3535   return a;
3536 }
3537 # 1405 "/usr/local/cuda/bin/../include/math_functions.h" 3
3538 extern __attribute__((weak)) float __internal_expf_kernel(float a, float scale); float __internal_expf_kernel(float a, float scale)
3539 {
3540   float j, z;
3541
3542   j = __cuda_truncf(a * 1.442695041f);
3543   z = a - j * 0.6931457519f;
3544   z = z - j * 1.4286067653e-6f;
3545   z = z * 1.442695041f;
3546   z = __cuda_exp2f(z) * __cuda_exp2f(j + scale);
3547   return z;
3548 }
3549
3550 extern __attribute__((weak)) float __internal_accurate_expf(float a); float __internal_accurate_expf(float a)
3551 {
3552   float z;
3553   z = __internal_expf_kernel(a, 0.0f);
3554   if (a < -105.0f) z = 0.0f;
3555   if (a > 105.0f) z = __int_as_float(0x7f800000);
3556   return z;
3557 }
3558
3559 extern __attribute__((weak)) float __internal_accurate_exp10f(float a); float __internal_accurate_exp10f(float a)
3560 {
3561   float j, z;
3562   j = __cuda_truncf(a * 3.321928094f);
3563   z = a - j * 3.0102920532226563e-001f;
3564   z = z - j * 7.9034171557301747e-007f;
3565   z = z * 3.321928094f;
3566   z = __cuda_exp2f(z) * __cuda_exp2f(j);
3567   if (a < -46.0f) z = 0.0f;
3568   if (a > 46.0f) z = __int_as_float(0x7f800000);
3569   return z;
3570 }
3571
3572 extern __attribute__((weak)) float __internal_lgammaf_pos(float a); float __internal_lgammaf_pos(float a)
3573 {
3574   float sum;
3575   float s, t;
3576
3577   if (__cuda___isinff(a)) {
3578     return a;
3579   }
3580   if (a >= 3.0f) {
3581     if (a >= 7.8f) {
3582
3583
3584
3585       s = 1.0f / a;
3586       t = s * s;
3587       sum = 0.77783067e-3f;
3588       sum = sum * t - 0.2777655457e-2f;
3589       sum = sum * t + 0.83333273853e-1f;
3590       sum = sum * s + 0.918938533204672f;
3591       s = 0.5f * __internal_accurate_logf(a);
3592       t = a - 0.5f;
3593       s = s * t;
3594       t = s - a;
3595       s = s + sum;
3596       t = t + s;
3597       return t;
3598     } else {
3599       a = a - 3.0f;
3600       s = - 7.488903254816711E+002f;
3601       s = s * a - 1.234974215949363E+004f;
3602       s = s * a - 4.106137688064877E+004f;
3603       s = s * a - 4.831066242492429E+004f;
3604       s = s * a - 1.430333998207429E+005f;
3605       t = a - 2.592509840117874E+002f;
3606       t = t * a - 1.077717972228532E+004f;
3607       t = t * a - 9.268505031444956E+004f;
3608       t = t * a - 2.063535768623558E+005f;
3609       t = s / t;
3610       t = t + a;
3611       return t;
3612     }
3613   } else if (a >= 1.5f) {
3614     a = a - 2.0f;
3615     t = + 4.959849168282574E-005f;
3616     t = t * a - 2.208948403848352E-004f;
3617     t = t * a + 5.413142447864599E-004f;
3618     t = t * a - 1.204516976842832E-003f;
3619     t = t * a + 2.884251838546602E-003f;
3620     t = t * a - 7.382757963931180E-003f;
3621     t = t * a + 2.058131963026755E-002f;
3622     t = t * a - 6.735248600734503E-002f;
3623     t = t * a + 3.224670187176319E-001f;
3624     t = t * a + 4.227843368636472E-001f;
3625     t = t * a;
3626     return t;
3627   } else if (a >= 0.7f) {
3628     a = 1.0f - a;
3629     t = + 4.588266515364258E-002f;
3630     t = t * a + 1.037396712740616E-001f;
3631     t = t * a + 1.228036339653591E-001f;
3632     t = t * a + 1.275242157462838E-001f;
3633     t = t * a + 1.432166835245778E-001f;
3634     t = t * a + 1.693435824224152E-001f;
3635     t = t * a + 2.074079329483975E-001f;
3636     t = t * a + 2.705875136435339E-001f;
3637     t = t * a + 4.006854436743395E-001f;
3638     t = t * a + 8.224669796332661E-001f;
3639     t = t * a + 5.772156651487230E-001f;
3640     t = t * a;
3641     return t;
3642   } else {
3643     t = + 3.587515669447039E-003f;
3644     t = t * a - 5.471285428060787E-003f;
3645     t = t * a - 4.462712795343244E-002f;
3646     t = t * a + 1.673177015593242E-001f;
3647     t = t * a - 4.213597883575600E-002f;
3648     t = t * a - 6.558672843439567E-001f;
3649     t = t * a + 5.772153712885004E-001f;
3650     t = t * a;
3651     t = t * a + a;
3652     return -__internal_accurate_logf(t);
3653   }
3654 }
3655
3656
3657 extern __attribute__((weak)) float __internal_sin_kernel(float x); float __internal_sin_kernel(float x)
3658 {
3659   float x2, z;
3660
3661   x2 = x * x;
3662   z = - 1.95152959e-4f;
3663   z = z * x2 + 8.33216087e-3f;
3664   z = z * x2 - 1.66666546e-1f;
3665   z = z * x2;
3666   z = z * x + x;
3667
3668   return z;
3669 }
3670
3671
3672 extern __attribute__((weak)) float __internal_cos_kernel(float x); float __internal_cos_kernel(float x)
3673 {
3674   float x2, z;
3675
3676   x2 = x * x;
3677   z = 2.44331571e-5f;
3678   z = z * x2 - 1.38873163e-3f;
3679   z = z * x2 + 4.16666457e-2f;
3680   z = z * x2 - 5.00000000e-1f;
3681   z = z * x2 + 1.00000000e+0f;
3682   return z;
3683 }
3684
3685 extern __attribute__((weak)) float __internal_accurate_sinf(float a); float __internal_accurate_sinf(float a)
3686 {
3687   float z;
3688   int i;
3689
3690   if (__cuda___isinff(a)) {
3691     return __int_as_float(0x7fffffff);
3692   }
3693   if (a == 0.0f) {
3694     return a;
3695   }
3696   z = __internal_trig_reduction_kernel(a, &i);
3697
3698   if (i & 1) {
3699     z = __internal_cos_kernel(z);
3700   } else {
3701     z = __internal_sin_kernel(z);
3702   }
3703   if (i & 2) {
3704     z = -z;
3705   }
3706   return z;
3707 }
3708
3709
3710
3711
3712
3713
3714
3715 extern __attribute__((weak)) float __cuda_rintf(float a); float __cuda_rintf(float a)
3716 {
3717   return __cuda_nearbyintf(a);
3718 }
3719
3720 extern __attribute__((weak)) float __cuda_sinf(float a); float __cuda_sinf(float a)
3721 {
3722
3723
3724
3725   return __internal_accurate_sinf(a);
3726
3727 }
3728
3729 extern __attribute__((weak)) float __cuda_cosf(float a); float __cuda_cosf(float a)
3730 {
3731
3732
3733
3734   float z;
3735   int i;
3736
3737   if (__cuda___isinff(a)) {
3738     return __int_as_float(0x7fffffff);
3739   }
3740   z = __internal_trig_reduction_kernel(a, &i);
3741
3742   i++;
3743   if (i & 1) {
3744     z = __internal_cos_kernel(z);
3745   } else {
3746     z = __internal_sin_kernel(z);
3747   }
3748   if (i & 2) {
3749     z = -z;
3750   }
3751   return z;
3752
3753 }
3754
3755 extern __attribute__((weak)) float __cuda_tanf(float a); float __cuda_tanf(float a)
3756 {
3757
3758
3759
3760   float z;
3761   int i;
3762
3763   if (__cuda___isinff(a)) {
3764     return __int_as_float(0x7fffffff);
3765   }
3766   z = __internal_trig_reduction_kernel(a, &i);
3767
3768   z = __internal_tan_kernel(z);
3769   if (i & 1) {
3770     z = -1.0f / z;
3771   }
3772   return z;
3773
3774 }
3775
3776 extern __attribute__((weak)) float __cuda_log2f(float a); float __cuda_log2f(float a)
3777 {
3778
3779
3780
3781   return __internal_accurate_log2f(a);
3782
3783 }
3784
3785 extern __attribute__((weak)) float __cuda_expf(float a); float __cuda_expf(float a)
3786 {
3787
3788
3789
3790   return __internal_accurate_expf(a);
3791
3792 }
3793
3794 extern __attribute__((weak)) float __cuda_exp10f(float a); float __cuda_exp10f(float a)
3795 {
3796
3797
3798
3799   return __internal_accurate_exp10f(a);
3800
3801 }
3802
3803 extern __attribute__((weak)) float __cuda_coshf(float a); float __cuda_coshf(float a)
3804 {
3805   float z;
3806
3807   a = __cuda_fabsf(a);
3808   z = __internal_expf_kernel(a, -2.0f);
3809   z = 2.0f * z + 0.125f / z;
3810   if (a >= 90.0f) {
3811     z = __int_as_float(0x7f800000);
3812   }
3813   return z;
3814 }
3815
3816 extern __attribute__((weak)) float __cuda_sinhf(float a); float __cuda_sinhf(float a)
3817 {
3818   float s, z;
3819
3820   s = a;
3821   a = __cuda_fabsf(a);
3822   if (a < 1.0f) {
3823     float a2 = a * a;
3824
3825     z = 2.816951222e-6f;
3826     z = z * a2 + 1.983615978e-4f;
3827     z = z * a2 + 8.333350058e-3f;
3828     z = z * a2 + 1.666666650e-1f;
3829     z = z * a2;
3830     z = z * a + a;
3831   } else {
3832     z = __internal_expf_kernel(a, -2.0f);
3833     z = 2.0f * z - 0.125f / z;
3834     if (a >= 90.0f) {
3835       z = __int_as_float(0x7f800000);
3836     }
3837   }
3838   return __cuda_copysignf(z, s);
3839 }
3840
3841 extern __attribute__((weak)) float __cuda_tanhf(float a); float __cuda_tanhf(float a)
3842 {
3843   float t;
3844
3845   t = __cuda_fabsf(a);
3846   if (t < 0.55f) {
3847     float z, z2;
3848     z = t;
3849     z2 = z * z;
3850     t = 1.643758066599993e-2f;
3851     t = t * z2 - 5.267181327760551e-2f;
3852     t = t * z2 + 1.332072505223051e-1f;
3853     t = t * z2 - 3.333294663641083e-1f;
3854     t = t * z2;
3855     t = t * z + z;
3856   }
3857   else if (t < 88.0f) {
3858     t = 1.0f - 2.0f / (__internal_expf_kernel(2.0f * t, 0.0f) + 1.0f);
3859   }
3860   else if (t >= 88.0f) {
3861     t = 1.0f;
3862   }
3863   return __cuda_copysignf(t, a);
3864 }
3865
3866 extern __attribute__((weak)) float __cuda_atan2f(float a, float b); float __cuda_atan2f(float a, float b)
3867 {
3868   float t0, t1, t3;
3869
3870
3871
3872   t3 = __cuda_fabsf(b);
3873   t1 = __cuda_fabsf(a);
3874
3875   if (t3 == 0.0f && t1 == 0.0f) {
3876     t3 = __cuda___signbitf(b) ? 3.141592654f : 0;
3877   } else if (__cuda___isinff(t3) && __cuda___isinff(t1)) {
3878     t3 = __cuda___signbitf(b) ? 2.356194490f : 0.785398163f;
3879   } else {
3880
3881     if (t3 < t1) {
3882       t0 = t1;
3883       t1 = t3;
3884     } else {
3885       t0 = t3;
3886       t1 = t1;
3887     }
3888     t3 = __internal_accurate_fdividef(t1, t0);
3889     t3 = __internal_atanf_kernel(t3);
3890
3891     if (__cuda_fabsf(a) > __cuda_fabsf(b)) t3 = 1.570796327f - t3;
3892     if (b < 0.0f) t3 = 3.141592654f - t3;
3893   }
3894   t3 = __cuda_copysignf(t3, a);
3895
3896   return t3;
3897 }
3898
3899 extern __attribute__((weak)) float __cuda_atanf(float a); float __cuda_atanf(float a)
3900 {
3901   float t0, t1;
3902
3903
3904   t0 = __cuda_fabsf(a);
3905   t1 = t0;
3906   if (t0 > 1.0f) {
3907     t1 = 1.0f / t1;
3908   }
3909
3910   t1 = __internal_atanf_kernel(t1);
3911
3912   if (t0 > 1.0f) {
3913     t1 = 1.570796327f - t1;
3914   }
3915   return __cuda_copysignf(t1, a);
3916 }
3917
3918
3919 extern __attribute__((weak)) float __internal_asinf_kernel(float a); float __internal_asinf_kernel(float a)
3920 {
3921   float t2, t3, t4;
3922
3923   t2 = a * a;
3924   t3 = - 0.501162291f;
3925   t3 = t3 * t2 + 0.915201485f;
3926   t3 = t3 * t2;
3927   t3 = t3 * a;
3928   t4 = t2 - 5.478654385f;
3929   t4 = t4 * t2 + 5.491230488f;
3930   t4 = 1.0f / t4;
3931   a = t3 * t4 + a;
3932   return a;
3933 }
3934
3935 extern __attribute__((weak)) float __cuda_asinf(float a); float __cuda_asinf(float a)
3936 {
3937   float t0, t1, t2;
3938
3939   t0 = __cuda_fabsf(a);
3940   t2 = 1.0f - t0;
3941   t2 = 0.5f * t2;
3942   t2 = __cuda_sqrtf(t2);
3943   t1 = t0 > 0.575f ? t2 : t0;
3944   t1 = __internal_asinf_kernel(t1);
3945   t2 = -2.0f * t1 + 1.570796327f;
3946   if (t0 > 0.575f) {
3947     t1 = t2;
3948   }
3949   return __cuda_copysignf(t1, a);
3950 }
3951
3952 extern __attribute__((weak)) float __cuda_acosf(float a); float __cuda_acosf(float a)
3953 {
3954   float t0, t1, t2;
3955
3956   t0 = __cuda_fabsf(a);
3957   t2 = 1.0f - t0;
3958   t2 = 0.5f * t2;
3959   t2 = __cuda_sqrtf(t2);
3960   t1 = t0 > 0.575f ? t2 : t0;
3961   t1 = __internal_asinf_kernel(t1);
3962   t1 = t0 > 0.575f ? 2.0f * t1 : 1.570796327f - t1;
3963   if (__cuda___signbitf(a)) {
3964     t1 = 3.141592654f - t1;
3965   }
3966   return t1;
3967 }
3968
3969 extern __attribute__((weak)) float __cuda_logf(float a); float __cuda_logf(float a)
3970 {
3971
3972
3973
3974   return __internal_accurate_logf(a);
3975
3976 }
3977
3978 extern __attribute__((weak)) float __cuda_log10f(float a); float __cuda_log10f(float a)
3979 {
3980
3981
3982
3983   return 0.434294482f * __internal_accurate_logf(a);
3984
3985 }
3986
3987 extern __attribute__((weak)) float __cuda_log1pf(float a); float __cuda_log1pf(float a)
3988 {
3989   float t;
3990
3991
3992
3993
3994   if (a >= -0.394f && a <= 0.65f) {
3995
3996     t = a + 2.0f;
3997     t = a / t;
3998     t = -a * t;
3999     t = __internal_atanhf_kernel (a, t);
4000   } else {
4001     t = __internal_accurate_logf (1.0f + a);
4002   }
4003   return t;
4004 }
4005
4006 extern __attribute__((weak)) float __cuda_acoshf(float a); float __cuda_acoshf(float a)
4007 {
4008   float s, t;
4009
4010   t = a - 1.0f;
4011   if (__cuda_fabsf(t) > 8388608.0f) {
4012
4013     return 0.693147181f + __internal_accurate_logf(a);
4014   } else {
4015     s = a + 1.0f;
4016     t = t + __cuda_sqrtf(s * t);
4017     return __cuda_log1pf(t);
4018   }
4019 }
4020
4021 extern __attribute__((weak)) float __cuda_asinhf(float a); float __cuda_asinhf(float a)
4022 {
4023   float fa, oofa, t;
4024
4025   fa = __cuda_fabsf(a);
4026   if (fa > 8.507059173e37f) {
4027     t = 0.693147181f + __logf(fa);
4028   } else {
4029     oofa = 1.0f / fa;
4030     t = fa + fa / (oofa + __cuda_sqrtf(1.0f + oofa * oofa));
4031     t = __cuda_log1pf(t);
4032   }
4033   return __cuda_copysignf(t, a);
4034 }
4035
4036 extern __attribute__((weak)) float __cuda_atanhf(float a); float __cuda_atanhf(float a)
4037 {
4038   float fa, t;
4039
4040   fa = __cuda_fabsf(a);
4041   t = (2.0f * fa) / (1.0f - fa);
4042   t = 0.5f * __cuda_log1pf(t);
4043   return __cuda_copysignf(t, a);
4044 }
4045
4046 extern __attribute__((weak)) float __cuda_expm1f(float a); float __cuda_expm1f(float a)
4047 {
4048   float t, z, j, u;
4049
4050   t = __cuda_rintf (a * 1.442695041f);
4051   z = a - t * 0.6931457519f;
4052   z = z - t * 1.4286067653e-6f;
4053
4054   if (__cuda_fabsf(a) < 0.41f) {
4055     z = a;
4056     t = 0.0f;
4057   }
4058
4059   j = t;
4060   if (t == 128.0f) j = j - 1.0f;
4061
4062   u = 1.38795078474044430E-003f;
4063   u = u * z + 8.38241261853264930E-003f;
4064   u = u * z + 4.16678317762833940E-002f;
4065   u = u * z + 1.66663978874356580E-001f;
4066   u = u * z + 4.99999940395997040E-001f;
4067   u = u * z;
4068   u = u * z + z;
4069   if (a == 0.0f) u = a;
4070
4071   z = __cuda_exp2f (j);
4072   a = z - 1.0f;
4073   if (a != 0.0f) u = u * z + a;
4074   if (t == 128.0f) u = u + u;
4075
4076   if (j > 128.0f) u = __int_as_float(0x7f800000);
4077   if (j < -25.0f) u = -1.0f;
4078   return u;
4079 }
4080
4081 extern __attribute__((weak)) float __cuda_hypotf(float a, float b); float __cuda_hypotf(float a, float b)
4082 {
4083   float v, w, t;
4084
4085   a = __cuda_fabsf(a);
4086   b = __cuda_fabsf(b);
4087
4088   if (a > b) {
4089     v = a;
4090     w = b;
4091   } else {
4092     v = b;
4093     w = a;
4094   }
4095   t = __internal_accurate_fdividef(w, v);
4096   t = 1.0f + t * t;
4097   t = v * __cuda_sqrtf(t);
4098   if (v == 0.0f) {
4099     t = v + w;
4100   }
4101   if ((v == __int_as_float(0x7f800000)) || (w == __int_as_float(0x7f800000))) {
4102     t = __int_as_float(0x7f800000);
4103   }
4104   return t;
4105 }
4106
4107 extern __attribute__((weak)) float __cuda_cbrtf(float a); float __cuda_cbrtf(float a)
4108 {
4109   float s, t;
4110   if (a == 0.0f || __cuda___isinff(a)) {
4111     return a;
4112   }
4113   s = __cuda_fabsf(a);
4114   t = __cuda_exp2f(0.333333333f * __log2f(s));
4115   t = t - (t - (s / (t * t))) * 0.333333333f;
4116   if (__cuda___signbitf(a)) {
4117      t = -t;
4118   }
4119   return t;
4120 }
4121
4122 extern __attribute__((weak)) float __cuda_erff(float a); float __cuda_erff(float a)
4123 {
4124   float t, r, q;
4125
4126   t = __cuda_fabsf(a);
4127   if (t < 1.0f) {
4128     t = t * t;
4129     r = -5.58510127926029810E-004f;
4130     r = r * t + 4.90688891415893070E-003f;
4131     r = r * t - 2.67027980930150640E-002f;
4132     r = r * t + 1.12799056505903940E-001f;
4133     r = r * t - 3.76122956138427440E-001f;
4134     r = r * t + 1.12837911712623450E+000f;
4135     a = a * r;
4136   } else if (t <= __int_as_float(0x7f800000)) {
4137
4138
4139
4140     q = 0.3275911f * t + 1.0f;
4141     q = 1.0f / q;
4142     r = 1.061405429f;
4143     r = r * q - 1.453152027f;
4144     r = r * q + 1.421413741f;
4145     r = r * q - 0.284496736f;
4146     r = r * q + 0.254829592f;
4147     r = r * q;
4148     q = __internal_expf_kernel(-a * a, 0.0f);
4149     r = 1.0f - q * r;
4150     if (t >= 5.5f) {
4151       r = 1.0f;
4152     }
4153     a = __int_as_float (__float_as_int(r) | (__float_as_int(a) & 0x80000000));
4154   }
4155   return a;
4156 }
4157
4158 extern __attribute__((weak)) float __cuda_erfcf(float a); float __cuda_erfcf(float a)
4159 {
4160   if (a <= 0.55f) {
4161     return 1.0f - __cuda_erff(a);
4162   } else if (a > 10.0f) {
4163     return 0.0f;
4164   } else {
4165     float p;
4166     float q;
4167     float h;
4168     float l;
4169
4170
4171
4172
4173     p = + 4.014893410762552E-006f;
4174     p = p * a + 5.640401259462436E-001f;
4175     p = p * a + 2.626649872281140E+000f;
4176     p = p * a + 5.486372652389673E+000f;
4177     p = p * a + 5.250714831459401E+000f;
4178     q = a + 4.651376250488319E+000f;
4179     q = q * a + 1.026302828878470E+001f;
4180     q = q * a + 1.140762166021288E+001f;
4181     q = q * a + 5.251211619089947E+000f;
4182
4183     h = 1.0f / q;
4184     q = 2.0f * h - q * h * h;
4185     p = p * q;
4186
4187     h = __int_as_float(__float_as_int(a) & 0xfffff000);
4188     l = a - h;
4189     q = -h * h;
4190     q = __internal_expf_kernel(q, 0.0f);
4191     if (l != 0.0f) {
4192       a = a + h;
4193       l = l * a;
4194       h = __internal_expf_kernel(-l, 0.0f);
4195       q = q * h;
4196     }
4197     p = p * q;
4198     return p;
4199   }
4200 }
4201
4202 extern __attribute__((weak)) float __cuda_lgammaf(float a); float __cuda_lgammaf(float a)
4203 {
4204   float t;
4205   float i;
4206   int quot;
4207   t = __internal_lgammaf_pos(__cuda_fabsf(a));
4208   if (a >= 0.0f) return t;
4209   a = __cuda_fabsf(a);
4210   i = __cuda_floorf(a);
4211   if (a == i) return __int_as_float(0x7f800000);
4212   if (a < 1e-19f) return -__internal_accurate_logf(a);
4213   i = __cuda_rintf (2.0f * a);
4214   quot = (int)i;
4215   i = a - 0.5f * i;
4216   i = i * 3.141592654f;
4217   if (quot & 1) {
4218     i = __internal_cos_kernel(i);
4219   } else {
4220     i = __internal_sin_kernel(i);
4221   }
4222   i = __cuda_fabsf(i);
4223   t = 1.144729886f - __internal_accurate_logf(i * a) - t;
4224   return t;
4225 }
4226
4227 extern __attribute__((weak)) float __cuda_ldexpf(float a, int b); float __cuda_ldexpf(float a, int b)
4228 {
4229   float fa = __cuda_fabsf(a);
4230
4231   if (fa == 0.0f || __cuda___isinff(fa) || b == 0) {
4232     return a;
4233   }
4234   else if (__cuda_abs(b) < 126) {
4235     return a * __cuda_exp2f((float)b);
4236   }
4237   else if (__cuda_abs(b) < 252) {
4238     int bhalf = b / 2;
4239     return a * __cuda_exp2f((float)bhalf) * __cuda_exp2f((float)(b - bhalf));
4240   }
4241   else {
4242     int bquarter = b / 4;
4243     float t = __cuda_exp2f((float)bquarter);
4244     return a * t * t * t * __cuda_exp2f((float)(b - 3 * bquarter));
4245   }
4246 }
4247
4248 extern __attribute__((weak)) float __cuda_scalbnf(float a, int b); float __cuda_scalbnf(float a, int b)
4249 {
4250
4251   return __cuda_ldexpf(a, b);
4252 }
4253
4254 extern __attribute__((weak)) float __cuda_scalblnf(float a, long int b); float __cuda_scalblnf(float a, long int b)
4255 {
4256   int t;
4257   if (b > 2147483647L) {
4258     t = 2147483647;
4259   } else if (b < (-2147483647 - 1)) {
4260     t = (-2147483647 - 1);
4261   } else {
4262     t = (int)b;
4263   }
4264   return __cuda_scalbnf(a, t);
4265 }
4266
4267 extern __attribute__((weak)) float __cuda_frexpf(float a, int *b); float __cuda_frexpf(float a, int *b)
4268 {
4269   float fa = __cuda_fabsf(a);
4270   unsigned int expo;
4271   unsigned int denorm;
4272
4273   if (fa < 1.175494351e-38f) {
4274     a *= 16777216.0f;
4275     denorm = 24;
4276   } else {
4277     denorm = 0;
4278   }
4279   expo = ((__float_as_int(a) >> 23) & 0xff);
4280   if ((fa == 0.0f) || (expo == 0xff)) {
4281     expo = 0;
4282     a = a + a;
4283   } else {
4284     expo = expo - denorm - 126;
4285     a = __int_as_float(((__float_as_int(a) & 0x807fffff) | 0x3f000000));
4286   }
4287   *b = expo;
4288   return a;
4289 }
4290
4291 extern __attribute__((weak)) float __cuda_modff(float a, float *b); float __cuda_modff(float a, float *b)
4292 {
4293   float t;
4294   if (__cuda___finitef(a)) {
4295     t = __cuda_truncf(a);
4296     *b = t;
4297     t = a - t;
4298     return __cuda_copysignf(t, a);
4299   } else if (__cuda___isinff(a)) {
4300     t = 0.0f;
4301     *b = a;
4302     return __cuda_copysignf(t, a);
4303   } else {
4304     *b = a;
4305     return a;
4306   }
4307 }
4308
4309 extern __attribute__((weak)) float __cuda_fmodf(float a, float b); float __cuda_fmodf(float a, float b)
4310 {
4311   float orig_a;
4312
4313   if (__cuda___isnanf(a) || __cuda___isnanf(b)) {
4314     return a + b;
4315   }
4316   orig_a = a;
4317   a = __cuda_fabsf(a);
4318   b = __cuda_fabsf(b);
4319   if (__cuda___isinff(a) || b == 0.0f) {
4320     return __int_as_float(0x7fffffff);
4321   } else if (a >= b) {
4322
4323
4324     int expoa = (a < 1.175494351e-38f) ?
4325         ((int)__log2f(a)) : (((__float_as_int(a) >> 23) & 0xff) - 127);
4326     int expob = (b < 1.175494351e-38f) ?
4327         ((int)__log2f(b)) : (((__float_as_int(b) >> 23) & 0xff) - 127);
4328     int scale = expoa - expob;
4329     float scaled_b = __cuda_ldexpf(b, scale);
4330     if (scaled_b <= 0.5f * a) {
4331       scaled_b *= 2.0f;
4332     }
4333
4334
4335
4336
4337
4338
4339
4340     while (scaled_b >= b) {
4341       if (a >= scaled_b) {
4342         a -= scaled_b;
4343       }
4344       scaled_b *= 0.5f;
4345     }
4346     return __cuda_copysignf(a, orig_a);
4347   } else {
4348     return orig_a;
4349   }
4350 }
4351
4352 extern __attribute__((weak)) float __cuda_remainderf(float a, float b); float __cuda_remainderf(float a, float b)
4353 {
4354   float orig_a;
4355   float twoa = 0.0f;
4356   unsigned int quot0 = 0;
4357
4358   if (__cuda___isnanf(a) || __cuda___isnanf(b)) {
4359     return a + b;
4360   }
4361   orig_a = a;
4362   a = __cuda_fabsf(a);
4363   b = __cuda_fabsf(b);
4364   if (__cuda___isinff(a) || (b == 0.0f)) {
4365     return __int_as_float(0x7fffffff);
4366   } else if (a >= b) {
4367
4368     int expoa = (a < 1.175494351e-38f) ?
4369         ((int)__log2f(a)) : (((__float_as_int(a) >> 23) & 0xff) - 127);
4370     int expob = (b < 1.175494351e-38f) ?
4371         ((int)__log2f(b)) : (((__float_as_int(b) >> 23) & 0xff) - 127);
4372     int scale = expoa - expob;
4373     float scaled_b = __cuda_ldexpf(b, scale);
4374     if (scaled_b <= 0.5f * a) {
4375       scaled_b *= 2.0f;
4376     }
4377 # 2255 "/usr/local/cuda/bin/../include/math_functions.h" 3
4378     while (scaled_b >= b) {
4379       quot0 = 0;
4380       if (a >= scaled_b) {
4381         twoa = (2.0f * a - scaled_b) - scaled_b;
4382         a -= scaled_b;
4383         quot0 = 1;
4384       }
4385       scaled_b *= 0.5f;
4386     }
4387   }
4388
4389
4390   twoa = 2.0f * a;
4391   if ((twoa > b) || ((twoa == b) && quot0)) {
4392     a -= b;
4393     a = __cuda_copysignf (a, -1.0f);
4394   }
4395 # 2287 "/usr/local/cuda/bin/../include/math_functions.h" 3
4396   a = __int_as_float((__float_as_int(orig_a) & 0x80000000)^
4397                      __float_as_int(a));
4398   return a;
4399 }
4400
4401 extern __attribute__((weak)) float __cuda_remquof(float a, float b, int* quo); float __cuda_remquof(float a, float b, int* quo)
4402 {
4403   float orig_a;
4404   float twoa = 0.0f;
4405   unsigned int quot = 0;
4406   unsigned int sign;
4407
4408   if (__cuda___isnanf(a) || __cuda___isnanf(b)) {
4409     *quo = quot;
4410     return a + b;
4411   }
4412   orig_a = a;
4413
4414   sign = 0 - (__cuda___signbitf(a) != __cuda___signbitf(b));
4415   a = __cuda_fabsf(a);
4416   b = __cuda_fabsf(b);
4417   if (__cuda___isinff(a) || (b == 0.0f)) {
4418     *quo = quot;
4419     return __int_as_float(0x7fffffff);
4420   } else if (a >= b) {
4421
4422
4423     int expoa = (a < 1.175494351e-38f) ?
4424         ((int)__log2f(a)) : (((__float_as_int(a) >> 23) & 0xff) - 127);
4425     int expob = (b < 1.175494351e-38f) ?
4426         ((int)__log2f(b)) : (((__float_as_int(b) >> 23) & 0xff) - 127);
4427     int scale = expoa - expob;
4428     float scaled_b = __cuda_ldexpf(b, scale);
4429     if (scaled_b <= 0.5f * a) {
4430       scaled_b *= 2.0f;
4431     }
4432 # 2340 "/usr/local/cuda/bin/../include/math_functions.h" 3
4433     while (scaled_b >= b) {
4434       quot <<= 1;
4435       if (a >= scaled_b) {
4436         twoa = (2.0f * a - scaled_b) - scaled_b;
4437         a -= scaled_b;
4438         quot += 1;
4439       }
4440       scaled_b *= 0.5f;
4441     }
4442   }
4443
4444
4445   twoa = 2.0f * a;
4446   if ((twoa > b) || ((twoa == b) && (quot & 1))) {
4447     quot++;
4448     a -= b;
4449     a = __cuda_copysignf (a, -1.0f);
4450   }
4451 # 2375 "/usr/local/cuda/bin/../include/math_functions.h" 3
4452   a = __int_as_float((__float_as_int(orig_a) & 0x80000000)^
4453                      __float_as_int(a));
4454   quot = quot & (~((~0)<<3));
4455   quot = quot ^ sign;
4456   quot = quot - sign;
4457   *quo = quot;
4458   return a;
4459 }
4460
4461 extern __attribute__((weak)) float __cuda_fmaf(float a, float b, float c); float __cuda_fmaf(float a, float b, float c)
4462 {
4463   unsigned int xx, yy, zz, ww;
4464   unsigned int temp, s, u;
4465   unsigned int expo_x, expo_y, expo_z;
4466
4467   xx = __float_as_int(a);
4468   yy = __float_as_int(b);
4469   zz = __float_as_int(c);
4470 # 2401 "/usr/local/cuda/bin/../include/math_functions.h" 3
4471   temp = 0xff;
4472   expo_x = temp & (xx >> 23);
4473   expo_x = expo_x - 1;
4474   expo_y = temp & (yy >> 23);
4475   expo_y = expo_y - 1;
4476   expo_z = temp & (zz >> 23);
4477   expo_z = expo_z - 1;
4478
4479   if (!((expo_x <= 0xFD) &&
4480         (expo_y <= 0xFD) &&
4481         (expo_z <= 0xFD))) {
4482
4483
4484
4485
4486     if ((yy << 1) > 0xff000000) {
4487       return __int_as_float(0x7fffffff);
4488     }
4489     if ((zz << 1) > 0xff000000) {
4490       return __int_as_float(0x7fffffff);
4491     }
4492     if ((xx << 1) > 0xff000000) {
4493       return __int_as_float(0x7fffffff);
4494     }
4495 # 2436 "/usr/local/cuda/bin/../include/math_functions.h" 3
4496     if ((((xx << 1) == 0) && ((yy << 1) == 0xff000000)) ||
4497         (((yy << 1) == 0) && ((xx << 1) == 0xff000000))) {
4498       return __int_as_float(0x7fffffff);
4499     }
4500     if ((zz << 1) == 0xff000000) {
4501       if (((yy << 1) == 0xff000000) || ((xx << 1) == 0xff000000)) {
4502         if ((int)(xx ^ yy ^ zz) < 0) {
4503           return __int_as_float(0x7fffffff);
4504         }
4505       }
4506     }
4507
4508
4509
4510
4511     if ((xx << 1) == 0xff000000) {
4512       xx = xx ^ (yy & 0x80000000);
4513       return __int_as_float(xx);
4514     }
4515     if ((yy << 1) == 0xff000000) {
4516       yy = yy ^ (xx & 0x80000000);
4517       return __int_as_float(yy);
4518     }
4519     if ((zz << 1) == 0xff000000) {
4520       return __int_as_float(zz);
4521     }
4522
4523
4524
4525
4526
4527     if (zz == 0x80000000) {
4528       if (((xx << 1) == 0) || ((yy << 1) == 0)) {
4529         if ((int)(xx ^ yy) < 0) {
4530           return __int_as_float(zz);
4531         }
4532       }
4533     }
4534
4535
4536
4537     if (((zz << 1) == 0) &&
4538         (((xx << 1) == 0) || ((yy << 1) == 0))) {
4539       zz &= 0x7fffffff;
4540       return __int_as_float(zz);
4541     }
4542
4543
4544
4545     if (((xx << 1) == 0) || ((yy << 1) == 0)) {
4546       return __int_as_float(zz);
4547     }
4548
4549     if (expo_x == (unsigned int)-1) {
4550       temp = xx & 0x80000000;
4551       xx = xx << 8;
4552       while (!(xx & 0x80000000)) {
4553         xx <<= 1;
4554         expo_x--;
4555       }
4556       expo_x++;
4557       xx = (xx >> 8) | temp;
4558     }
4559
4560     if (expo_y == (unsigned int)-1) {
4561       temp = yy & 0x80000000;
4562       yy = yy << 8;
4563       while (!(yy & 0x80000000)) {
4564         yy <<= 1;
4565         expo_y--;
4566       }
4567       expo_y++;
4568       yy = (yy >> 8) | temp;
4569     }
4570
4571     if ((expo_z == (unsigned int)-1) && ((zz << 1) != 0)) {
4572       temp = zz & 0x80000000;
4573       zz = zz << 8;
4574       while (!(zz & 0x80000000)) {
4575         zz <<= 1;
4576         expo_z--;
4577       }
4578       expo_z++;
4579       zz = (zz >> 8) | temp;
4580     }
4581   }
4582
4583   expo_x = expo_x + expo_y;
4584   expo_y = xx ^ yy;
4585   xx = xx & 0x00ffffff;
4586   yy = yy << 8;
4587   xx = xx | 0x00800000;
4588   yy = yy | 0x80000000;
4589
4590   s = __umulhi(xx, yy);
4591   yy = xx * yy;
4592   xx = s;
4593   expo_x = expo_x - 127 + 2;
4594   expo_y = expo_y & 0x80000000;
4595
4596
4597   if (xx < 0x00800000) {
4598       xx = (xx << 1) | (yy >> 31);
4599       yy = (yy << 1);
4600       expo_x--;
4601   }
4602   temp = 0;
4603   if ((zz << 1) != 0) {
4604     s = zz & 0x80000000;
4605     zz &= 0x00ffffff;
4606     zz |= 0x00800000;
4607     ww = 0;
4608
4609     if ((int)expo_z > (int)expo_x) {
4610       temp = expo_z;
4611       expo_z = expo_x;
4612       expo_x = temp;
4613       temp = zz;
4614       zz = xx;
4615       xx = temp;
4616       temp = ww;
4617       ww = yy;
4618       yy = temp;
4619       temp = expo_y;
4620       expo_y = s;
4621       s = temp;
4622     }
4623
4624
4625     expo_z = expo_x - expo_z;
4626     u = expo_y ^ s;
4627     if (expo_z <= 49) {
4628
4629       temp = 0;
4630       while (expo_z >= 32) {
4631         temp = ww | (temp != 0);
4632         ww = zz;
4633         zz = 0;
4634         expo_z -= 32;
4635       }
4636       if (expo_z) {
4637         temp = ((temp >> expo_z) | (ww << (32 - expo_z)) |
4638                 ((temp << (32 - expo_z)) != 0));
4639         ww = (ww >> expo_z) | (zz << (32 - expo_z));
4640         zz = (zz >> expo_z);
4641       }
4642     } else {
4643       temp = 1;
4644       ww = 0;
4645       zz = 0;
4646     }
4647     if ((int)u < 0) {
4648
4649       temp = (unsigned)(-(int)temp);
4650       s = (temp != 0);
4651       u = yy - s;
4652       s = u > yy;
4653       yy = u - ww;
4654       s += yy > u;
4655       xx = (xx - zz) - s;
4656       if (!(xx | yy | temp)) {
4657
4658         return __int_as_float(xx);
4659       }
4660       if ((int)xx < 0) {
4661
4662
4663
4664         temp = ~temp;
4665         yy = ~yy;
4666         xx = ~xx;
4667         if (++temp == 0) {
4668           if (++yy == 0) {
4669             ++xx;
4670           }
4671         }
4672         expo_y ^= 0x80000000;
4673       }
4674
4675       while (!(xx & 0x00800000)) {
4676         xx = (xx << 1) | (yy >> 31);
4677         yy = (yy << 1);
4678         expo_x--;
4679       }
4680     } else {
4681
4682       yy = yy + ww;
4683       s = yy < ww;
4684       xx = xx + zz + s;
4685       if (xx & 0x01000000) {
4686         temp = temp | (yy << 31);
4687         yy = (yy >> 1) | (xx << 31);
4688         xx = ((xx & 0x80000000) | (xx >> 1)) & ~0x40000000;
4689         expo_x++;
4690       }
4691     }
4692   }
4693   temp = yy | (temp != 0);
4694   if (expo_x <= 0xFD) {
4695
4696     xx |= expo_y;
4697     s = xx & 1;
4698     xx += (temp == 0x80000000) ? s : (temp >> 31);
4699     xx = xx + (expo_x << 23);
4700     return __int_as_float(xx);
4701   } else if ((int)expo_x >= 126) {
4702
4703     xx = expo_y | 0x7f800000;
4704     return __int_as_float(xx);
4705   }
4706
4707   expo_x = (unsigned int)(-(int)expo_x);
4708   if (expo_x > 25) {
4709
4710     return __int_as_float(expo_y);
4711   }
4712   yy = (xx << (32 - expo_x)) | ((yy) ? 1 : 0);
4713   xx = expo_y + (xx >> expo_x);
4714   xx = xx + ((yy==0x80000000) ? (xx & 1) : (yy >> 31));
4715   xx |= expo_y;
4716
4717
4718
4719
4720   return __int_as_float(xx);
4721 }
4722
4723 static float __cudart_A1[32] =
4724 {
4725   1.0000000000e+000f,
4726   1.0218971968e+000f,
4727   1.0442737341e+000f,
4728   1.0671404600e+000f,
4729   1.0905077457e+000f,
4730   1.1143867970e+000f,
4731   1.1387885809e+000f,
4732   1.1637248993e+000f,
4733   1.1892070770e+000f,
4734   1.2152473927e+000f,
4735   1.2418577671e+000f,
4736   1.2690509558e+000f,
4737   1.2968395948e+000f,
4738   1.3252366781e+000f,
4739   1.3542555571e+000f,
4740   1.3839099407e+000f,
4741   1.4142135382e+000f,
4742   1.4451807737e+000f,
4743   1.4768261909e+000f,
4744   1.5091644526e+000f,
4745   1.5422108173e+000f,
4746   1.5759809017e+000f,
4747   1.6104903221e+000f,
4748   1.6457555294e+000f,
4749   1.6817928553e+000f,
4750   1.7186193466e+000f,
4751   1.7562521696e+000f,
4752   1.7947090864e+000f,
4753   1.8340080976e+000f,
4754   1.8741676807e+000f,
4755   1.9152065516e+000f,
4756   1.9571441412e+000f
4757 };
4758
4759 static float __cudart_A2[32] =
4760 {
4761   0.0000000000e+000f,
4762  -4.8115598617e-008f,
4763   4.8334701575e-008f,
4764  -5.9337519787e-008f,
4765  -1.3077539940e-008f,
4766  -5.4355400181e-008f,
4767   5.3862223126e-008f,
4768  -4.0514414934e-008f,
4769   3.7976352729e-008f,
4770  -3.2673948880e-008f,
4771   4.4968381019e-008f,
4772   1.4193333175e-009f,
4773  -4.0189995332e-008f,
4774  -3.4963733242e-008f,
4775  -1.0123349270e-008f,
4776  -5.8755773580e-008f,
4777   2.4203234972e-008f,
4778   3.3241999375e-008f,
4779  -4.5008988536e-008f,
4780  -2.4959373235e-008f,
4781   8.0709048333e-009f,
4782  -5.6610254262e-008f,
4783   9.8362171741e-009f,
4784  -5.1249720912e-008f,
4785  -2.4755326677e-008f,
4786  -4.8496175964e-008f,
4787  -9.2357703707e-009f,
4788  -1.1415044909e-008f,
4789  -1.1239277953e-008f,
4790  -4.6630056261e-008f,
4791   9.8453281083e-009f,
4792  -1.7021804410e-008f
4793 };
4794
4795 static float __cudart_Ainv[32] =
4796 {
4797   1.0000000000e+000f,
4798   9.7857207060e-001f,
4799   9.5760327578e-001f,
4800   9.3708384037e-001f,
4801   9.1700404882e-001f,
4802   8.9735454321e-001f,
4803   8.7812608480e-001f,
4804   8.5930967331e-001f,
4805   8.4089642763e-001f,
4806   8.2287776470e-001f,
4807   8.0524516106e-001f,
4808   7.8799045086e-001f,
4809   7.7110540867e-001f,
4810   7.5458222628e-001f,
4811   7.3841309547e-001f,
4812   7.2259038687e-001f,
4813   7.0710676908e-001f,
4814   6.9195497036e-001f,
4815   6.7712777853e-001f,
4816   6.6261833906e-001f,
4817   6.4841979742e-001f,
4818   6.3452547789e-001f,
4819   6.2092888355e-001f,
4820   6.0762369633e-001f,
4821   5.9460353851e-001f,
4822   5.8186244965e-001f,
4823   5.6939429045e-001f,
4824   5.5719339848e-001f,
4825   5.4525387287e-001f,
4826   5.3357023001e-001f,
4827   5.2213686705e-001f,
4828   5.1094859838e-001f
4829 };
4830
4831 extern __attribute__((weak)) float __internal_accurate_powf(float a, float b); float __internal_accurate_powf(float a, float b)
4832 {
4833   int i;
4834   float t;
4835   int expo;
4836   float log_hi, log_lo;
4837   float b_hi, b_lo;
4838   float prod_hi, prod_lo;
4839
4840   if ((a > 0.707106781f) && (a < 1.414213562f)) {
4841     float f, g, u, v, q;
4842
4843
4844
4845
4846
4847
4848     f = a - 1.0f;
4849     g = a + 1.0f;
4850     g = 1.0f / g;
4851     u = 2.0f * f * g;
4852     v = u * u;
4853     q = 1.49356810919559350E-001f/64.0f;
4854     q = q * v + 1.99887797540072460E-001f/16.0f;
4855     q = q * v + 3.33333880955515580E-001f/4.0f;
4856     q = q * v;
4857     q = q * u;
4858     log_hi = __int_as_float(__float_as_int(u) & 0xfffff000);
4859     v = __int_as_float(__float_as_int(f) & 0xfffff000);
4860     u = 2.0f * (f - log_hi);
4861     f = f - v;
4862     u = u - log_hi * v;
4863     u = u - log_hi * f;
4864     u = g * u;
4865     log_lo = q + u;
4866
4867
4868     b_hi = __int_as_float(__float_as_int(b) & 0xfffff000);
4869     b_lo = b - b_hi;
4870     prod_lo = b_lo * log_lo;
4871     prod_lo += b_lo * log_hi;
4872     prod_lo += b_hi * log_lo;
4873     prod_hi = b_hi * log_hi;
4874
4875
4876     return __cuda_expf(prod_hi) * __cuda_expf(prod_lo);
4877   }
4878
4879
4880   if (a >= 1.175494351e-38f) {
4881     i = __float_as_int(a);
4882     expo = ((i >> 23) & 0xff) - 127;
4883   } else {
4884     a *= 16777216.0f;
4885     i = __float_as_int(a);
4886     expo = ((i >> 23) & 0xff) - 127 - 24;
4887   }
4888   i = (i & 0x007fffff) | (0x3f800000);
4889   t = __int_as_float(i);
4890
4891   i = 0;
4892   if (t >= __cudart_A1[i+16]) i += 16;
4893   if (t >= __cudart_A1[i+8]) i += 8;
4894   if (t >= __cudart_A1[i+4]) i += 4;
4895   if (t >= __cudart_A1[i+2]) i += 2;
4896   if (t >= __cudart_A1[i+1]) i += 1;
4897
4898   t = t - __cudart_A1[i];
4899   t = t - __cudart_A2[i];
4900
4901   t = t * __cudart_Ainv[i];
4902
4903
4904   log_hi = (float)expo + (float)i * 0.03125f;
4905
4906   log_lo = - 3.42338934684934650E-001f;
4907   log_lo = log_lo * t + 4.80524913518140690E-001f;
4908   log_lo = log_lo * t - 7.21345070621603800E-001f;
4909   log_lo = log_lo * t + 1.44269503837073180E+000f;
4910   log_lo = log_lo * t;
4911
4912
4913   b_hi = __int_as_float(__float_as_int(b) & 0xfffff000);
4914   b_lo = b - b_hi;
4915   prod_lo = b_lo * log_lo;
4916   prod_lo = prod_lo + b_lo * log_hi;
4917   prod_lo = prod_lo + b_hi * log_lo;
4918   prod_hi = b_hi * log_hi;
4919
4920
4921   if (prod_hi >= 256.0f) {
4922     return __int_as_float(0x7f800000);
4923   }
4924   if (prod_hi <= -256.0f) {
4925     return 0.0f;
4926   }
4927
4928
4929   b = __cuda_exp2f (0.5f * prod_hi);
4930   t = __cuda_exp2f (prod_lo);
4931   t = t * b;
4932   t = t * b;
4933   return t;
4934 }
4935
4936 extern __attribute__((weak)) float __cuda_powif(float a, int b); float __cuda_powif(float a, int b)
4937 {
4938   unsigned int e = __cuda_abs(b);
4939   float r = 1.0f;
4940
4941   while (1) {
4942     if ((e & 1) != 0) {
4943       r = r * a;
4944     }
4945     e = e >> 1;
4946     if (e == 0) {
4947       return b < 0 ? 1.0f/r : r;
4948     }
4949     a = a * a;
4950   }
4951 }
4952
4953 extern __attribute__((weak)) double __cuda_powi(double a, int b); double __cuda_powi(double a, int b)
4954 {
4955   unsigned int e = __cuda_abs(b);
4956   double r = 1.0;
4957
4958   while (1) {
4959     if ((e & 1) != 0) {
4960       r = r * a;
4961     }
4962     e = e >> 1;
4963     if (e == 0) {
4964       return b < 0 ? 1.0/r : r;
4965     }
4966     a = a * a;
4967   }
4968 }
4969
4970 extern __attribute__((weak)) float __cuda_powf(float a, float b); float __cuda_powf(float a, float b)
4971 {
4972
4973
4974
4975   int bIsOddInteger;
4976   float t;
4977   if (a == 1.0f || b == 0.0f) {
4978     return 1.0f;
4979   }
4980   if (__cuda___isnanf(a) || __cuda___isnanf(b)) {
4981     return a + b;
4982   }
4983   if (a == __int_as_float(0x7f800000)) {
4984     return __cuda___signbitf(b) ? 0.0f : __int_as_float(0x7f800000);
4985   }
4986   if (__cuda___isinff(b)) {
4987     if (a == -1.0f) {
4988       return 1.0f;
4989     }
4990     t = (__cuda_fabsf(a) > 1.0f) ? __int_as_float(0x7f800000) : 0.0f;
4991     if (b < 0.0f) {
4992       t = 1.0f / t;
4993     }
4994     return t;
4995   }
4996   bIsOddInteger = (b - (2.0f * floorf(0.5f * b))) == 1.0f;
4997   if (a == 0.0f) {
4998     t = bIsOddInteger ? a : 0.0f;
4999     if (b < 0.0f) {
5000       t = 1.0f / t;
5001     }
5002     return t;
5003   }
5004   if (a == -__int_as_float(0x7f800000)) {
5005     t = (b < 0.0f) ? -1.0f/a : -a;
5006     if (bIsOddInteger) {
5007       t = __int_as_float(__float_as_int(t) ^ 0x80000000);
5008     }
5009     return t;
5010   }
5011   if ((a < 0.0f) && (b != __cuda_truncf(b))) {
5012     return __int_as_float(0x7fffffff);
5013   }
5014   t = __cuda_fabsf(a);
5015   t = __internal_accurate_powf(t, b);
5016   if ((a < 0.0f) && bIsOddInteger) {
5017     t = __int_as_float(__float_as_int(t) ^ 0x80000000);
5018   }
5019   return t;
5020
5021 }
5022
5023
5024 extern __attribute__((weak)) float __internal_tgammaf_kernel(float a); float __internal_tgammaf_kernel(float a)
5025 {
5026   float t;
5027   t = - 1.05767296987211380E-003f;
5028   t = t * a + 7.09279059435508670E-003f;
5029   t = t * a - 9.65347121958557050E-003f;
5030   t = t * a - 4.21736613253687960E-002f;
5031   t = t * a + 1.66542401247154280E-001f;
5032   t = t * a - 4.20043267827838460E-002f;
5033   t = t * a - 6.55878234051332940E-001f;
5034   t = t * a + 5.77215696929794240E-001f;
5035   t = t * a + 1.00000000000000000E+000f;
5036   return t;
5037 }
5038
5039
5040
5041
5042
5043 extern __attribute__((weak)) float __cuda_tgammaf(float a); float __cuda_tgammaf(float a)
5044 {
5045   float s, xx, x=a;
5046   if (x >= 0.0f) {
5047     if (x > 36.0f) x = 36.0f;
5048     s = 1.0f;
5049     xx = x;
5050     if (x > 34.03f) {
5051       xx -= 1.0f;
5052     }
5053     while (xx > 1.5f) {
5054       xx = xx - 1.0f;
5055       s = s * xx;
5056     }
5057     if (x >= 0.5f) {
5058       xx = xx - 1.0f;
5059     }
5060     xx = __internal_tgammaf_kernel(xx);
5061     if (x < 0.5f) {
5062       xx = xx * x;
5063     }
5064     s = s / xx;
5065     if (x > 34.03f) {
5066
5067       xx = x - 1.0f;
5068       s = s * xx;
5069     }
5070     return s;
5071   } else {
5072     if (x == __cuda_floorf(x)) {
5073       x = __int_as_float(0x7fffffff);
5074
5075       return x;
5076
5077     }
5078     if (x < -41.1f) x = -41.1f;
5079     xx = x;
5080     if (x < -34.03f) {
5081       xx += 6.0f;
5082     }
5083     s = xx;
5084     while (xx < -0.5f) {
5085       xx = xx + 1.0f;
5086       s = s * xx;
5087     }
5088     xx = __internal_tgammaf_kernel(xx);
5089     s = s * xx;
5090     s = 1.0f / s;
5091     if (x < -34.03f) {
5092       xx = x;
5093       xx *= (x + 1.0f);
5094       xx *= (x + 2.0f);
5095       xx *= (x + 3.0f);
5096       xx *= (x + 4.0f);
5097       xx *= (x + 5.0f);
5098       xx = 1.0f / xx;
5099       s = s * xx;
5100       if ((a < -42.0f) && !(((int)a)&1)) {
5101         s = __int_as_float(0x80000000);
5102       }
5103     }
5104     return s;
5105   }
5106 }
5107
5108 extern __attribute__((weak)) float __cuda_roundf(float a); float __cuda_roundf(float a)
5109 {
5110   float fa = __cuda_fabsf(a);
5111   if (fa > 8388608.0f) {
5112     return a;
5113   } else {
5114     float u = __cuda_floorf(fa + 0.5f);
5115     if (fa < 0.5f) u = 0.0f;
5116     return __cuda_copysignf(u, a);
5117   }
5118 }
5119
5120 extern __attribute__((weak)) long long int __internal_llroundf_kernel(float a); long long int __internal_llroundf_kernel(float a)
5121 {
5122   unsigned long long int res, t = 0LL;
5123   int shift;
5124   unsigned int ia = __float_as_int(a);
5125   if ((ia << 1) > 0xff000000) return 0LL;
5126   if ((int)ia >= 0x5f000000) return 0x7fffffffffffffffLL;
5127   if (ia >= 0xdf000000) return 0x8000000000000000LL;
5128   shift = 189 - ((ia >> 23) & 0xff);
5129   res = ((long long int)(((ia << 8) | 0x80000000) >> 1)) << 32;
5130   if (shift >= 64) {
5131     t = res;
5132     res = 0;
5133   } else if (shift) {
5134     t = res << (64 - shift);
5135     res = res >> shift;
5136   }
5137   if (t >= 0x8000000000000000LL) {
5138       res++;
5139   }
5140   if ((int)ia < 0) res = (unsigned long long int)(-(long long int)res);
5141   return (long long int)res;
5142 }
5143
5144 extern __attribute__((weak)) long long int __cuda_llroundf(float a); long long int __cuda_llroundf(float a)
5145 {
5146   return __internal_llroundf_kernel(a);
5147 }
5148
5149 extern __attribute__((weak)) long int __cuda_lroundf(float a); long int __cuda_lroundf(float a)
5150 {
5151
5152   return (long int)__cuda_llroundf(a);
5153 # 3101 "/usr/local/cuda/bin/../include/math_functions.h" 3
5154 }
5155
5156 extern __attribute__((weak)) float __cuda_fdimf(float a, float b); float __cuda_fdimf(float a, float b)
5157 {
5158   float t;
5159   t = a - b;
5160   if (a <= b) {
5161     t = 0.0f;
5162   }
5163   return t;
5164 }
5165
5166 extern __attribute__((weak)) int __cuda_ilogbf(float a); int __cuda_ilogbf(float a)
5167 {
5168   unsigned int i;
5169   int expo;
5170   a = __cuda_fabsf(a);
5171   if (a <= 1.175494351e-38f) {
5172
5173     if (a == 0.0f) {
5174       expo = -((int)((unsigned int)-1 >> 1))-1;
5175     } else {
5176       expo = -126;
5177       i = __float_as_int(a);
5178       i = i << 8;
5179       while ((int)i >= 0) {
5180         expo--;
5181         i = i + i;
5182       }
5183     }
5184   } else {
5185     i = __float_as_int(a);
5186     expo = ((int)((i >> 23) & 0xff)) - 127;
5187     if ((i == 0x7f800000)) {
5188       expo = ((int)((unsigned int)-1 >> 1));
5189     }
5190     if ((i > 0x7f800000)) {
5191       expo = -((int)((unsigned int)-1 >> 1))-1;
5192     }
5193   }
5194   return expo;
5195 }
5196
5197 extern __attribute__((weak)) float __cuda_logbf(float a); float __cuda_logbf(float a)
5198 {
5199   unsigned int i;
5200   int expo;
5201   float res;
5202
5203   if (__cuda___isnanf(a)) return a + a;
5204
5205   a = __cuda_fabsf(a);
5206   if (a <= 1.175494351e-38f) {
5207
5208     if (a == 0.0f) {
5209       res = -__int_as_float(0x7f800000);
5210     } else {
5211       expo = -126;
5212       i = __float_as_int(a);
5213       i = i << 8;
5214       while ((int)i >= 0) {
5215         expo--;
5216         i = i + i;
5217       }
5218       res = (float)expo;
5219     }
5220   } else {
5221     i = __float_as_int(a);
5222     expo = ((int)((i >> 23) & 0xff)) - 127;
5223     res = (float)expo;
5224     if ((i >= 0x7f800000)) {
5225
5226       res = a + a;
5227     }
5228   }
5229   return res;
5230 }
5231
5232 extern __attribute__((weak)) void __cuda_sincosf(float a, float *sptr, float *cptr); void __cuda_sincosf(float a, float *sptr, float *cptr)
5233 {
5234
5235
5236
5237   float t, u, s, c;
5238   int quadrant;
5239   if (__cuda___isinff(a)) {
5240     *sptr = __int_as_float(0x7fffffff);
5241     *cptr = __int_as_float(0x7fffffff);
5242     return;
5243   }
5244   if (a == 0.0f) {
5245     *sptr = a;
5246     *cptr = 1.0f;
5247     return;
5248   }
5249   t = __internal_trig_reduction_kernel(a, &quadrant);
5250   u = __internal_cos_kernel(t);
5251   t = __internal_sin_kernel(t);
5252   if (quadrant & 1) {
5253     s = u;
5254     c = t;
5255   } else {
5256     s = t;
5257     c = u;
5258   }
5259   if (quadrant & 2) {
5260     s = -s;
5261   }
5262   quadrant++;
5263   if (quadrant & 2) {
5264     c = -c;
5265   }
5266   *sptr = s;
5267   *cptr = c;
5268
5269 }
5270 # 3226 "/usr/local/cuda/bin/../include/math_functions.h" 3
5271 extern __attribute__((weak)) double rsqrt(double a); double rsqrt(double a)
5272 {
5273   return 1.0 / sqrt(a);
5274 }
5275
5276 extern __attribute__((weak)) float rsqrtf(float a); float rsqrtf(float a)
5277 {
5278   return (float)rsqrt((double)a);
5279 }
5280 # 4167 "/usr/local/cuda/bin/../include/math_functions.h" 3
5281 # 1 "/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h" 1 3
5282 # 45 "/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h" 3
5283 extern __attribute__((weak)) double __cuda_fabs(double a); double __cuda_fabs(double a)
5284 {
5285   return (float)__cuda_fabsf((float)a);
5286 }
5287
5288 extern __attribute__((weak)) double __cuda_fmax(double a, double b); double __cuda_fmax(double a, double b)
5289 {
5290   return (float)__cuda_fmaxf((float)a, (float)b);
5291 }
5292
5293 extern __attribute__((weak)) double __cuda_fmin(double a, double b); double __cuda_fmin(double a, double b)
5294 {
5295   return (float)__cuda_fminf((float)a, (float)b);
5296 }
5297
5298 extern __attribute__((weak)) int __cuda___finite(double a); int __cuda___finite(double a)
5299 {
5300   return __cuda___finitef((float)a);
5301 }
5302
5303 extern __attribute__((weak)) int __cuda___isinf(double a); int __cuda___isinf(double a)
5304 {
5305   return __cuda___isinff((float)a);
5306 }
5307
5308 extern __attribute__((weak)) int __cuda___isnan(double a); int __cuda___isnan(double a)
5309 {
5310   return __cuda___isnanf((float)a);
5311 }
5312
5313 extern __attribute__((weak)) int __cuda___signbit(double a); int __cuda___signbit(double a)
5314 {
5315   return __cuda___signbitf((float)a);
5316 }
5317
5318 extern __attribute__((weak)) double __cuda_sqrt(double a); double __cuda_sqrt(double a)
5319 {
5320   return (double)__cuda_sqrtf((float)a);
5321 }
5322
5323 extern __attribute__((weak)) double __cuda_rsqrt(double a); double __cuda_rsqrt(double a)
5324 {
5325   return (double)__cuda_rsqrtf((float)a);
5326 }
5327
5328 extern __attribute__((weak)) double __cuda_ceil(double a); double __cuda_ceil(double a)
5329 {
5330   return (double)__cuda_ceilf((float)a);
5331 }
5332
5333 extern __attribute__((weak)) double __cuda_trunc(double a); double __cuda_trunc(double a)
5334 {
5335   return (double)__cuda_truncf((float)a);
5336 }
5337
5338 extern __attribute__((weak)) double __cuda_floor(double a); double __cuda_floor(double a)
5339 {
5340   return (double)__cuda_floorf((float)a);
5341 }
5342
5343 extern __attribute__((weak)) double __cuda_copysign(double a, double b); double __cuda_copysign(double a, double b)
5344 {
5345   return (double)__cuda_copysignf((float)a, (float)b);
5346 }
5347
5348 extern __attribute__((weak)) double __cuda_sin(double a); double __cuda_sin(double a)
5349 {
5350   return (double)__cuda_sinf((float)a);
5351 }
5352
5353 extern __attribute__((weak)) double __cuda_cos(double a); double __cuda_cos(double a)
5354 {
5355   return (double)__cuda_cosf((float)a);
5356 }
5357
5358 extern __attribute__((weak)) void __cuda_sincos(double a, double *sptr, double *cptr); void __cuda_sincos(double a, double *sptr, double *cptr)
5359 {
5360   float fs, fc;
5361
5362   __cuda_sincosf((float)a, &fs, &fc);
5363
5364   *sptr = (double)fs;
5365   *cptr = (double)fc;
5366 }
5367
5368 extern __attribute__((weak)) double __cuda_tan(double a); double __cuda_tan(double a)
5369 {
5370   return (double)__cuda_tanf((float)a);
5371 }
5372
5373 extern __attribute__((weak)) double __cuda_exp(double a); double __cuda_exp(double a)
5374 {
5375   return (double)__cuda_expf((float)a);
5376 }
5377
5378 extern __attribute__((weak)) double __cuda_exp2(double a); double __cuda_exp2(double a)
5379 {
5380   return (double)__cuda_exp2f((float)a);
5381 }
5382
5383 extern __attribute__((weak)) double __cuda_exp10(double a); double __cuda_exp10(double a)
5384 {
5385   return (double)__cuda_exp10f((float)a);
5386 }
5387
5388 extern __attribute__((weak)) double __cuda_expm1(double a); double __cuda_expm1(double a)
5389 {
5390   return (double)__cuda_expm1f((float)a);
5391 }
5392
5393 extern __attribute__((weak)) double __cuda_cosh(double a); double __cuda_cosh(double a)
5394 {
5395   return (double)__cuda_coshf((float)a);
5396 }
5397
5398 extern __attribute__((weak)) double __cuda_sinh(double a); double __cuda_sinh(double a)
5399 {
5400   return (double)__cuda_sinhf((float)a);
5401 }
5402
5403 extern __attribute__((weak)) double __cuda_tanh(double a); double __cuda_tanh(double a)
5404 {
5405   return (double)__cuda_tanhf((float)a);
5406 }
5407
5408 extern __attribute__((weak)) double __cuda_asin(double a); double __cuda_asin(double a)
5409 {
5410   return (double)__cuda_asinf((float)a);
5411 }
5412
5413 extern __attribute__((weak)) double __cuda_acos(double a); double __cuda_acos(double a)
5414 {
5415   return (double)__cuda_acosf((float)a);
5416 }
5417
5418 extern __attribute__((weak)) double __cuda_atan(double a); double __cuda_atan(double a)
5419 {
5420   return (double)__cuda_atanf((float)a);
5421 }
5422
5423 extern __attribute__((weak)) double __cuda_atan2(double a, double b); double __cuda_atan2(double a, double b)
5424 {
5425   return (double)__cuda_atan2f((float)a, (float)b);
5426 }
5427
5428 extern __attribute__((weak)) double __cuda_log(double a); double __cuda_log(double a)
5429 {
5430   return (double)__cuda_logf((float)a);
5431 }
5432
5433 extern __attribute__((weak)) double __cuda_log2(double a); double __cuda_log2(double a)
5434 {
5435   return (double)__cuda_log2f((float)a);
5436 }
5437
5438 extern __attribute__((weak)) double __cuda_log10(double a); double __cuda_log10(double a)
5439 {
5440   return (double)__cuda_log10f((float)a);
5441 }
5442
5443 extern __attribute__((weak)) double __cuda_log1p(double a); double __cuda_log1p(double a)
5444 {
5445   return (double)__cuda_log1pf((float)a);
5446 }
5447
5448 extern __attribute__((weak)) double __cuda_acosh(double a); double __cuda_acosh(double a)
5449 {
5450   return (double)__cuda_acoshf((float)a);
5451 }
5452
5453 extern __attribute__((weak)) double __cuda_asinh(double a); double __cuda_asinh(double a)
5454 {
5455   return (double)__cuda_asinhf((float)a);
5456 }
5457
5458 extern __attribute__((weak)) double __cuda_atanh(double a); double __cuda_atanh(double a)
5459 {
5460   return (double)__cuda_atanhf((float)a);
5461 }
5462
5463 extern __attribute__((weak)) double __cuda_hypot(double a, double b); double __cuda_hypot(double a, double b)
5464 {
5465   return (double)__cuda_hypotf((float)a, (float)b);
5466 }
5467
5468 extern __attribute__((weak)) double __cuda_cbrt(double a); double __cuda_cbrt(double a)
5469 {
5470   return (double)__cuda_cbrtf((float)a);
5471 }
5472
5473 extern __attribute__((weak)) double __cuda_erf(double a); double __cuda_erf(double a)
5474 {
5475   return (double)__cuda_erff((float)a);
5476 }
5477
5478 extern __attribute__((weak)) double __cuda_erfc(double a); double __cuda_erfc(double a)
5479 {
5480   return (double)__cuda_erfcf((float)a);
5481 }
5482
5483 extern __attribute__((weak)) double __cuda_lgamma(double a); double __cuda_lgamma(double a)
5484 {
5485   return (double)__cuda_lgammaf((float)a);
5486 }
5487
5488 extern __attribute__((weak)) double __cuda_tgamma(double a); double __cuda_tgamma(double a)
5489 {
5490   return (double)__cuda_tgammaf((float)a);
5491 }
5492
5493 extern __attribute__((weak)) double __cuda_ldexp(double a, int b); double __cuda_ldexp(double a, int b)
5494 {
5495   return (double)__cuda_ldexpf((float)a, b);
5496 }
5497
5498 extern __attribute__((weak)) double __cuda_scalbn(double a, int b); double __cuda_scalbn(double a, int b)
5499 {
5500   return (double)__cuda_scalbnf((float)a, b);
5501 }
5502
5503 extern __attribute__((weak)) double __cuda_scalbln(double a, long b); double __cuda_scalbln(double a, long b)
5504 {
5505   return (double)__cuda_scalblnf((float)a, b);
5506 }
5507
5508 extern __attribute__((weak)) double __cuda_frexp(double a, int *b); double __cuda_frexp(double a, int *b)
5509 {
5510   return (double)__cuda_frexpf((float)a, b);
5511 }
5512
5513 extern __attribute__((weak)) double __cuda_modf(double a, double *b); double __cuda_modf(double a, double *b)
5514 {
5515   float fb;
5516   float fa = __cuda_modff((float)a, &fb);
5517
5518   *b = (double)fb;
5519
5520   return (double)fa;
5521 }
5522
5523 extern __attribute__((weak)) double __cuda_fmod(double a, double b); double __cuda_fmod(double a, double b)
5524 {
5525   return (double)__cuda_fmodf((float)a, (float)b);
5526 }
5527
5528 extern __attribute__((weak)) double __cuda_remainder(double a, double b); double __cuda_remainder(double a, double b)
5529 {
5530   return (double)__cuda_remainderf((float)a, (float)b);
5531 }
5532
5533 extern __attribute__((weak)) double __cuda_remquo(double a, double b, int *c); double __cuda_remquo(double a, double b, int *c)
5534 {
5535   return (double)__cuda_remquof((float)a, (float)b, c);
5536 }
5537
5538 extern __attribute__((weak)) double __cuda_nextafter(double a, double b); double __cuda_nextafter(double a, double b)
5539 {
5540   return (double)__cuda_nextafterf((float)a, (float)b);
5541 }
5542
5543 extern __attribute__((weak)) double __cuda_nan(const char *tagp); double __cuda_nan(const char *tagp)
5544 {
5545   return (double)__cuda_nanf(tagp);
5546 }
5547
5548 extern __attribute__((weak)) double __cuda_pow(double a, double b); double __cuda_pow(double a, double b)
5549 {
5550   return (double)__cuda_powf((float)a, (float)b);
5551 }
5552
5553 extern __attribute__((weak)) double __cuda_round(double a); double __cuda_round(double a)
5554 {
5555   return (double)__cuda_roundf((float)a);
5556 }
5557
5558 extern __attribute__((weak)) long __cuda_lround(double a); long __cuda_lround(double a)
5559 {
5560   return __cuda_lroundf((float)a);
5561 }
5562
5563 extern __attribute__((weak)) long long __cuda_llround(double a); long long __cuda_llround(double a)
5564 {
5565   return __cuda_llroundf((float)a);
5566 }
5567
5568 extern __attribute__((weak)) double __cuda_rint(double a); double __cuda_rint(double a)
5569 {
5570   return (double)__cuda_rintf((float)a);
5571 }
5572
5573 extern __attribute__((weak)) long __cuda_lrint(double a); long __cuda_lrint(double a)
5574 {
5575   return __cuda_lrintf((float)a);
5576 }
5577
5578 extern __attribute__((weak)) long long __cuda_llrint(double a); long long __cuda_llrint(double a)
5579 {
5580   return __cuda_llrintf((float)a);
5581 }
5582
5583 extern __attribute__((weak)) double __cuda_nearbyint(double a); double __cuda_nearbyint(double a)
5584 {
5585   return (double)__cuda_nearbyintf((float)a);
5586 }
5587
5588 extern __attribute__((weak)) double __cuda_fdim(double a, double b); double __cuda_fdim(double a, double b)
5589 {
5590   return (double)__cuda_fdimf((float)a, (float)b);
5591 }
5592
5593 extern __attribute__((weak)) int __cuda_ilogb(double a); int __cuda_ilogb(double a)
5594 {
5595   return __cuda_ilogbf((float)a);
5596 }
5597
5598 extern __attribute__((weak)) double __cuda_logb(double a); double __cuda_logb(double a)
5599 {
5600   return (double)__cuda_logbf((float)a);
5601 }
5602
5603 extern __attribute__((weak)) double __cuda_fma(double a, double b, double c); double __cuda_fma(double a, double b, double c)
5604 {
5605   return (double)__cuda_fmaf((float)a, (float)b, (float)c);
5606 }
5607 # 4168 "/usr/local/cuda/bin/../include/math_functions.h" 2 3
5608 # 89 "/usr/local/cuda/bin/../include/common_functions.h" 2
5609 # 196 "/usr/local/cuda/bin/../include/crt/host_runtime.h" 2
5610 # 6 "/tmp/tmpxft_00001ecc_00000000-0.stub.c" 2
5611 extern void __sti____cudaRegisterAll_29_tmpxft_00001ecc_00000000_2_ii_91788a12(void) __attribute__((__constructor__));
5612 void __sti____cudaRegisterAll_29_tmpxft_00001ecc_00000000_2_ii_91788a12(void){__cudaFatCubinHandle = __cudaRegisterFatBinary((void*)(&__fatDeviceText));}
5613 # 475 "y.cu" 2