| kspaceFirstOrder3D-CUDA 1.1 The CUDA/C++ implementation of the k-wave toolbox for the time-domain simulation of acoustic wave fields in 3D

- Main Page
- Related Pages
- Namespaces
- Classes
- Files

- File List
- File Members

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
36 #include <Parameters/CUDADeviceConstants.cuh>
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,
85 (i / cudaDeviceConstants.nx) % cudaDeviceConstants.ny,
86 i / ( cudaDeviceConstants.nx * cudaDeviceConstants.ny));
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,
100 (i / cudaDeviceConstants.nxComplex) % cudaDeviceConstants.nyComplex,
101 i / ( cudaDeviceConstants.nxComplex * cudaDeviceConstants.nyComplex));
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 */
operator+
__device__ float2 operator+(const float2 a, const float2 b)
Operator + for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:199
TCUDADeviceConstants::nx
unsigned int nx
size of X dimension.
Definition: CUDADeviceConstants.cuh:48
CUDADeviceConstants.cuh
The header file for the class for storing constants residing in CUDA constant memory.
TCUDADeviceConstants
Structure for CUDA parameters to be placed in constant memory. Only 32b values are used...
Definition: CUDADeviceConstants.cuh:42
TCUDADeviceConstants::nxComplex
unsigned int nxComplex
size of complex X dimension.
Definition: CUDADeviceConstants.cuh:58
GetIndex
__device__ unsigned int GetIndex()
Get global 1D coordinate for 1D CUDA block.
Definition: CUDAUtils.cuh:57
operator*
__device__ float2 operator*(const float2 a, const float2 b)
Operator * for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:118
GetReal3DCoords
__device__ dim3 GetReal3DCoords(const unsigned int i)
Get 3D coordinates for a real matrix form a 1D index.
Definition: CUDAUtils.cuh:82
operator+=
__device__ void operator+=(float2 &a, const float2 b)
Operator += for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:245
operator*=
__device__ void operator*=(float2 &a, const float2 b)
Operator *= for float2 datatype (per element multiplication).
Definition: CUDAUtils.cuh:164
GetStride
__device__ unsigned int GetStride()
Get x-stride for 3D CUDA block (for processing multiple grid points by a single thread).
Definition: CUDAUtils.cuh:69
cudaDeviceConstants
__constant__ TCUDADeviceConstants cudaDeviceConstants
This variable holds basic simulation constants for GPU.
Definition: CUDADeviceConstants.cu:54
TCUDADeviceConstants::nyComplex
unsigned int nyComplex
size of complex Y dimension.
Definition: CUDADeviceConstants.cuh:60
TCUDADeviceConstants::ny
unsigned int ny
size of Y dimension.
Definition: CUDADeviceConstants.cuh:50
GetComplex3DCoords
__device__ dim3 GetComplex3DCoords(const unsigned int i)
Get a 3D coordinates for a complex matrix form a 1D index.
Definition: CUDAUtils.cuh:97