Commit ef5cc17a authored by Tomas Sobotik's avatar Tomas Sobotik
Browse files

Working 3D Fast Sweep, Refactor of parrallel Jacobi-Hamilton

parent e1dbcfbf

examples/fast-sweeping/main.h

+3 −3
Original line number Diff line number Diff line
@@ -17,9 +17,9 @@ 


#include "MainBuildConfig.h"

	//for HOST versions:

//#include "tnlFastSweeping.h"

#include "tnlFastSweeping.h"

	//for DEVICE versions:

#include "tnlFastSweeping_CUDA.h"

//#include "tnlFastSweeping_CUDA.h"

#include "fastSweepingConfig.h"

#include <solvers/tnlConfigTags.h>
@@ -43,7 +43,7 @@ int main( int argc, char* argv[] ) 
   if( ! parseCommandLine( argc, argv, configDescription, parameters ) )

      return false;



    tnlFastSweeping<tnlGrid<2,double,tnlHost, int>, double, int> solver;

    tnlFastSweeping<tnlGrid<3,double,tnlHost, int>, double, int> solver;

    if(!solver.init(parameters))

   {

    	cerr << "Solver failed to initialize." << endl;

examples/fast-sweeping/tnlFastSweeping.h

+78 −0
Original line number Diff line number Diff line
@@ -107,9 +107,87 @@ protected: 


};



















template< typename MeshReal,

          typename Device,

          typename MeshIndex,

          typename Real,

          typename Index >

class tnlFastSweeping< tnlGrid< 3,MeshReal, Device, MeshIndex >, Real, Index >

{



public:

	typedef Real RealType;

	typedef Device DeviceType;

	typedef Index IndexType;

	typedef tnlGrid< 3, Real, Device, Index > MeshType;

	typedef tnlVector< RealType, DeviceType, IndexType> DofVectorType;

	typedef typename MeshType::CoordinatesType CoordinatesType;







	static tnlString getType();

	bool init( const tnlParameterContainer& parameters );



	bool initGrid();

	bool run();



	//for single core version use this implementation:

	void updateValue(const Index i, const Index j, const Index k);

	//for parallel version use this one instead:

//	void updateValue(const Index i, const Index j, DofVectorType* grid);





	void setupSquare1000(Index i, Index j);

	void setupSquare1100(Index i, Index j);

	void setupSquare1010(Index i, Index j);

	void setupSquare1001(Index i, Index j);

	void setupSquare1110(Index i, Index j);

	void setupSquare1101(Index i, Index j);

	void setupSquare1011(Index i, Index j);

	void setupSquare1111(Index i, Index j);

	void setupSquare0000(Index i, Index j);

	void setupSquare0100(Index i, Index j);

	void setupSquare0010(Index i, Index j);

	void setupSquare0001(Index i, Index j);

	void setupSquare0110(Index i, Index j);

	void setupSquare0101(Index i, Index j);

	void setupSquare0011(Index i, Index j);

	void setupSquare0111(Index i, Index j);



	Real fabsMin(const Real x, const Real y);





protected:



	MeshType Mesh;



	bool exactInput;



	DofVectorType dofVector, dofVector2;



	RealType h;



#ifdef HAVE_OPENMP

//	omp_lock_t* gridLock;

#endif





};





	//for single core version use this implementation:

#include "tnlFastSweeping2D_impl.h"

	//for parallel version use this one instead:

// #include "tnlFastSweeping2D_openMP_impl.h"



#include "tnlFastSweeping3D_impl.h"



#endif /* TNLFASTSWEEPING_H_ */

examples/fast-sweeping/tnlFastSweeping3D_CUDA_impl.h

0 → 100644
+1253 −0

File added.

Preview size limit exceeded, changes collapsed.

examples/fast-sweeping/tnlFastSweeping3D_impl.h

0 → 100644
+946 −0

File added.

Preview size limit exceeded, changes collapsed.

examples/fast-sweeping/tnlFastSweeping_CUDA.h

+80 −1
Original line number Diff line number Diff line
@@ -105,12 +105,90 @@ protected: 


















template< typename MeshReal,

          typename Device,

          typename MeshIndex,

          typename Real,

          typename Index >

class tnlFastSweeping< tnlGrid< 3,MeshReal, Device, MeshIndex >, Real, Index >

{



public:

	typedef Real RealType;

	typedef Device DeviceType;

	typedef Index IndexType;

	typedef tnlGrid< 3, Real, Device, Index > MeshType;

	typedef tnlVector< RealType, DeviceType, IndexType> DofVectorType;

	typedef typename MeshType::CoordinatesType CoordinatesType;







	__host__ static tnlString getType();

	__host__ bool init( const tnlParameterContainer& parameters );

	__host__ bool run();



#ifdef HAVE_CUDA

	__device__ bool initGrid();

	__device__ void updateValue(const Index i, const Index j, const Index k);

	__device__ void updateValue(const Index i, const Index j, const Index k, double** sharedMem, const int k3);

	__device__ Real fabsMin(const Real x, const Real y);



	tnlFastSweeping< tnlGrid< 3,MeshReal, Device, MeshIndex >, Real, Index >* cudaSolver;

	double* cudaDofVector;

	double* cudaDofVector2;

	int counter;

	__device__ void setupSquare1000(Index i, Index j);

	__device__ void setupSquare1100(Index i, Index j);

	__device__ void setupSquare1010(Index i, Index j);

	__device__ void setupSquare1001(Index i, Index j);

	__device__ void setupSquare1110(Index i, Index j);

	__device__ void setupSquare1101(Index i, Index j);

	__device__ void setupSquare1011(Index i, Index j);

	__device__ void setupSquare1111(Index i, Index j);

	__device__ void setupSquare0000(Index i, Index j);

	__device__ void setupSquare0100(Index i, Index j);

	__device__ void setupSquare0010(Index i, Index j);

	__device__ void setupSquare0001(Index i, Index j);

	__device__ void setupSquare0110(Index i, Index j);

	__device__ void setupSquare0101(Index i, Index j);

	__device__ void setupSquare0011(Index i, Index j);

	__device__ void setupSquare0111(Index i, Index j);

#endif



	MeshType Mesh;



protected:







	bool exactInput;



	DofVectorType dofVector;



	RealType h;





};















#ifdef HAVE_CUDA

//template<int sweep_t>

__global__ void runCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver, int sweep, int i);

//__global__ void runCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver, int sweep, int i);

__global__ void runCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, double, int >* solver, int sweep, int i);



__global__ void initCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver);

__global__ void initCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, double, int >* solver);

#endif



/*various implementtions.... choose one*/
@@ -119,5 +197,6 @@ __global__ void initCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, dou 
//#include "tnlFastSweeping2D_CUDA_v3_impl.h"

#include "tnlFastSweeping2D_CUDA_v4_impl.h"

//#include "tnlFastSweeping2D_CUDA_v5_impl.h"

#include "tnlFastSweeping3D_CUDA_impl.h"



#endif /* TNLFASTSWEEPING_H_ */