OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
 
Loading...
Searching...
No Matches
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
15template<unsigned int dim>
16struct ite_gpu_dist
17{
18 dim3 thr;
19 dim3 wthr;
20
23
24 grid_key_dx<dim,int> start_base;
25
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
77template<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
83template<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.
grid_key_dx is the key to access any element in the grid
Definition grid_key.hpp:19