diff --git a/src/Grid/grid_dist_id.hpp b/src/Grid/grid_dist_id.hpp index 4b5954dbb05438fcddecef3d7766325306c30e70..182080a0c7984b2664c8c66c767b4fb9138f9cee 100644 --- a/src/Grid/grid_dist_id.hpp +++ b/src/Grid/grid_dist_id.hpp @@ -2636,6 +2636,35 @@ public: } } + /*! \brief apply a convolution using the stencil N + * + * + */ + template<unsigned int prop_src, unsigned int prop_dst, unsigned int stencil_size, typename lambda_f, typename ... ArgsT > + void conv_cross_b(grid_key_dx<3> start, grid_key_dx<3> stop , lambda_f func, ArgsT ... args) + { + for (int i = 0 ; i < loc_grid.size() ; i++) + { + Box<dim,long int> inte; + + Box<dim,long int> base; + for (int j = 0 ; j < dim ; j++) + { + base.setLow(j,(long int)start.get(j) - (long int)gdb_ext.get(i).origin.get(j)); + base.setHigh(j,(long int)stop.get(j) - (long int)gdb_ext.get(i).origin.get(j)); + } + + Box<dim,long int> dom = gdb_ext.get(i).Dbox; + + bool overlap = dom.Intersect(base,inte); + + if (overlap == true) + { + loc_grid.get(i).template conv_cross_b<prop_src,prop_dst,stencil_size>(inte.getKP1(),inte.getKP2(),func,args...); + } + } + } + /*! \brief apply a convolution using the stencil N * * diff --git a/src/Grid/tests/sgrid_dist_id_gpu_unit_tests.cu b/src/Grid/tests/sgrid_dist_id_gpu_unit_tests.cu index 432cc219110021621867a3e9d9d69f120f3bbce4..9b79ea1b56f6f3aefa79afe68cc62402530a30b4 100644 --- a/src/Grid/tests/sgrid_dist_id_gpu_unit_tests.cu +++ b/src/Grid/tests/sgrid_dist_id_gpu_unit_tests.cu @@ -631,6 +631,101 @@ BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv2_test_3d ) BOOST_REQUIRE_EQUAL(match,true); } +BOOST_AUTO_TEST_CASE( sgrid_gpu_test_conv_cross_block_test_3d ) +{ + #ifdef CUDA_ON_CPU + size_t sz[3] = {20,20,20}; + #else + size_t sz[3] = {60,60,60}; + #endif + periodicity<3> bc = {PERIODIC,PERIODIC,PERIODIC}; + + Ghost<3,long int> g(1); + + Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0}); + + sgrid_dist_id_gpu<3,float,aggregate<float,float,float,float>> gdist(sz,domain,g,bc); + + gdist.template setBackgroundValue<0>(666); + gdist.template setBackgroundValue<1>(666); + gdist.template setBackgroundValue<2>(666); + gdist.template setBackgroundValue<3>(666); + + /////// GPU insert + flush + + Box<3,size_t> box({1,1,1},{sz[0]-1,sz[1]-1,sz[2]-1}); + + /////// GPU Run kernel + + float c = 5.0; + + typedef typename GetAddBlockType<decltype(gdist)>::type InsertBlockT; + + gdist.addPoints(box.getKP1(),box.getKP2(), + [] __device__ (int i, int j, int k) + { + return true; + }, + [c] __device__ (InsertBlockT & data, int i, int j, int k) + { + data.template get<0>() = c + i + j + k; + data.template get<1>() = c + 1000 + i + j + k; + } + ); + + gdist.template flush<smax_<0>,smax_<1>>(flush_type::FLUSH_ON_DEVICE); + + gdist.template ghost_get<0,1>(RUN_ON_DEVICE); + + for (int i = 0 ; i < 10 ; i++) + { + gdist.template ghost_get<0,1>(RUN_ON_DEVICE); + } + + // Now run the convolution + + typedef typename GetCpBlockType<decltype(gdist),0,1>::type CpBlockType; + + gdist.template conv_cross_b<0,1,1>({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2},[] __device__ (CpBlockType & u, auto & block, int offset,int i, int j, int k){ + return u(i+1,j,k) - u(i-1,j,k) + u(i,j+1,k) - u(i,j-1,k) + u(i,j,k+1) - u(i,j,k-1) + block.template get<0>()[offset]; + }); + + gdist.deviceToHost<0,1,2,3>(); + + // Now we check that ghost is correct + + auto it3 = gdist.getSubDomainIterator({2,2,2},{(int)sz[0]-2,(int)sz[1]-2,(int)sz[2]-2}); + + bool match = true; + + while (it3.isNext()) + { + auto p = it3.get(); + + auto p_xp1 = p.move(0,1); + auto p_xm1 = p.move(0,-1); + auto p_yp1 = p.move(1,1); + auto p_ym1 = p.move(1,-1); + auto p_zp1 = p.move(2,1); + auto p_zm1 = p.move(2,-1); + + float sub1 = gdist.template get<1>(p); + + if (sub1 != 6.0 + gdist.template get<0>(p)) + { + std::cout << sub1 << std::endl; + std::cout << gdist.template get<0>(p_xp1) << " " << gdist.template get<0>(p_xm1) << std::endl; + std::cout << gdist.template get<1>(p_xp1) << " " << gdist.template get<1>(p_xm1) << std::endl; + match = false; + break; + } + + ++it3; + } + + BOOST_REQUIRE_EQUAL(match,true); +} + BOOST_AUTO_TEST_CASE( sgrid_gpu_test_ghost_point_remove ) { size_t sz[3] = {60,60,60};