COMPASS  5.0.0
End-to-end AO simulation tool using GPU acceleration
carma_utils.cuh
1 // -----------------------------------------------------------------------------
2 // This file is part of COMPASS <https://anr-compass.github.io/compass/>
3 //
4 // Copyright (C) 2012-2019 COMPASS Team <https://github.com/ANR-COMPASS>
5 // All rights reserved.
6 // Distributed under GNU - LGPL
7 //
8 // COMPASS is free software: you can redistribute it and/or modify it under the terms of the GNU Lesser
9 // General Public License as published by the Free Software Foundation, either version 3 of the License,
10 // or any later version.
11 //
12 // COMPASS: End-to-end AO simulation tool using GPU acceleration
13 // The COMPASS platform was designed to meet the need of high-performance for the simulation of AO systems.
14 //
15 // The final product includes a software package for simulating all the critical subcomponents of AO,
16 // particularly in the context of the ELT and a real-time core based on several control approaches,
17 // with performances consistent with its integration into an instrument. Taking advantage of the specific
18 // hardware architecture of the GPU, the COMPASS tool allows to achieve adequate execution speeds to
19 // conduct large simulation campaigns called to the ELT.
20 //
21 // The COMPASS platform can be used to carry a wide variety of simulations to both testspecific components
22 // of AO of the E-ELT (such as wavefront analysis device with a pyramid or elongated Laser star), and
23 // various systems configurations such as multi-conjugate AO.
24 //
25 // COMPASS is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the
26 // implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
27 // See the GNU Lesser General Public License for more details.
28 //
29 // You should have received a copy of the GNU Lesser General Public License along with COMPASS.
30 // If not, see <https://www.gnu.org/licenses/lgpl-3.0.txt>.
31 // -----------------------------------------------------------------------------
32 
33 //! \file carma_utils.cuh
34 //! \ingroup libcarma
35 //! \brief this file provides tools to CarmaObj
36 //! \author COMPASS Team <https://github.com/ANR-COMPASS>
37 //! \version 5.0.0
38 //! \date 2011/01/28
39 //! \copyright GNU Lesser General Public License
40 
41 #ifndef _CARMA_UTILS_CUH_
42 #define _CARMA_UTILS_CUH_
43 
44 #include <cufft.h>
45 #include <stdio.h>
46 #include <stdlib.h>
47 #include <string.h>
48 
49 #include <driver_types.h>
50 #include <vector_types.h>
51 
52 #include <cub/cub.cuh>
53 #include <cuda.h>
54 #include <cuda_runtime_api.h>
55 
56 template <class T> struct SharedMemory {
57  __device__ inline operator T *() {
58  extern __shared__ int __smem[];
59  return (T *)__smem;
60  }
61 
62  __device__ inline operator const T *() const {
63  extern __shared__ int __smem[];
64  return (T *)__smem;
65  }
66 };
67 
68 // specialize for double to avoid unaligned memory
69 // access compile errors
70 template <> struct SharedMemory<double> {
71  __device__ inline operator double *() {
72  extern __shared__ double __smem_d[];
73  return (double *)__smem_d;
74  }
75 
76  __device__ inline operator const double *() const {
77  extern __shared__ double __smem_d[];
78  return (double *)__smem_d;
79  }
80 };
81 template <> struct SharedMemory<float> {
82  __device__ inline operator float *() {
83  extern __shared__ float __smem_f[];
84  return (float *)__smem_f;
85  }
86 
87  __device__ inline operator const float *() const {
88  extern __shared__ float __smem_f[];
89  return (float *)__smem_f;
90  }
91 };
92 
93 template <class T> __device__ inline void mswap(T &a, T &b) {
94  T tmp = a;
95  a = b;
96  b = tmp;
97 }
98 
99 template <class T>
100 __inline__ __device__ void reduce_krnl(T *sdata, int size, int n) {
101  if (size & (size - 1)) { // test if size is not a power of 2
102  unsigned int s;
103  if ((size & 1) != 0)
104  s = size / 2 + 1; //(size&1)==size%2
105  else
106  s = size / 2;
107  unsigned int s_old = size;
108  while (s > 0) {
109  if ((n < s) && (n + s < s_old)) {
110  sdata[n] += sdata[n + s];
111  }
112  //__threadfence_block();
113  //__threadfence();
114  __syncthreads();
115  s_old = s;
116  s /= 2;
117  if ((2 * s < s_old) && (s != 0))
118  s += 1;
119  }
120  } else {
121  // do reduction in shared mem
122  for (unsigned int s = size / 2; s > 0; s >>= 1) {
123  if (n < s) {
124  sdata[n] += sdata[n + s];
125  }
126  //__threadfence_block();
127  //__threadfence();
128  __syncthreads();
129  }
130  }
131 }
132 
133 template <class T_data>
134 __inline__ __device__ T_data carma_clip(T_data n, T_data min, T_data max) {
135  return n > max ? max : (n < min ? min : n);
136 }
137 
138 #endif //_CARMA_UTILS_CUH_