Loading examples/fast-sweeping/main.h +12 −12 Original line number Diff line number Diff line Loading @@ -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/tnlBuildConfigTags.h> Loading Loading @@ -60,16 +60,16 @@ int main( int argc, char* argv[] ) } else if(dim == "3") { // tnlFastSweeping<tnlGrid<3,double,tnlHost, int>, double, int> solver; // if(!solver.init(parameters)) // { // cerr << "Solver failed to initialize." << endl; // return EXIT_FAILURE; // } // checkCudaDevice; // cout << "-------------------------------------------------------------" << endl; // cout << "Starting solver..." << endl; // solver.run(); tnlFastSweeping<tnlGrid<3,double,tnlHost, int>, double, int> solver; if(!solver.init(parameters)) { cerr << "Solver failed to initialize." << endl; return EXIT_FAILURE; } checkCudaDevice; cout << "-------------------------------------------------------------" << endl; cout << "Starting solver..." << endl; solver.run(); } else { Loading examples/fast-sweeping/tnlFastSweeping2D_CUDA_v4_impl.h +6 −8 Original line number Diff line number Diff line Loading @@ -119,7 +119,7 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: dim3 threadsPerBlock(16, 16); dim3 numBlocks(n/16 + 1 ,n/16 +1); // setEntityGridCUDA<<<dim3(1,1),dim3(1,1)>>>(this->cudaSolver); initCUDA<<<numBlocks,threadsPerBlock>>>(this->cudaSolver); cudaDeviceSynchronize(); checkCudaDevice; Loading Loading @@ -149,13 +149,15 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: cudaDeviceSynchronize(); checkCudaDevice; data.setLike(dofVector.getData()); cudaMemcpy(data.getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); //data.setLike(dofVector.getData()); //cudaMemcpy(data.getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); cudaMemcpy(dofVector.getData().getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); cudaFree(cudaDofVector); cudaFree(cudaDofVector2); cudaFree(cudaSolver); data.save("u-00001.tnl"); //data.save("u-00001.tnl"); dofVector.save("u-00001.tnl"); cudaDeviceSynchronize(); return true; } Loading @@ -174,13 +176,9 @@ template< typename MeshReal, __device__ void tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: updateValue( Index i, Index j) { if(i >= Mesh.getDimensions().x() || j >= Mesh.getDimensions().y() || i<0 || j<0 ) printf("i = %d, j = %d",i,j); tnlGridEntity< tnlGrid< 2,double, tnlHost, int >, 2, tnlGridEntityNoStencilStorage > Entity(Mesh); Entity.setCoordinates(CoordinatesType(i,j)); Entity.refresh(); // printf("index: %d\n",Entity.getIndex()); // neighbourEntities.refresh(Mesh,Entity.getIndex()); tnlNeighbourGridEntityGetter<tnlGridEntity< MeshType, 2, tnlGridEntityNoStencilStorage >,2> neighbourEntities(Entity); Real value = cudaDofVector2[Entity.getIndex()]; Real a,b, tmp; Loading examples/fast-sweeping/tnlFastSweeping2D_impl.h +8 −8 Original line number Diff line number Diff line Loading @@ -318,10 +318,10 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: // // dofVector[Mesh.getCellIndex(CoordinatesType(i,j))] = tmp*INT_MAX; data.setLike(dofVector2.getData()); data=dofVector2.getData(); cout << data.getType() << endl; data.save("u-00000.tnl"); //data.setLike(dofVector2.getData()); //data=dofVector2.getData(); //cout << data.getType() << endl; dofVector2.save("u-00000.tnl"); //dofVector2.getData().save("u-00000.tnl"); return true; Loading Loading @@ -377,10 +377,10 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: /*---------------------------------------------------------------------------------------------------------------------------*/ data.setLike(dofVector2.getData()); data = dofVector2.getData(); cout << data.getType() << endl; data.save("u-00001.tnl"); // data.setLike(dofVector2.getData()); // data = dofVector2.getData(); // cout << data.getType() << endl; dofVector2.save("u-00001.tnl"); //dofVector2.getData().save("u-00001.tnl"); return true; Loading examples/fast-sweeping/tnlFastSweeping3D_CUDA_impl.h +464 −458 File changed.Preview size limit exceeded, changes collapsed. Show changes examples/fast-sweeping/tnlFastSweeping_CUDA.h +2 −4 Original line number Diff line number Diff line Loading @@ -158,7 +158,7 @@ protected: bool exactInput; DofVectorType dofVector; tnlMeshFunction<MeshType> dofVector; DofVectorType data; RealType h; Loading @@ -179,8 +179,6 @@ __global__ void runCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, doub __global__ void initCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver); __global__ void initCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, double, int >* solver); __global__ void setEntityGridCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver); #endif /*various implementtions.... choose one*/ Loading @@ -191,6 +189,6 @@ __global__ void setEntityGridCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, i //#include "tnlFastSweeping2D_CUDA_v5_impl.h" // #include "tnlFastSweeping3D_CUDA_impl.h" #include "tnlFastSweeping3D_CUDA_impl.h" #endif /* TNLFASTSWEEPING_H_ */ Loading
examples/fast-sweeping/main.h +12 −12 Original line number Diff line number Diff line Loading @@ -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/tnlBuildConfigTags.h> Loading Loading @@ -60,16 +60,16 @@ int main( int argc, char* argv[] ) } else if(dim == "3") { // tnlFastSweeping<tnlGrid<3,double,tnlHost, int>, double, int> solver; // if(!solver.init(parameters)) // { // cerr << "Solver failed to initialize." << endl; // return EXIT_FAILURE; // } // checkCudaDevice; // cout << "-------------------------------------------------------------" << endl; // cout << "Starting solver..." << endl; // solver.run(); tnlFastSweeping<tnlGrid<3,double,tnlHost, int>, double, int> solver; if(!solver.init(parameters)) { cerr << "Solver failed to initialize." << endl; return EXIT_FAILURE; } checkCudaDevice; cout << "-------------------------------------------------------------" << endl; cout << "Starting solver..." << endl; solver.run(); } else { Loading
examples/fast-sweeping/tnlFastSweeping2D_CUDA_v4_impl.h +6 −8 Original line number Diff line number Diff line Loading @@ -119,7 +119,7 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: dim3 threadsPerBlock(16, 16); dim3 numBlocks(n/16 + 1 ,n/16 +1); // setEntityGridCUDA<<<dim3(1,1),dim3(1,1)>>>(this->cudaSolver); initCUDA<<<numBlocks,threadsPerBlock>>>(this->cudaSolver); cudaDeviceSynchronize(); checkCudaDevice; Loading Loading @@ -149,13 +149,15 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: cudaDeviceSynchronize(); checkCudaDevice; data.setLike(dofVector.getData()); cudaMemcpy(data.getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); //data.setLike(dofVector.getData()); //cudaMemcpy(data.getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); cudaMemcpy(dofVector.getData().getData(), cudaDofVector2, this->dofVector.getData().getSize()*sizeof(double), cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); cudaFree(cudaDofVector); cudaFree(cudaDofVector2); cudaFree(cudaSolver); data.save("u-00001.tnl"); //data.save("u-00001.tnl"); dofVector.save("u-00001.tnl"); cudaDeviceSynchronize(); return true; } Loading @@ -174,13 +176,9 @@ template< typename MeshReal, __device__ void tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: updateValue( Index i, Index j) { if(i >= Mesh.getDimensions().x() || j >= Mesh.getDimensions().y() || i<0 || j<0 ) printf("i = %d, j = %d",i,j); tnlGridEntity< tnlGrid< 2,double, tnlHost, int >, 2, tnlGridEntityNoStencilStorage > Entity(Mesh); Entity.setCoordinates(CoordinatesType(i,j)); Entity.refresh(); // printf("index: %d\n",Entity.getIndex()); // neighbourEntities.refresh(Mesh,Entity.getIndex()); tnlNeighbourGridEntityGetter<tnlGridEntity< MeshType, 2, tnlGridEntityNoStencilStorage >,2> neighbourEntities(Entity); Real value = cudaDofVector2[Entity.getIndex()]; Real a,b, tmp; Loading
examples/fast-sweeping/tnlFastSweeping2D_impl.h +8 −8 Original line number Diff line number Diff line Loading @@ -318,10 +318,10 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: // // dofVector[Mesh.getCellIndex(CoordinatesType(i,j))] = tmp*INT_MAX; data.setLike(dofVector2.getData()); data=dofVector2.getData(); cout << data.getType() << endl; data.save("u-00000.tnl"); //data.setLike(dofVector2.getData()); //data=dofVector2.getData(); //cout << data.getType() << endl; dofVector2.save("u-00000.tnl"); //dofVector2.getData().save("u-00000.tnl"); return true; Loading Loading @@ -377,10 +377,10 @@ bool tnlFastSweeping< tnlGrid< 2,MeshReal, Device, MeshIndex >, Real, Index > :: /*---------------------------------------------------------------------------------------------------------------------------*/ data.setLike(dofVector2.getData()); data = dofVector2.getData(); cout << data.getType() << endl; data.save("u-00001.tnl"); // data.setLike(dofVector2.getData()); // data = dofVector2.getData(); // cout << data.getType() << endl; dofVector2.save("u-00001.tnl"); //dofVector2.getData().save("u-00001.tnl"); return true; Loading
examples/fast-sweeping/tnlFastSweeping3D_CUDA_impl.h +464 −458 File changed.Preview size limit exceeded, changes collapsed. Show changes
examples/fast-sweeping/tnlFastSweeping_CUDA.h +2 −4 Original line number Diff line number Diff line Loading @@ -158,7 +158,7 @@ protected: bool exactInput; DofVectorType dofVector; tnlMeshFunction<MeshType> dofVector; DofVectorType data; RealType h; Loading @@ -179,8 +179,6 @@ __global__ void runCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, doub __global__ void initCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver); __global__ void initCUDA(tnlFastSweeping< tnlGrid< 3,double, tnlHost, int >, double, int >* solver); __global__ void setEntityGridCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, int >, double, int >* solver); #endif /*various implementtions.... choose one*/ Loading @@ -191,6 +189,6 @@ __global__ void setEntityGridCUDA(tnlFastSweeping< tnlGrid< 2,double, tnlHost, i //#include "tnlFastSweeping2D_CUDA_v5_impl.h" // #include "tnlFastSweeping3D_CUDA_impl.h" #include "tnlFastSweeping3D_CUDA_impl.h" #endif /* TNLFASTSWEEPING_H_ */