Commit 9609830d authored by Tomáš Oberhuber's avatar Tomáš Oberhuber
Browse files

Fixes after merge.

parent d001c75e

src/TNL/Experimental/Hamilton-Jacobi/Solvers/hamilton-jacobi/tnlDirectEikonalMethodsBase_impl.h

+6 −6
Original line number Diff line number Diff line
@@ -716,8 +716,8 @@ __global__ void CudaInitCaller( const Functions::MeshFunction< Meshes::Grid< 1, 




        output[ cind ] =

               input( cell ) >= 0 ? std::numeric_limits< RealType >::max() :

                                    - std::numeric_limits< RealType >::max();

               input( cell ) >= 0 ? std::numeric_limits< Real >::max() :

                                    - std::numeric_limits< Real >::max();

        interfaceMap[ cind ] = false; 



        const Real& h = mesh.getSpaceSteps().x();
@@ -768,8 +768,8 @@ __global__ void CudaInitCaller( const Functions::MeshFunction< Meshes::Grid< 2, 




        output[ cind ] =

               input( cell ) >= 0 ? std::numeric_limits< RealType >::max() :

                                    - std::numeric_limits< RealType >::max();

               input( cell ) >= 0 ? std::numeric_limits< Real >::max() :

                                    - std::numeric_limits< Real >::max();

        interfaceMap[ cind ] = false; 



        const Real& hx = mesh.getSpaceSteps().x();
@@ -841,8 +841,8 @@ __global__ void CudaInitCaller3d( const Functions::MeshFunction< Meshes::Grid< 3 




        output[ cind ] =

               input( cell ) >= 0 ? std::numeric_limits< RealType >::max() :

                                    - std::numeric_limits< RealType >::max();

               input( cell ) >= 0 ? std::numeric_limits< Real >::max() :

                                    - std::numeric_limits< Real >::max();

        interfaceMap[ cind ] = false; 

        cell.refresh();

src/TNL/Experimental/Hamilton-Jacobi/Solvers/hamilton-jacobi/tnlFastSweepingMethod1D_impl.h

+3 −3
Original line number Diff line number Diff line
@@ -170,7 +170,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
    }

    

    __shared__ volatile Real sArray[ 18 ];

    sArray[thri] = TypeInfo< Real >::getMaxValue();

    sArray[thri] = std::numeric_limits<  Real >::max();

    

    //filling sArray edges

    int dimX = mesh.getDimensions().x();
@@ -191,7 +191,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimX > (blIdx+1) * blockDim.x )

            sArray[xkolik] = aux[ blIdx*blockDim.x - 1 + xkolik ];

        else

            sArray[xkolik] = TypeInfo< Real >::getMaxValue();

            sArray[xkolik] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 1 )
@@ -199,7 +199,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( blIdx != 0 )

            sArray[0] = aux[ blIdx*blockDim.x - 1 ];

        else

            sArray[0] = TypeInfo< Real >::getMaxValue();

            sArray[0] = std::numeric_limits< Real >::max();

    }

src/TNL/Experimental/Hamilton-Jacobi/Solvers/hamilton-jacobi/tnlFastSweepingMethod2D_impl.h

+5 −5
Original line number Diff line number Diff line
@@ -393,7 +393,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
    

    //__shared__ volatile Real sArray[ blockDim.y+2 ][ blockDim.x+2 ];

    __shared__ volatile Real sArray[18][18];

    sArray[thrj][thri] = TypeInfo< Real >::getMaxValue();

    sArray[thrj][thri] = std::numeric_limits< Real >::max();

    

    //filling sArray edges

    int dimX = mesh.getDimensions().x(); int dimY = mesh.getDimensions().y();
@@ -422,7 +422,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimX > (blIdx+1) * blockDim.x  && thrj+1 < ykolik )

            sArray[thrj+1][xkolik] = aux[ blIdy*blockDim.y*dimX - dimX + blIdx*blockDim.x - 1 + (thrj+1)*dimX + xkolik ];

        else

            sArray[thrj+1][xkolik] = TypeInfo< Real >::getMaxValue();

            sArray[thrj+1][xkolik] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 1 )
@@ -430,7 +430,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( blIdx != 0 && thrj+1 < ykolik )

            sArray[thrj+1][0] = aux[ blIdy*blockDim.y*dimX - dimX + blIdx*blockDim.x - 1 + (thrj+1)*dimX ];

        else

            sArray[thrj+1][0] = TypeInfo< Real >::getMaxValue();

            sArray[thrj+1][0] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 2 )
@@ -438,7 +438,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimY > (blIdy+1) * blockDim.y  && thri+1 < xkolik )

            sArray[ykolik][thrj+1] = aux[ blIdy*blockDim.y*dimX - dimX + blIdx*blockDim.x - 1 + ykolik*dimX + thrj+1 ];

        else

           sArray[ykolik][thrj+1] = TypeInfo< Real >::getMaxValue();

           sArray[ykolik][thrj+1] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 3 )
@@ -446,7 +446,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( blIdy != 0 && thrj+1 < xkolik )

            sArray[0][thrj+1] = aux[ blIdy*blockDim.y*dimX - dimX + blIdx*blockDim.x - 1 + thrj+1 ];

        else

            sArray[0][thrj+1] = TypeInfo< Real >::getMaxValue();

            sArray[0][thrj+1] = std::numeric_limits< Real >::max();

    }

src/TNL/Experimental/Hamilton-Jacobi/Solvers/hamilton-jacobi/tnlFastSweepingMethod3D_impl.h

