decomposition_cuda_tests.cu

#define BOOST_TEST_DYN_LINK
#include <boost/test/unit_test.hpp>
 
#include "VCluster/VCluster.hpp"
#include "Decomposition/CartDecomposition.hpp"
 
#define SUB_UNIT_FACTOR 1024
 
template<typename dec_type>
__global__ void test_proc_idbc(Point<3,double> p1 ,Point<3,double> p2 , dec_type dec, unsigned int * pr_id)
{
    pr_id[0] = dec.processorIDBC(p1);
    pr_id[1] = dec.processorIDBC(p2);
}
 
template<typename dec_type>
__global__ void test_ghost_n(Point<3,double> p1 ,Point<3,double> p2 , dec_type dec, unsigned int * ng_id)
{
    ng_id[0] = dec.ghost_processorID_N(p1);
    ng_id[1] = dec.ghost_processorID_N(p2);
}
 
template<typename dec_type, typename output_type>
__global__ void test_ghost(Point<3,double> p1 ,Point<3,double> p2 , dec_type dec, unsigned int * ng_id , output_type g_id)
{
    for (unsigned int i = 0 ; i < ng_id[0] ; i++)
    {
        dec.ghost_processor_ID(p1,g_id,0,i);
    }
 
    for (unsigned int i = 0 ; i < ng_id[1] ; i++)
    {
        dec.ghost_processor_ID(p2,g_id,ng_id[0],i);
    }
}
 
BOOST_AUTO_TEST_SUITE( decomposition_to_gpu_test )
 
BOOST_AUTO_TEST_CASE( CartDecomposition_check_cross_consistency_between_proc_idbc_and_ghost2_gpu )
{
    // Vcluster
    Vcluster<> & vcl = create_vcluster();
 
    CartDecomposition<3, double, CudaMemory,memory_traits_inte> dec(vcl);
 
    size_t bc[3] = {PERIODIC,PERIODIC,PERIODIC};
 
    // Physical domain
    Box<3, double> box( { -0.01, -0.01, 0.0 }, { 0.01, 0.01, 0.003 });
 
    Ghost<3,double> g(0.0015);
 
    dec.setGoodParameters(box, bc, g, 512);
 
    dec.decompose();
 
    // Now we check the point
 
    for (size_t j = 0 ; j < 3 ; j++ )
    {
        for (size_t i = 0 ; i < dec.getNSubDomain() ; i++)
        {
            Point<3,double> p1;
            Point<3,double> p2;
 
            p1.get(0) = SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(0);
            p1.get(1) = SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(1);
            p1.get(2) = SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(2);
 
            p2 = p1;
 
//          p2.get(j) = std::nextafter(SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(j),-1.0);
 
            auto gpudec = dec.toKernel();
 
            CudaMemory mem;
            mem.allocate(2*sizeof(unsigned int));
 
            CUDA_LAUNCH_DIM3((test_proc_idbc<decltype(gpudec)>),1,1,p1,p2,gpudec,(unsigned int *)mem.getDevicePointer());
 
            mem.deviceToHost();
 
            BOOST_REQUIRE(((unsigned int *)mem.getPointer())[0] < vcl.size());
            BOOST_REQUIRE(((unsigned int *)mem.getPointer())[1] < vcl.size());
 
            CudaMemory mem2;
            mem2.allocate(2*sizeof(unsigned int));
            CUDA_LAUNCH_DIM3((test_ghost_n<decltype(gpudec)>),1,1,p1,p2,gpudec,(unsigned int *)mem2.getDevicePointer());
 
            mem2.deviceToHost();
 
            unsigned int tot = ((unsigned int *)mem2.getPointer())[0] + ((unsigned int *)mem2.getPointer())[1];
 
            openfpm::vector_gpu<aggregate<int,int>> vd;
            vd.resize(tot);
            CUDA_LAUNCH_DIM3((test_ghost<decltype(gpudec),decltype(vd.toKernel())>),1,1,p1,p2,gpudec,(unsigned int *)mem2.getDevicePointer(),vd.toKernel());
 
            if (((unsigned int *)mem.getPointer())[0] != ((unsigned int *)mem.getPointer())[1])
            {
                if (vcl.rank() == ((unsigned int *)mem.getPointer())[1] )
                {
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[1] != 0);
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[0] == 0);
                }
 
                if (vcl.rank() == ((unsigned int *)mem.getPointer())[0])
                {
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[1] == 0 );
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[0] != 0 );
                }
            }
 
 
            p1.get(0) = std::nextafter(SpaceBox<3,double>(dec.getSubDomains().get(i)).getHigh(0),SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(0));
            p1.get(1) = std::nextafter(SpaceBox<3,double>(dec.getSubDomains().get(i)).getHigh(1),SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(1));
            p1.get(2) = std::nextafter(SpaceBox<3,double>(dec.getSubDomains().get(i)).getHigh(2),SpaceBox<3,double>(dec.getSubDomains().get(i)).getLow(2));
 
            p2 = p1;
 
            p2.get(j) = std::nextafter(SpaceBox<3,double>(dec.getSubDomains().get(i)).getHigh(j),1.0);
 
            CUDA_LAUNCH_DIM3((test_proc_idbc<decltype(gpudec)>),1,1,p1,p2,gpudec,(unsigned int *)mem.getDevicePointer());
 
            mem.deviceToHost();
 
            BOOST_REQUIRE(((unsigned int *)mem.getPointer())[0] < vcl.size());
            BOOST_REQUIRE(((unsigned int *)mem.getPointer())[1] < vcl.size());
 
            mem2.allocate(2*sizeof(unsigned int));
            CUDA_LAUNCH_DIM3((test_ghost_n<decltype(gpudec)>),1,1,p1,p2,gpudec,(unsigned int *)mem2.getDevicePointer());
 
            mem2.deviceToHost();
 
            tot = ((unsigned int *)mem2.getPointer())[0] + ((unsigned int *)mem2.getPointer())[1];
 
            vd.resize(tot);
            CUDA_LAUNCH_DIM3((test_ghost<decltype(gpudec),decltype(vd.toKernel())>),1,1,p1,p2,gpudec,(unsigned int *)mem2.getDevicePointer(),vd.toKernel());
 
            if (((unsigned int *)mem.getPointer())[0] != ((unsigned int *)mem.getPointer())[1])
            {
                if (vcl.rank() == ((unsigned int *)mem.getPointer())[1])
                {
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[1] != 0);
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[0] == 0);
                }
 
                if (vcl.rank() == ((unsigned int *)mem.getPointer())[0])
                {
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[1] == 0 );
                    BOOST_REQUIRE(((unsigned int *)mem2.getPointer())[0] != 0 );
                }
            }
 
        }
    }
}
 
BOOST_AUTO_TEST_SUITE_END()