sgrid_dist_id_gpu_unit_tests.cu 5.31 KB
Newer Older
1 2 3 4 5 6 7 8 9 10
#define BOOST_TEST_DYN_LINK

#include <boost/test/unit_test.hpp>
#include "Grid/grid_dist_id.hpp"

BOOST_AUTO_TEST_SUITE( sgrid_gpu_test_suite )

template<unsigned int p>
struct insert_kernel
{
11 12
	template<typename SparseGridGpu_type, typename ite_type>
	__device__ void operator()(SparseGridGpu_type & sg, ite_type & ite, float c)
13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32
	{
	    sg.init();

	    const auto bDimX = blockDim.x;
	    const auto bDimY = blockDim.y;
	    const auto bDimZ = blockDim.z;
	    const auto bIdX = blockIdx.x;
	    const auto bIdY = blockIdx.y;
	    const auto bIdZ = blockIdx.z;
	    const auto tIdX = threadIdx.x;
	    const auto tIdY = threadIdx.y;
	    const auto tIdZ = threadIdx.z;
	    int x = bIdX * bDimX + tIdX;
	    int y = bIdY * bDimY + tIdY;
	    int z = bIdZ * bDimZ + tIdZ;

	    if (x+ite.start.get(0) > ite.stop.get(0))
	    {return;}
	    if (SparseGridGpu_type::d >= 2 && y+ite.start.get(1) > ite.stop.get(1))
	    {return;}
33
	    if (SparseGridGpu_type::d >= 3 && z+ite.start.get(2) > ite.stop.get(2))
34 35
	    {return;}

36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60
	    grid_key_dx<SparseGridGpu_type::d, size_t> coord;
	    grid_key_dx<SparseGridGpu_type::d, size_t> coord_glob;

	    if (SparseGridGpu_type::d >= 2)
	    {
	    	coord.set_d(0,x+ite.start.get(0));
	    	coord_glob.set_d(0,x+ite.start.get(0)+ite.origin.get(0));
	    	coord.set_d(1,y+ite.start.get(1));
	    	coord_glob.set_d(1,y+ite.start.get(1)+ite.origin.get(1));
	    }
	    else if (SparseGridGpu_type::d >= 3)
	    {
		    coord.set_d(0,x+ite.start.get(0));
		    coord_glob.set_d(0,x+ite.start.get(0)+ite.origin.get(0));
		    coord.set_d(1,y+ite.start.get(1));
		    coord_glob.set_d(1,y+ite.start.get(1)+ite.origin.get(1));
		    coord.set_d(2,z+ite.start.get(2));
		    coord_glob.set_d(2,z+ite.start.get(2)+ite.origin.get(2));
	    }


	    if (SparseGridGpu_type::d >= 2)
	    {sg.template insert<p>(coord) = c + coord_glob.get(0) + coord_glob.get(1);}
	    else
	    {sg.template insert<p>(coord) = c + coord_glob.get(0) + coord_glob.get(1) + coord_glob.get(2);}
61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101

	    __syncthreads();

	    sg.flush_block_insert();

	    // Compiler avoid warning
	    y++;
	    z++;
	}
};

template<unsigned int p>
struct stencil_kernel
{
	template<typename SparseGridGpu_type>
	__device__ void operator()(SparseGridGpu_type & sg, ite_gpu<SparseGridGpu_type::d> & ite, float c)
	{
		// TODO
	}
};

BOOST_AUTO_TEST_CASE( sgrid_gpu_test_base )
{
	size_t sz[2] = {17,17};
	periodicity<2> bc = {PERIODIC,PERIODIC};

	Ghost<2,long int> g(1);

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

	sgrid_dist_id_gpu<2,float,aggregate<float>> gdist(sz,domain,g,bc);

	gdist.template setBackgroundValue<0>(666);

	/////// GPU insert + flush

	Box<2,size_t> box({1,1},{1,1});
	auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());

	/////// GPU Run kernel

102
	gdist.setInsertBuffer(1);
103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122

	float c = 5.0;

	gdist.template iterateGridGPU<insert_kernel<0>>(it,c);
	gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);

	gdist.template deviceToHost<0>();

	{
		Box<2,size_t> box2({0,0},{15,15});

		auto it = gdist.getGridIterator(box2.getKP1(),box2.getKP2());

		while (it.isNext())
		{
			auto p = it.get_dist();
			auto p2 = it.get();

			if (p2.get(0) == box.getLow(0) && p2.get(1) == box.getLow(1))
			{
123
				BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 7.0);
124 125 126
			}
			else
			{
127 128 129 130 131 132
				if (gdist.template get<0>(p) != 666.0)
				{
					float f = gdist.template get<0>(p);
					std::cout << "ERROR: " << gdist.template get<0>(p) << std::endl;
				}

133 134 135 136 137 138 139
				BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 666.0);
			}

			++it;
		}
	}

140 141
	return;

142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182
	//

	c = 3.0;

	Box<2,size_t> box3({3,3},{11,11});

	auto it3 = gdist.getGridIterator(box3.getKP1(),box3.getKP2());
	gdist.setInsertBuffer(128);

	gdist.template iterateGridGPU<insert_kernel<0>>(it3,c);
	gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
	gdist.template deviceToHost<0>();

	{
		Box<2,size_t> box2({0,0},{15,15});

		auto it = gdist.getGridIterator(box2.getKP1(),box2.getKP2());

		while (it.isNext())
		{
			auto p = it.get_dist();
			auto p2 = it.get();

			Point<2,size_t> p2_ = p2.toPoint();

			if (box.isInside(p2_))
			{
				BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 5.0);
			}
			else if (box3.isInside(p2_))
			{
				BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 3.0);
			}
			else
			{
				BOOST_REQUIRE_EQUAL(gdist.template get<0>(p), 666.0);
			}

			++it;
		}
	}
183 184
}

185

186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201 202 203 204 205 206 207 208
BOOST_AUTO_TEST_CASE( sgrid_gpu_test_output )
{
	auto & v_cl = create_vcluster();

	if (v_cl.size() > 3){return;}

	size_t sz[2] = {17,17};
	periodicity<2> bc = {PERIODIC,PERIODIC};

	Ghost<2,long int> g(1);

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

	sgrid_dist_id_gpu<2,float,aggregate<float>> gdist(sz,domain,g,bc);

	gdist.template setBackgroundValue<0>(666);

	/////// GPU insert + flush

	Box<2,size_t> box({1,1},{15,15});
	auto it = gdist.getGridIterator(box.getKP1(),box.getKP2());

	/////// GPU Run kernel
209 210 211

	gdist.setInsertBuffer(128);

212
	float c = 5.0;
213

214 215
	gdist.template iterateGridGPU<insert_kernel<0>>(it,c);
	gdist.template flush<smax_<0>>(flush_type::FLUSH_ON_DEVICE);
216

217 218 219 220 221 222 223 224 225 226 227
	gdist.template deviceToHost<0>();

	gdist.write("sgrid_gpu_output");

	std::string file_test("sgrid_gpu_output_" + std::to_string(v_cl.size()) + "_" + std::to_string(v_cl.rank())  + ".vtk");
	std::string file("sgrid_gpu_output_" + std::to_string(v_cl.rank()) + ".vtk");

	bool test = compare(file,"test_data/" + file_test);

	BOOST_REQUIRE_EQUAL(true,test);
}
228 229

BOOST_AUTO_TEST_SUITE_END()