+13 −13
Original line number Diff line number Diff line
@@ -328,15 +328,15 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        hz = mesh.getSpaceSteps().z();

    }

    __shared__ volatile Real sArray[10][10][10];

    sArray[thrk][thrj][thri] = TypeInfo< Real >::getMaxValue();

    sArray[thrk][thrj][thri] = std::numeric_limits< Real >::max();

    if(thri == 0 )

    {

        sArray[8][thrj+1][thrk+1] = TypeInfo< Real >::getMaxValue();

        sArray[9][thrj+1][thrk+1] = TypeInfo< Real >::getMaxValue();

        sArray[thrk+1][thrj+1][8] = TypeInfo< Real >::getMaxValue();

        sArray[thrk+1][thrj+1][9] = TypeInfo< Real >::getMaxValue();

        sArray[thrj+1][8][thrk+1] = TypeInfo< Real >::getMaxValue();

        sArray[thrj+1][9][thrk+1] = TypeInfo< Real >::getMaxValue();

        sArray[8][thrj+1][thrk+1] = std::numeric_limits< Real >::max();

        sArray[9][thrj+1][thrk+1] = std::numeric_limits< Real >::max();

        sArray[thrk+1][thrj+1][8] = std::numeric_limits< Real >::max();

        sArray[thrk+1][thrj+1][9] = std::numeric_limits< Real >::max();

        sArray[thrj+1][8][thrk+1] = std::numeric_limits< Real >::max();

        sArray[thrj+1][9][thrk+1] = std::numeric_limits< Real >::max();

    }

            

    //filling sArray edges
@@ -374,7 +374,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( blIdx != 0 && thrj+1 < ykolik && thrk+1 < zkolik )

            sArray[thrk+1][thrj+1][0] = aux[ blIdz*blockDim.z * dimX * dimY + blIdy * blockDim.y*dimX + blIdx*blockDim.x + thrj * dimX -1 + thrk*dimX*dimY ];

        else

            sArray[thrk+1][thrj+1][0] = TypeInfo< Real >::getMaxValue();

            sArray[thrk+1][thrj+1][0] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 1 )
@@ -382,14 +382,14 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimX > (blIdx+1) * blockDim.x && thrj+1 < ykolik && thrk+1 < zkolik )

            sArray[thrk+1][thrj+1][9] = aux[ blIdz*blockDim.z * dimX * dimY + blIdy *blockDim.y*dimX+ blIdx*blockDim.x + blockDim.x + thrj * dimX + thrk*dimX*dimY ];

        else

            sArray[thrk+1][thrj+1][9] = TypeInfo< Real >::getMaxValue();

            sArray[thrk+1][thrj+1][9] = std::numeric_limits< Real >::max();

    }

    if( thri == 2 )

    {        

        if( blIdy != 0 && thrj+1 < xkolik && thrk+1 < zkolik )

            sArray[thrk+1][0][thrj+1] = aux[ blIdz*blockDim.z * dimX * dimY + blIdy * blockDim.y*dimX + blIdx*blockDim.x - dimX + thrj + thrk*dimX*dimY ];

        else

            sArray[thrk+1][0][thrj+1] = TypeInfo< Real >::getMaxValue();

            sArray[thrk+1][0][thrj+1] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 3 )
@@ -397,14 +397,14 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimY > (blIdy+1) * blockDim.y && thrj+1 < xkolik && thrk+1 < zkolik )

            sArray[thrk+1][9][thrj+1] = aux[ blIdz*blockDim.z * dimX * dimY + (blIdy+1) * blockDim.y*dimX + blIdx*blockDim.x + thrj + thrk*dimX*dimY ];

        else

            sArray[thrk+1][9][thrj+1] = TypeInfo< Real >::getMaxValue();

            sArray[thrk+1][9][thrj+1] = std::numeric_limits< Real >::max();

    }

    if( thri == 4 )

    {        

        if( blIdz != 0 && thrj+1 < ykolik && thrk+1 < xkolik )

            sArray[0][thrj+1][thrk+1] = aux[ blIdz*blockDim.z * dimX * dimY + blIdy * blockDim.y*dimX + blIdx*blockDim.x - dimX * dimY + thrj * dimX + thrk ];

        else

            sArray[0][thrj+1][thrk+1] = TypeInfo< Real >::getMaxValue();

            sArray[0][thrj+1][thrk+1] = std::numeric_limits< Real >::max();

    }

    

    if( thri == 5 )
@@ -412,7 +412,7 @@ __global__ void CudaUpdateCellCaller( tnlDirectEikonalMethodsBase< Meshes::Grid< 
        if( dimZ > (blIdz+1) * blockDim.z && thrj+1 < ykolik && thrk+1 < xkolik )

            sArray[9][thrj+1][thrk+1] = aux[ (blIdz+1)*blockDim.z * dimX * dimY + blIdy * blockDim.y*dimX + blIdx*blockDim.x + thrj * dimX + thrk ];

        else

            sArray[9][thrj+1][thrk+1] = TypeInfo< Real >::getMaxValue();

            sArray[9][thrj+1][thrk+1] = std::numeric_limits< Real >::max();

    }

    

    if( i < mesh.getDimensions().x() && j < mesh.getDimensions().y() && k < mesh.getDimensions().z() )

src/TNL/Math.h

+6 −1
Original line number Diff line number Diff line
@@ -72,7 +72,12 @@ template< class T > 
__cuda_callable__ inline

T abs( const T& n )

{

#if defined(__MIC__)

#if defined(__CUDA_ARCH__)

   if( std::is_integral< T >::value )

      return ::abs( n );

   else

      return ::fabs( n );

#elif defined(__MIC__)

   if( n < ( T ) 0 )

      return -n;

   return n;