Commit 7810dfdd authored by Tomas Sobotik's avatar Tomas Sobotik
Browse files

Functional CUDA version, not yet optimized.

parent f8ac5799

examples/hamilton-jacobi-parallel/tnlParallelEikonalSolver_impl.h

+45 −58
Original line number Diff line number Diff line
@@ -1474,7 +1474,7 @@ void /*tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int>:: 
	//printf("hurewrwr %f \n", u[l]);

	if(u[0] * u[l] <= 0.0)

	{

		printf("contains %d \n",i);

		//printf("contains %d \n",i);

		atomicMax( &containsCurve, 1);

	}
@@ -1508,97 +1508,87 @@ void /*tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int>:: 
	int i = blockIdx.y * gridDim.x + blockIdx.x;

	int l = threadIdx.y * blockDim.x + threadIdx.x;



	if(i+l == 0)

		printf("a");

	if(caller->getSubgridValueCUDA(i) != INT_MAX)

	{

		double a;

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

		caller->getSubgridCUDA(i,caller, &u[l]);

		int bound = caller->getBoundaryConditionCUDA(i);

		if(l == 0)

		//if(l == 0)

			//printf("i = %d, bound = %d\n",i,caller->getSubgridValueCUDA(i));

		if(bound & 1)

		{

			caller->runSubgridCUDA(1,u,i);

			//this->calculationsCount[i]++;

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if(bound & 2)

		{

			caller->runSubgridCUDA(2,u,i);

			//this->calculationsCount[i]++;

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if(bound & 4)

		{

			caller->runSubgridCUDA(4,u,i);

			//this->calculationsCount[i]++;

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if(bound & 8)

		{

			caller->runSubgridCUDA(8,u,i);

			//this->calculationsCount[i]++;

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}







		if( ((bound & 2) ))

		{

			caller->runSubgridCUDA(3,u,i);

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if( ((bound & 4) ))

		{

			caller->runSubgridCUDA(5,u,i);

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if( ((bound & 2) ))

		{

			caller->runSubgridCUDA(10,u,i);

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

		caller->getSubgridCUDA(i,caller, &a);

		u[l] = a;

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}

		if(   (bound & 4) )

		{

			caller->runSubgridCUDA(12,u,i);

		}

			__syncthreads();

			caller->insertSubgridCUDA(u[l],i);

			__syncthreads();

			caller->getSubgridCUDA(i,caller, &u[l]);

			__syncthreads();

		}





		caller->setBoundaryConditionCUDA(i, 0);
@@ -1607,9 +1597,6 @@ void /*tnlParallelEikonalSolver<SchemeHost, SchemeDevice, Device, double, int>:: 


	}





	if(i+l == 0)

		printf("b");

}



#endif /*HAVE_CUDA*/

src/operators/godunov-eikonal/parallelGodunovEikonal2D_impl.h

+3 −3
Original line number Diff line number Diff line
@@ -331,7 +331,7 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 


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

#ifdef HAVE_CUDA

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

	//printf("%d   :    %d ;;;; %d   :   %d  , %f \n",threadIdx.x, mesh.getDimensions().x() , threadIdx.y,mesh.getDimensions().y(), epsilon );

#endif

	//if(fabs(u[cellIndex]) < acc) return 0.0;
@@ -351,7 +351,7 @@ Real parallelGodunovEikonalScheme< tnlGrid< 2, MeshReal, Device, MeshIndex >, Re 
		   else *//*if(boundaryCondition & 2)

			   xb = 0.0;

		   else /**/if(coordinates.x() == 0)

			   xb = positivePart((u[cellIndex] - u[mesh.template getCellNextToCell<+1,0>( cellIndex )])/hx);

			   xb = positivePart((u[cellIndex] - u[mesh.template getCellNextToCell<1,0>( cellIndex )])/hx);

		   else

			   xb = positivePart((u[cellIndex] - u[mesh.template getCellNextToCell<-1,0>( cellIndex )])/hx);