OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
isolation.cu
1 #include <iostream>
2 #include <thread>
3 
4 size_t debug_tot_call = 0;
5 
6 #define PRINT_STACKTRACE
7 #define CHECKFOR_POSNAN
8 #define CHECKFOR_POSINF
9 #define CHECKFOR_PROPNAN
10 #define CHECKFOR_PROPINF
11 
12 #define NO_WARNING
13 #include "Graph/CartesianGraphFactory.hpp"
14 
15 void timeout_cycle()
16 {
17  // 6 seconds
18  std::this_thread::sleep_for (std::chrono::seconds(900));
19 
20  std::cout << "Time Out" << std::endl;
21  std::exit(1);
22 }
23 
24 
25 #define BOOST_DISABLE_ASSERTS
26 
27 
28 #include "config.h"
29 #undef VERSION
30 
31 #define BOOST_TEST_DYN_LINK
32 #include <boost/test/unit_test.hpp>
33 #include "VCluster/VCluster.hpp"
34 #include <Vector/vector_dist.hpp>
35 #include "Vector/tests/vector_dist_util_unit_tests.hpp"
36 
37 // initialization function:
38 bool init_unit_test()
39 {
40 // std::thread to (timeout_cycle);
41 // to.detach();
42  return true;
43 }
44 
45 // entry point
46 int main(int argc, char* argv[])
47 {
48  openfpm_init(&argc,&argv);
49 
50  return boost::unit_test::unit_test_main( &init_unit_test, argc, argv );
51 }
52 
53 
54 
55 BOOST_AUTO_TEST_CASE( vector_dist_ghost_put_gpu )
56 {
57 
58 
59  Vcluster<> & v_cl = create_vcluster();
60 
61  long int k = 25*25*25*create_vcluster().getProcessingUnits();
62  k = std::pow(k, 1/3.);
63 
64  if (v_cl.getProcessingUnits() > 48)
65  return;
66 
67  BOOST_TEST_CHECKPOINT( "Testing 3D periodic ghost put k=" << k );
68 
69  long int big_step = k / 30;
70  big_step = (big_step == 0)?1:big_step;
71  long int small_step = 21;
72 
73  // 3D test
74  for ( ; k >= 2 ; k-= (k > 2*big_step)?big_step:small_step )
75  {
76  float r_cut = 1.3 / k;
77  float r_g = 1.5 / k;
78 
79  Box<3,float> box({0.0,0.0,0.0},{1.0,1.0,1.0});
80 
81  // Boundary conditions
82  size_t bc[3]={PERIODIC,PERIODIC,PERIODIC};
83 
84  // ghost
85  Ghost<3,float> ghost(r_g);
86 
87  typedef aggregate<float,float,float> part_prop;
88 
89  // Distributed vector
90  vector_dist_gpu<3,float, part_prop > vd(0,box,bc,ghost);
91 
92  auto it = vd.getGridIterator({(size_t)k,(size_t)k,(size_t)k});
93 
94  while (it.isNext())
95  {
96  auto key = it.get();
97 
98  vd.add();
99 
100  vd.getLastPosWrite()[0] = key.get(0)*it.getSpacing(0);
101  vd.getLastPosWrite()[1] = key.get(1)*it.getSpacing(1);
102  vd.getLastPosWrite()[2] = key.get(2)*it.getSpacing(2);
103 
104  // Fill some properties randomly
105 
106  vd.getLastPropWrite<0>() = 0.0;
107 
108  vd.getLastPropWrite<2>() = 0.0;
109 
110  ++it;
111  }
112 
113  vd.map();
114 
115  vd.hostToDevicePos();
116  vd.template hostToDeviceProp<0,2>();
117  // sync the ghost
118  vd.ghost_get<0,2>(RUN_ON_DEVICE);
119  vd.template deviceToHostProp<0,2>();
120  vd.deviceToHostPos();
121 
122  {
123  auto NN = vd.getCellList(r_cut);
124  float a = 1.0f*k*k;
125 
126  // run trough all the particles + ghost
127 
128  auto it2 = vd.getDomainIterator();
129 
130  while (it2.isNext())
131  {
132  // particle p
133  auto p = it2.get();
134  Point<3,float> xp = vd.getPos(p);
135 
136  // Get an iterator over the neighborhood particles of p
137  auto Np = NN.getNNIterator<NO_CHECK>(NN.getCell(xp));
138 
139  // For each neighborhood particle ...
140  while (Np.isNext())
141  {
142  auto q = Np.get();
143  Point<3,float> xq = vd.getPosRead(q);
144 
145  float dist = xp.distance(xq);
146 
147  if (dist < r_cut)
148  {
149  vd.getPropWrite<0>(q) += a*(-dist*dist+r_cut*r_cut);
150  vd.getPropWrite<2>(q) += a*(-dist*dist+r_cut*r_cut);
151  }
152 
153  ++Np;
154  }
155 
156  ++it2;
157  }
158 
159  vd.hostToDevicePos();
160  vd.template hostToDeviceProp<0,2>();
161  vd.template ghost_put<add_atomic_,0,2>(RUN_ON_DEVICE);
162  vd.template deviceToHostProp<0,2>();
163  vd.deviceToHostPos();
164 
165  bool ret = true;
166  auto it3 = vd.getDomainIterator();
167 
168  float constant = vd.getProp<0>(it3.get());
169  float eps = 0.001;
170 
171  while (it3.isNext())
172  {
173  float constant2 = vd.getProp<0>(it3.get());
174  float constant3 = vd.getProp<2>(it3.get());
175  if (fabs(constant - constant2)/constant > eps || fabs(constant - constant3)/constant > eps)
176  {
177  Point<3,float> p = vd.getPosRead(it3.get());
178 
179  std::cout << p.toString() << " " << constant2 << "/" << constant << "/" << constant3 << " " << v_cl.getProcessUnitID() << std::endl;
180  ret = false;
181  break;
182  }
183 
184  ++it3;
185  }
186  BOOST_REQUIRE_EQUAL(ret,true);
187  }
188 
189  auto itp = vd.getDomainAndGhostIterator();
190  while (itp.isNext())
191  {
192  auto key = itp.get();
193 
194  vd.getPropWrite<0>(key) = 0.0;
195  vd.getPropWrite<2>(key) = 0.0;
196 
197  ++itp;
198  }
199 
200  {
201  auto NN = vd.getCellList(r_cut);
202  float a = 1.0f*k*k;
203 
204  // run trough all the particles + ghost
205 
206  auto it2 = vd.getDomainIterator();
207 
208  while (it2.isNext())
209  {
210  // particle p
211  auto p = it2.get();
212  Point<3,float> xp = vd.getPosRead(p);
213 
214  // Get an iterator over the neighborhood particles of p
215  auto Np = NN.getNNIterator<NO_CHECK>(NN.getCell(xp));
216 
217  // For each neighborhood particle ...
218  while (Np.isNext())
219  {
220  auto q = Np.get();
221  Point<3,float> xq = vd.getPosRead(q);
222 
223  float dist = xp.distance(xq);
224 
225  if (dist < r_cut)
226  {
227  vd.getPropWrite<0>(q) += a*(-dist*dist+r_cut*r_cut);
228  vd.getPropWrite<2>(q) += a*(-dist*dist+r_cut*r_cut);
229  }
230 
231  ++Np;
232  }
233 
234  ++it2;
235  }
236 
237  vd.hostToDevicePos();
238  vd.template hostToDeviceProp<0,2>();
239  vd.template ghost_put<add_atomic_,0>(RUN_ON_DEVICE);
240  vd.template ghost_put<add_atomic_,2>(RUN_ON_DEVICE);
241  vd.template deviceToHostProp<0,2>();
242  vd.deviceToHostPos();
243 
244  bool ret = true;
245  auto it3 = vd.getDomainIterator();
246 
247  float constant = vd.getPropRead<0>(it3.get());
248  float eps = 0.001;
249 
250  while (it3.isNext())
251  {
252  float constant2 = vd.getPropRead<0>(it3.get());
253  float constant3 = vd.getPropRead<0>(it3.get());
254  if (fabs(constant - constant2)/constant > eps || fabs(constant - constant3)/constant > eps)
255  {
256  Point<3,float> p = vd.getPosRead(it3.get());
257 
258  std::cout << p.toString() << " " << constant2 << "/" << constant << "/" << constant3 << " " << it3.get().getKey() << " " << v_cl.getProcessUnitID() << std::endl;
259  ret = false;
260  break;
261  }
262 
263  ++it3;
264  }
265  BOOST_REQUIRE_EQUAL(ret,true);
266  }
267  }
268 
269  openfpm_finalize();
270 }
size_t getProcessUnitID()
Get the process unit id.
std::string toString() const
Return the string with the point coordinate.
Definition: Point.hpp:398
Definition: Ghost.hpp:39
Implementation of VCluster class.
Definition: VCluster.hpp:58
__device__ __host__ const T & get(unsigned int i) const
Get coordinate.
Definition: Point.hpp:172
size_t getProcessingUnits()
Get the total number of processors.
__device__ __host__ T distance(const Point< dim, T > &q) const
It calculate the distance between 2 points.
Definition: Point.hpp:250
Distributed vector.
aggregate of properties, from a list of object if create a struct that follow the OPENFPM native stru...
Definition: aggregate.hpp:214