OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
grid_dist_id_kernels.cuh
1 /*
2  * grid_dist_id_kernels.cuh
3  *
4  * Created on: Jun 25, 2019
5  * Author: i-bird
6  */
7 
8 #ifndef GRID_DIST_ID_KERNELS_CUH_
9 #define GRID_DIST_ID_KERNELS_CUH_
10 
11 #include "config.h"
12 
13 #ifdef CUDA_GPU
14 
15 template<unsigned int dim>
16 struct ite_gpu_dist
17 {
18  dim3 thr;
19  dim3 wthr;
20 
23 
24  grid_key_dx<dim,int> start_base;
25 
26  grid_key_dx<dim,int> origin;
27 
28  ite_gpu_dist()
29  {}
30 
31  ite_gpu_dist(ite_gpu<dim> & ite)
32  {
33  thr = ite.thr;
34  wthr = ite.wthr;
35 
36  start = ite.start;
37  stop = ite.stop;
38  }
39 
40  size_t nblocks()
41  {
42  return wthr.x * wthr.y * wthr.z;
43  }
44 };
45 
46 #define GRID_ID_3_GLOBAL(ite_gpu) grid_key_dx<3,int> key;\
47  grid_key_dx<3,int> keyg;\
48  key.set_d(0,threadIdx.x + blockIdx.x * blockDim.x + ite_gpu.start.get(0));\
49  key.set_d(1,threadIdx.y + blockIdx.y * blockDim.y + ite_gpu.start.get(1));\
50  key.set_d(2,threadIdx.z + blockIdx.z * blockDim.z + ite_gpu.start.get(2));\
51  \
52  bool inactive = false;\
53  \
54  keyg.set_d(0,key.get(0) + ite_gpu.origin.get(0));\
55  keyg.set_d(1,key.get(1) + ite_gpu.origin.get(1));\
56  keyg.set_d(2,key.get(2) + ite_gpu.origin.get(2));\
57  \
58  if (key.get(0) > ite_gpu.stop.get(0) || key.get(1) > ite_gpu.stop.get(1) || key.get(2) > ite_gpu.stop.get(2))\
59  {inactive = true;}
60 
61 
62 #define GRID_ID_2_GLOBAL(ite_gpu) grid_key_dx<2,int> key;\
63  grid_key_dx<2,int> keyg;\
64  key.set_d(0,threadIdx.x + blockIdx.x * blockDim.x + ite_gpu.start.get(0));\
65  key.set_d(1,threadIdx.y + blockIdx.y * blockDim.y + ite_gpu.start.get(1));\
66  \
67  bool inactive = false;\
68  \
69  keyg.set_d(0,key.get(0) + ite_gpu.origin.get(0));\
70  keyg.set_d(1,key.get(1) + ite_gpu.origin.get(1));\
71  \
72  if (key.get(0) > ite_gpu.stop.get(0) || key.get(1) > ite_gpu.stop.get(1))\
73  {inactive = true;}
74 
75 #endif
76 
77 template<typename grid_type, typename ite_gpu_type,typename func_t,typename ... args_t>
78 __global__ void grid_apply_functor(grid_type g, ite_gpu_type ite, func_t f, args_t ... args)
79 {
80  f(g,ite,args...);
81 }
82 
83 template<typename grid_type, typename ite_gpu_type,typename func_t,typename ... args_t>
84 __global__ void grid_apply_functor_shared_bool(grid_type g, ite_gpu_type ite, func_t f, args_t ... args)
85 {
86  __shared__ bool is_empty_block;
87 
88  f(g,ite,is_empty_block,args...);
89 }
90 
91 #endif /* GRID_DIST_ID_KERNELS_CUH_ */
This is a distributed grid.