1 // -----------------------------------------------------------------------------
2 // This file is part of COMPASS <https://anr-compass.github.io/compass/>
4 // Copyright (C) 2012-2019 COMPASS Team <https://github.com/ANR-COMPASS>
5 // All rights reserved.
6 // Distributed under GNU - LGPL
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.
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.
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.
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.
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.
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 // -----------------------------------------------------------------------------
33 //! \file carma_utils.cuh
35 //! \brief this file provides tools to CarmaObj
36 //! \author COMPASS Team <https://github.com/ANR-COMPASS>
39 //! \copyright GNU Lesser General Public License
41 #ifndef _CARMA_UTILS_CUH_
42 #define _CARMA_UTILS_CUH_
49 #include <driver_types.h>
50 #include <vector_types.h>
52 #include <cub/cub.cuh>
54 #include <cuda_runtime_api.h>
56 template <class T> struct SharedMemory {
57 __device__ inline operator T *() {
58 extern __shared__ int __smem[];
62 __device__ inline operator const T *() const {
63 extern __shared__ int __smem[];
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;
76 __device__ inline operator const double *() const {
77 extern __shared__ double __smem_d[];
78 return (double *)__smem_d;
81 template <> struct SharedMemory<float> {
82 __device__ inline operator float *() {
83 extern __shared__ float __smem_f[];
84 return (float *)__smem_f;
87 __device__ inline operator const float *() const {
88 extern __shared__ float __smem_f[];
89 return (float *)__smem_f;
93 template <class T> __device__ inline void mswap(T &a, T &b) {
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
104 s = size / 2 + 1; //(size&1)==size%2
107 unsigned int s_old = size;
109 if ((n < s) && (n + s < s_old)) {
110 sdata[n] += sdata[n + s];
112 //__threadfence_block();
117 if ((2 * s < s_old) && (s != 0))
121 // do reduction in shared mem
122 for (unsigned int s = size / 2; s > 0; s >>= 1) {
124 sdata[n] += sdata[n + s];
126 //__threadfence_block();
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);
138 #endif //_CARMA_UTILS_CUH_