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

Improved CUDA support, still some problems remaining.

parent b71adb55

examples/hamilton-jacobi-parallel/tnlParallelEikonalSolver.h

+2 −1
Original line number Diff line number Diff line
@@ -112,6 +112,7 @@ public: 
	int*boundaryConditions_cuda;

	int* unusedCell_cuda;

	int* calculationsCount_cuda;

	double* tmpw;

	//MeshTypeCUDA mesh_cuda, subMesh_cuda;

	//SchemeDevice scheme_cuda;

	//double delta_cuda, tau0_cuda, stopTime_cuda,cflCondition_cuda;
@@ -156,7 +157,7 @@ template <typename SchemeHost, typename SchemeDevice, typename Device> 
__global__ void initRunCUDA(tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int >* caller);



template <typename SchemeHost, typename SchemeDevice, typename Device>

__global__ void initCUDA( tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int >* cudaSolver, double* ptr);

__global__ void initCUDA( tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int >* cudaSolver, double* ptr, bool * ptr2);



template <typename SchemeHost, typename SchemeDevice, typename Device>

__global__ void synchronizeCUDA(tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int >* cudaSolver);

examples/hamilton-jacobi-parallel/tnlParallelEikonalSolver_impl.h

+278 −127

File changed.

Preview size limit exceeded, changes collapsed.

src/operators/godunov-eikonal/parallelGodunovEikonal.h

+4 −4
Original line number Diff line number Diff line
@@ -149,11 +149,11 @@ public: 
                   const RealType& time,

                   const IndexType boundaryCondition ) const;



    template< typename Vector >



 #ifdef HAVE_CUDA

    __device__ __host__

    __device__

 #endif

    Real getValue( const MeshType& mesh,

    Real getValueDev( const MeshType& mesh,

                   const IndexType cellIndex,

                   const CoordinatesType& coordinates,

                   const RealType* u,

src/operators/godunov-eikonal/parallelGodunovEikonal2D_impl.h

+15 −10
Original line number Diff line number Diff line
@@ -145,6 +145,7 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 
          	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 const IndexType boundaryCondition ) const

{





	if ( ((coordinates.x() == 0 && (boundaryCondition & 4)) or

		 (coordinates.x() == mesh.getDimensions().x() - 1 && (boundaryCondition & 2)) or

		 (coordinates.y() == 0 && (boundaryCondition & 8)) or
@@ -165,6 +166,7 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 


	signui = sign(u[cellIndex],epsilon);





	if(fabs(u[cellIndex]) < acc) return 0.0;



	   if(signui > 0.0)
@@ -296,11 +298,11 @@ template< typename MeshReal, 
          typename MeshIndex,

          typename Real,

          typename Index >

template< typename Vector >



#ifdef HAVE_CUDA

__device__ __host__

__device__

#endif

Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Real, Index >:: getValue( const MeshType& mesh,

Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Real, Index >:: getValueDev( const MeshType& mesh,

          	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 const IndexType cellIndex,

          	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 const CoordinatesType& coordinates,

          	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 const Real* u,
@@ -308,6 +310,7 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 
          	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	  	 const IndexType boundaryCondition) const

{





	if ( ((coordinates.x() == 0 && (boundaryCondition & 4)) or

		 (coordinates.x() == mesh.getDimensions().x() - 1 && (boundaryCondition & 2)) or

		 (coordinates.y() == 0 && (boundaryCondition & 8)) or
@@ -327,8 +330,10 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 
	RealType nabla, xb, xf, yb, yf, signui;



	signui = sign(u[cellIndex],epsilon);



	if(fabs(u[cellIndex]) < acc) return 0.0;

#ifdef HAVE_CUDA

	//printf("%d   :    %d ;;;; %d   :   %d\n",threadIdx.x, coordinates.x(), threadIdx.y,coordinates.y());

#endif

	//if(fabs(u[cellIndex]) < acc) return 0.0;



	   if(signui > 0.0)

	   {
@@ -379,8 +384,8 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 
			   yb = 0.0;



		   nabla = sqrt (xf*xf + xb*xb + yf*yf + yb*yb );

		   if(fabs(1.0-nabla) < acc)

			   return 0.0;

		  // if(fabs(1.0-nabla) < acc)

			//   return 0.0;

		   return signui*(1.0 - nabla);

	   }

	   else if (signui < 0.0)
@@ -435,8 +440,8 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 


		   nabla = sqrt (xf*xf + xb*xb + yf*yf + yb*yb );



		   if(fabs(1.0-nabla) < acc)

			   return 0.0;

		  // if(fabs(1.0-nabla) < acc)

		//	   return 0.0;

		   return signui*(1.0 - nabla);

	   }

	   else