Alsvinn  0.5.3
The fast FVM simulator with UQ support
cuda_utils.hpp
Go to the documentation of this file.
1 /* Copyright (c) 2018 ETH Zurich, Kjetil Olsen Lye
2  * This program is free software: you can redistribute it and/or modify
3  * it under the terms of the GNU General Public License as published by
4  * the Free Software Foundation, either version 3 of the License, or
5  * (at your option) any later version.
6  *
7  * This program is distributed in the hope that it will be useful,
8  * but WITHOUT ANY WARRANTY; without even the implied warranty of
9  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
10  * GNU General Public License for more details.
11  *
12  * You should have received a copy of the GNU General Public License
13  * along with this program. If not, see <http://www.gnu.org/licenses/>.
14  */
15 
16 #pragma once
17 #include "cuda.h"
18 #include "cuda_runtime.h"
19 #include <iostream>
20 #include <exception>
21 #include "alsutils/types.hpp"
24 
25 #ifdef NDEBUG
26 #define CUDA_CHECK_IF_DEBUG
27 #else
28 #define CUDA_CHECK_IF_DEBUG { \
31  CUDA_SAFE_CALL(cudaGetLastError()); \
32 CUDA_SAFE_CALL(cudaDeviceSynchronize()); \
33 CUDA_SAFE_CALL(cudaGetLastError()); \
34 }
35 #endif
36 namespace alsfvm {
37 namespace cuda {
38 inline dim3 calculateBlockDimensions(size_t numberOfXCells,
39  size_t numberOfYCells, size_t numberOfZCells) {
40  const size_t blockSize = 1024;
41  return dim3(blockSize, numberOfYCells > 1 ? blockSize : 1,
42  numberOfZCells > 1 ? blockSize : 1);
43 }
44 
45 inline dim3 calculateGridDimensions(size_t numberOfXCells,
46  size_t numberOfYCells, size_t numberOfZCells, dim3 blockDimensions) {
47  return dim3((numberOfXCells + blockDimensions.x - 1) / blockDimensions.x,
48  (numberOfYCells + blockDimensions.y - 1) / blockDimensions.y,
49  (numberOfZCells + blockDimensions.z - 1) / blockDimensions.z);
50 }
51 
52 
57 inline ivec3 __device__ getCoordinates(dim3 threadIdx, dim3 blockIdx,
58  dim3 blockDim,
59  size_t numberOfXCells, size_t numberOfYCells,
60  size_t numberOfZCells, ivec3 directionVector) {
61  const int index = threadIdx.x + blockIdx.x * blockDim.x;
62 
63  const size_t xInternalFormat = index % numberOfXCells;
64  const size_t yInternalFormat = (index / numberOfXCells) % numberOfYCells;
65  const size_t zInternalFormat = (index) / (numberOfXCells * numberOfYCells);
66 
67  if (xInternalFormat >= numberOfXCells || yInternalFormat >= numberOfYCells
68  || zInternalFormat >= numberOfZCells) {
69  return {-1, -1, -1};
70  }
71 
72  const int x = xInternalFormat + directionVector[0];
73  const int y = yInternalFormat + directionVector[1];
74  const int z = zInternalFormat + directionVector[2];
75 
76  return {x, y, z};
77 }
78 
79 
81 inline std::tuple<int, ivec3> makeKernelLaunchParameters(ivec3 start, ivec3 end,
82  size_t blockSize) {
83  const ivec3 numberOfCellsPerDimension = end - start;
84 
85  const size_t totalNumberOfCells = size_t(numberOfCellsPerDimension.x) *
86  size_t(numberOfCellsPerDimension.y) *
87  size_t(numberOfCellsPerDimension.z);
88 
89  const int gridSize = (totalNumberOfCells + blockSize - 1 ) / blockSize;
90 
91  return std::make_tuple(gridSize, numberOfCellsPerDimension);
92 }
93 }
94 }
T z
Definition: vec3.hpp:28
#define __device__
Definition: types.hpp:45
std::tuple< int, ivec3 > makeKernelLaunchParameters(ivec3 start, ivec3 end, size_t blockSize)
Gets teh kernel launch paramemters.
Definition: cuda_utils.hpp:81
ivec3 __device__ getCoordinates(dim3 threadIdx, dim3 blockIdx, dim3 blockDim, size_t numberOfXCells, size_t numberOfYCells, size_t numberOfZCells, ivec3 directionVector)
Definition: cuda_utils.hpp:57
T y
Definition: vec3.hpp:27
dim3 calculateGridDimensions(size_t numberOfXCells, size_t numberOfYCells, size_t numberOfZCells, dim3 blockDimensions)
Definition: cuda_utils.hpp:45
Various utility functions to implement the tecno flux.
Definition: types.hpp:30
dim3 calculateBlockDimensions(size_t numberOfXCells, size_t numberOfYCells, size_t numberOfZCells)
Definition: cuda_utils.hpp:38
T x
Definition: vec3.hpp:26