cuda_grid_gpu_int.cu

#include "config.h"
#define BOOST_TEST_DYN_LINK
#include <boost/test/unit_test.hpp>
#include "Space/Shape/Box.hpp"
#include "Vector/map_vector.hpp"
#include "NN/CellList/cuda/CellList_gpu.hpp"
#define DISABLE_MPI_WRITTERS
#include "VTKWriter/VTKWriter.hpp"

template<typename vector_pos_type, typename vector_prop_type>
__global__ void test_launch(vector_pos_type set_points, vector_prop_type prop, Box<3,float> domain)
{
    int p = threadIdx.x + blockIdx.x * blockDim.x;

    if (p >= set_points.size())
    {return;}

    float pos[3];
    pos[0] = set_points.template get<0>(p)[0];
    pos[1] = set_points.template get<0>(p)[1];
    pos[2] = set_points.template get<0>(p)[1];

    float scalar = prop.template get<0>(p);

    float v[3];

    v[0] = prop.template get<1>(p)[0];
    v[1] = prop.template get<1>(p)[1];
    v[2] = prop.template get<1>(p)[2];
}

template<typename grid_type>
__global__ void test_launch_grid(grid_type grid, Box<3,float> domain, ite_gpu<3> ite_gpu)
{
    GRID_ID_3(ite_gpu);

    float scalar = grid.template get<0>(key);

    float v[3];

    v[0] = grid.template get<1>(key)[0];
    v[1] = grid.template get<1>(key)[1];
    v[2] = grid.template get<1>(key)[2];

    printf("Grid point %d %d %d     scalar: %f  vector: %f %f %f \n",(int)key.get(0),(int)key.get(1),(int)key.get(2),scalar,v[0],v[1],v[2]);
}

__global__ void test_launch_cuda_native(float * scalar, float * vector, int sxy, int sx , int sy , int sz , int stride)
{
    int id[3];

    id[0] = threadIdx.x + blockIdx.x * blockDim.x;
    id[1] = threadIdx.y + blockIdx.y * blockDim.y;
    id[2] = threadIdx.z + blockIdx.z * blockDim.z;

    if (id[0] >= sx) {return;}
    if (id[1] >= sy) {return;}
    if (id[2] >= sz) {return;}

    float s = scalar[id[2]*sxy+id[1]*sx+id[0]];

    float v[3];

    v[0] = vector[id[2]*sxy+id[1]*sx+id[0] + 0*stride];
    v[1] = vector[id[2]*sxy+id[1]*sx+id[0] + 1*stride];
    v[2] = vector[id[2]*sxy+id[1]*sx+id[0] + 2*stride];

    printf("Grid point from CUDA %d %d %d     scalar: %f  vector: %f %f %f \n",id[0],id[1],id[2],s,v[0],v[1],v[2]);
}

constexpr int NN_num = 4;

/*template<typename celllist_type>
__global__ void test_launch_cell_list(celllist_type cell, ite_gpu<3> ite_gpu)
{
    GRID_ID_3(ite_gpu)

    int nn_part = 0;

    int idx = 0;
    int NN[NN_num];

    auto NN_it = cell.template getNNIteratorBox<2>(key);

    while (NN_it.isNext())
    {
        auto q = NN_it.get();

        if (idx < NN_num)
        {
            NN[idx] = q;
            idx++;
        }

        nn_part++;

        ++NN_it;
    }

    printf("CELLLIST %d %d %d  nn_part: %d NN: %d %d %d %d \n",(int)key.get(0),(int)key.get(1),(int)key.get(2),nn_part,NN[0],NN[1],NN[2],NN[3]);
}*/

BOOST_AUTO_TEST_SUITE( grid_gpu_func_interp )

BOOST_AUTO_TEST_CASE (gpu_p2m)
{
    openfpm::vector_gpu<Point<3,float>> pos;
    openfpm::vector_gpu<aggregate<float,float[3]>> prop;

    pos.resize(100);
    prop.resize(100);

    for (size_t i = 0 ; i < 100 ; i++)
    {
        pos.template get<0>(i)[0] = (float)rand() / RAND_MAX;
        pos.template get<0>(i)[1] = (float)rand() / RAND_MAX;
        pos.template get<0>(i)[2] = (float)rand() / RAND_MAX;

        prop.template get<0>(i) = 5.0;
        prop.template get<1>(i)[0] = 3.0;
        prop.template get<1>(i)[1] = 4.0;
        prop.template get<1>(i)[2] = 5.0;
    }

    pos.template hostToDevice<0>();
    prop.template hostToDevice<0,1>();

    Box<3,float> domain({0.0,0.0,0.0},{1.0,1.0,1.0});

    auto ite = pos.getGPUIterator();

    test_launch<<<ite.wthr,ite.thr>>>(pos.toKernel(),prop.toKernel(),domain);

    grid_gpu<3,aggregate<float,float[3]>> grid;

    grid.resize({10,10,10});

    auto it = grid.getIterator();

    while (it.isNext())
    {
        auto key = it.get();

        grid.template get<0>(key) = key.get(0) + key.get(1) + key.get(2);
        grid.template get<1>(key)[0] = key.get(0);
        grid.template get<1>(key)[1] = key.get(1);
        grid.template get<1>(key)[2] = key.get(2);

        ++it;
    }

    grid_key_dx<3> start({0,0,0});
    grid_key_dx<3> stop({9,9,9});

    auto ite_g = grid.getGPUIterator(start,stop);

    grid.template hostToDevice<0,1>();

    test_launch_grid<<<ite_g.wthr,ite_g.thr>>>(grid.toKernel(),domain,ite_g);


    test_launch_cuda_native<<<ite_g.wthr,ite_g.thr>>>((float *)grid.template getDeviceBuffer<0>(),(float *)grid.template getDeviceBuffer<1>(),100,10,10,10,grid.size());


    openfpm::vector_gpu<Point<3,float>> pos_sort;
    openfpm::vector_gpu<aggregate<float,float[3]>> prop_sort;

    pos_sort.resize(pos.size());
    prop_sort.resize(prop.size());

    size_t g_m = pos.size();

    mgpu::ofp_context_t context(false);

    const size_t (& sz)[3] = grid.getGrid().getSize();

    CellList_gpu<3,float,CudaMemory,shift_only<3,float>> cl(domain,sz,2);

    cl.template construct(pos,pos_sort,prop,prop_sort,context,g_m);

    grid_key_dx<3> start_c({2,2,2});
    grid_key_dx<3> stop_c({11,11,11});

    auto ite_gpu = getGPUIterator_impl(cl.getGrid(),start_c,stop_c);

    test_launch_cell_list<<<ite_gpu.wthr,ite_gpu.thr>>>(cl.toKernel(),ite_gpu);


    // VTKWriter for a set of points
    VTKWriter<boost::mpl::pair<openfpm::vector_gpu<Point<3,float>>,
                               openfpm::vector_gpu<aggregate<float,float[3]>>>,
                               VECTOR_POINTS> vtk_writer;
    vtk_writer.add(pos,prop,pos.size());

    openfpm::vector<std::string> prp_names;
    prp_names.add("scalar");
    prp_names.add("vector");

    file_type ft = file_type::ASCII;

    pos.template deviceToHost<0>();
    prop.template deviceToHost<0,1>();

    // Write the VTK file
    vtk_writer.write("particles.vtk",prp_names,"particles",ft);
}

BOOST_AUTO_TEST_CASE (gpu_m2p)
{

}

BOOST_AUTO_TEST_SUITE_END()