OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
 
Loading...
Searching...
No Matches
isolation.cu
1#include <iostream>
2#include <thread>
3
4size_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
15void 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:
38bool init_unit_test()
39{
40// std::thread to (timeout_cycle);
41// to.detach();
42 return true;
43}
44
45// entry point
46int 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
55BOOST_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}
This class represent an N-dimensional box.
Definition Box.hpp:61
This class implement the point shape in an N-dimensional space.
Definition Point.hpp:28
__device__ __host__ T distance(const Point< dim, T > &q) const
It calculate the distance between 2 points.
Definition Point.hpp:250
__device__ __host__ const T & get(unsigned int i) const
Get coordinate.
Definition Point.hpp:172
std::string toString() const
Return the string with the point coordinate.
Definition Point.hpp:398
size_t getProcessUnitID()
Get the process unit id.
size_t getProcessingUnits()
Get the total number of processors.
Implementation of VCluster class.
Definition VCluster.hpp:59
Distributed vector.
aggregate of properties, from a list of object if create a struct that follow the OPENFPM native stru...