kspaceFirstOrder3D-CUDA  1.1
The CUDA/C++ implementation of the k-wave toolbox for the time-domain simulation of acoustic wave fields in 3D
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
CUDAUtils.cuh
Go to the documentation of this file.
1 /**
2  * @file CUDAUtils.cuh
3  *
4  * @author Jiri Jaros \n
5  * Faculty of Information Technology \n
6  * Brno University of Technology \n
7  * jarosjir@fit.vutbr.cz
8  *
9  * @brief The header file with CUDA utility functions. These routines are to be inlined.
10  *
11  * @version kspaceFirstOrder3D 3.4
12  *
13  * @date 22 March 2016, 15:25 (created) \n
14  * 25 July 2016, 10:56 (revised)
15  *
16  * @section License
17  * This file is part of the C++ extension of the k-Wave Toolbox
18  * (http://www.k-wave.org).\n Copyright (C) 2016 Jiri Jaros and Bradley Treeby.
19  *
20  * This file is part of the k-Wave. k-Wave is free software: you can redistribute it and/or modify
21  * it under the terms of the GNU Lesser General Public License as published by the Free Software
22  * Foundation, either version 3 of the License, or (at your option) any later version.
23  *
24  * k-Wave is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even
25  * the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser
26  * General Public License for more details.
27  *
28  * You should have received a copy of the GNU Lesser General Public License along with k-Wave.
29  * If not, see http://www.gnu.org/licenses/.
30  */
31 
32 
33 #ifndef CUDA_UTILS_CUH
34 #define CUDA_UTILS_CUH
35 
37 
38 //------------------------------------------------------------------------------------------------//
39 //------------------------------------------ Variables -------------------------------------------//
40 //------------------------------------------------------------------------------------------------//
41 /**
42  * This variable holds necessary simulation constants in the CUDA GPU memory. The variable is
43  * defined in CUDADeviceConstants.cu
44  */
45 extern __constant__ TCUDADeviceConstants cudaDeviceConstants;
46 
47 //------------------------------------------------------------------------------------------------//
48 //--------------------------------------- Index routines -----------------------------------------//
49 //------------------------------------------------------------------------------------------------//
50 
51 /**
52  * @brief Get global 1D coordinate for 1D CUDA block.
53  * @details Get global 1D coordinate for 1D CUDA block.
54  *
55  * @return x-coordinate for 1D CUDA block
56  */
57 inline __device__ unsigned int GetIndex()
58 {
59  return threadIdx.x + blockIdx.x * blockDim.x;
60 }// end of GetIndex()
61 //--------------------------------------------------------------------------------------------------
62 
63 /**
64  * @brief Get x-stride for 3D CUDA block (for processing multiple grid points by a single thread).
65  * @details Get x-stride for 3D CUDA block (for processing multiple grid points by a single thread).
66  *
67  * @return x stride for 3D CUDA block
68  */
69 inline __device__ unsigned int GetStride()
70 {
71  return blockDim.x * gridDim.x;
72 }// end of GetStride
73 //--------------------------------------------------------------------------------------------------
74 
75 /**
76  * @brief Get 3D coordinates for a real matrix form a 1D index.
77  * @details Get 3D coordinates for a real matrix form a 1D index.
78  *
79  * @param [in] i - index
80  * @return 3D coordinates
81  */
82 inline __device__ dim3 GetReal3DCoords(const unsigned int i)
83 {
84  return dim3( i % cudaDeviceConstants.nx,
87 }// end of GetReal3DCoords
88 //-------------------------------------------------------------------------------------------------
89 
90 /**
91  * @brief Get a 3D coordinates for a complex matrix form a 1D index.
92  * @details Get a 3D coordinates for a complex matrix form a 1D index.
93  *
94  * @param [in] i - index
95  * @return 3D coordinates
96  */
97 inline __device__ dim3 GetComplex3DCoords(const unsigned int i)
98 {
99  return dim3( i % cudaDeviceConstants.nxComplex,
102 }// end of GetComplex3DCoords
103 //--------------------------------------------------------------------------------------------------
104 
105 
106 //------------------------------------------------------------------------------------------------//
107 //----------------------------- Multiplication operators for float2 ------------------------------//
108 //------------------------------------------------------------------------------------------------//
109 
110 /**
111  * @brief Operator * for float2 datatype (per element multiplication).
112  * @details Operator * for float2 datatype (per element multiplication).
113  *
114  * @param [in] a
115  * @param [in] b
116  * @return a.x * b.x, a.y * b.y
117  */
118 inline __device__ float2 operator*(const float2 a,
119  const float2 b)
120 {
121  return make_float2(a.x * b.x, a.y * b.y);
122 }// end of operator*
123 //---------------------------------------------------------------------------------------------------
124 
125 /**
126  * @brief Operator * for float2 datatype (per element multiplication).
127  * @details Operator * for float2 datatype (per element multiplication).
128  *
129  * @param [in] a
130  * @param [in] b
131  * @return a.x * b, a.y * b
132  */
133 inline __device__ float2 operator*(const float2 a,
134  const float b)
135 {
136  return make_float2(a.x * b, a.y * b);
137 }// end of operator*
138 //--------------------------------------------------------------------------------------------------
139 
140 /**
141  * @brief Operator * for float2 datatype (per element multiplication).
142  * @details Operator * for float2 datatype (per element multiplication).
143  *
144  * @param [in] a
145  * @param [in] b
146  * @return a * b.x, a * b.y
147  */
148 inline __device__ float2 operator*(const float b,
149  const float2 a)
150 {
151  return make_float2(b * a.x, b * a.y);
152 }// end of operator*
153 //--------------------------------------------------------------------------------------------------
154 
155 
156 /**
157  * @brief Operator *= for float2 datatype (per element multiplication).
158  * @details Operator *= for float2 datatype (per element multiplication).
159  *
160  * @param [in,out] a
161  * @param [in] b
162  * @return a.x *= b.x, a.y *= b.y
163  */
164 inline __device__ void operator*=(float2& a,
165  const float2 b)
166 {
167  a.x *= b.x;
168  a.y *= b.y;
169 }// end of operator*=
170 //--------------------------------------------------------------------------------------------------
171 
172 /**
173  * @brief Operator *= for float2 datatype (per element multiplication).
174  * @details Operator *= for float2 datatype (per element multiplication).
175  *
176  * @param [in,out] a
177  * @param [in] b
178  * @return a.x =* b, a.y =* b
179  */
180 inline __device__ void operator*=(float2& a,
181  const float b)
182 {
183  a.x *= b;
184  a.y *= b;
185 }// end of operator*=
186 //--------------------------------------------------------------------------------------------------
187 
188 //------------------------------------------------------------------------------------------------//
189 //-------------------------------- Addition operators for float2 ---------------------------------//
190 //------------------------------------------------------------------------------------------------//
191 
192 /**
193  * @brief Operator + for float2 datatype (per element multiplication).
194  * @details Operator + for float2 datatype (per element multiplication).
195  * @param [in] a
196  * @param [in] b
197  * @return a.x + b.x, a.y + b.y
198  */
199 inline __device__ float2 operator+(const float2 a,
200  const float2 b)
201 {
202  return make_float2(a.x + b.x, a.y + b.y);
203 }// end of operator+
204 //--------------------------------------------------------------------------------------------------
205 
206 /**
207  * @brief Operator + for float2 datatype (per element multiplication)
208  * @details Operator + for float2 datatype (per element multiplication).
209  *
210  * @param [in] a
211  * @param [in] b
212  * @return a.x + b, a.y + b
213  */
214 inline __device__ float2 operator+(const float2 a,
215  const float b)
216 {
217  return make_float2(a.x + b, a.y + b);
218 }// end of operator+
219 //--------------------------------------------------------------------------------------------------
220 
221 /**
222  * @brief Operator + for float2 datatype (per element multiplication).
223  * @details Operator + for float2 datatype (per element multiplication).
224  *
225  * @param [in] a
226  * @param [in] b
227  * @return a + b.x, a + b.y
228  */
229 inline __device__ float2 operator+(const float b,
230  const float2 a)
231 {
232  return make_float2(b + a.x, b + a.y);
233 }// end of operator+
234 //--------------------------------------------------------------------------------------------------
235 
236 
237 /**
238  * @brief Operator += for float2 datatype (per element multiplication).
239  * @details Operator += for float2 datatype (per element multiplication)
240  *
241  * @param [in,out] a
242  * @param [in] b
243  * @return a.x += b.x, a.y += b.y
244  */
245 inline __device__ void operator+=(float2& a,
246  const float2 b)
247 {
248  a.x += b.x;
249  a.y += b.y;
250 }// end of operator+=
251 //--------------------------------------------------------------------------------------------------
252 
253 /**
254  * @brief Operator += for float2 datatype (per element multiplication).
255  * @details Operator += for float2 datatype (per element multiplication).
256  * @param [in,out] a
257  * @param [in] b
258  * @return a.x += b, a.y += b
259  */
260 inline __device__ void operator+=(float2& a,
261  const float b)
262 {
263  a.x += b;
264  a.y += b;
265 }// end of operator+=
266 //--------------------------------------------------------------------------------------------------
267 
268 #endif /* CUDA_UTILS_CUH */
__device__ float2 operator+(const float2 a, const float2 b)
Operator + for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:199
unsigned int nx
size of X dimension.
The header file for the class for storing constants residing in CUDA constant memory.
Structure for CUDA parameters to be placed in constant memory. Only 32b values are used...
unsigned int nxComplex
size of complex X dimension.
__device__ unsigned int GetIndex()
Get global 1D coordinate for 1D CUDA block.
Definition: CUDAUtils.cuh:57
__device__ float2 operator*(const float2 a, const float2 b)
Operator * for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:118
__device__ dim3 GetReal3DCoords(const unsigned int i)
Get 3D coordinates for a real matrix form a 1D index.
Definition: CUDAUtils.cuh:82
__device__ void operator+=(float2 &a, const float2 b)
Operator += for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:245
__device__ void operator*=(float2 &a, const float2 b)
Operator *= for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:164
__device__ unsigned int GetStride()
Get x-stride for 3D CUDA block (for processing multiple grid points by a single thread).
Definition: CUDAUtils.cuh:69
__constant__ TCUDADeviceConstants cudaDeviceConstants
This variable holds basic simulation constants for GPU.
unsigned int nyComplex
size of complex Y dimension.
unsigned int ny
size of Y dimension.
__device__ dim3 GetComplex3DCoords(const unsigned int i)
Get a 3D coordinates for a complex matrix form a 1D index.
Definition: CUDAUtils.cuh:97