Commit a9a566f7 authored by Tomáš Oberhuber's avatar Tomáš Oberhuber Committed by Jakub Klinkovský
Browse files

ArrayTestHip can be compiled, HIP tests are failing.

parent 0841dab5

CMakeLists.txt

+2 −1
Original line number Diff line number Diff line
@@ -263,9 +263,10 @@ if( ${WITH_HIP} ) 
   if(HIP_FOUND)

      set(BUILD_HIP TRUE)

      set(CMAKE_HIPCXX_FLAGS ${CMAKE_HIPCXX_FLAGS} -I${HIP_ROOT_DIR}/include -DHAVE_HIP -DHIP_PLATFORM=nvcc )

      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -I${HIP_ROOT_DIR}/include -DHAVE_HIP -I/usr/lib/llvm-8/include/openmp -L/usr/lib/llvm-8/lib")

      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -I${HIP_ROOT_DIR}/include -DHAVE_HIP -I/usr/lib/llvm-8/include/openmp -L/usr/lib/llvm-8/lib -lomp -fopenmp")

      #set(CMAKE_SHARED_LINKER_FLAGS ${CMAKE_SHARED_LINKER_FLAGS_INIT} -L/usr/lib/llvm-8/lib} )

      set(CMAKE_CXX_COMPILER "${HIP_ROOT_DIR}/bin/hipcc" )

      set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -L/usr/lib/llvm-8/lib -lomp")

   endif()

endif()

src/Examples/Hamilton-Jacobi/Solvers/hamilton-jacobi/tnlDirectEikonalMethodBase1D_impl.h

+1 −0
Original line number Diff line number Diff line
@@ -89,6 +89,7 @@ template< typename Real, 
        typename Index >

template< typename MeshEntity >

void

__device_callable__

tnlDirectEikonalMethodsBase< Meshes::Grid< 1, Real, Device, Index > >::

updateCell( MeshFunctionType& u,

        const MeshEntity& cell,

src/TNL/Algorithms/MemoryOperationsHip.hpp

+3 −3
Original line number Diff line number Diff line
@@ -132,7 +132,7 @@ compare( const Element1* destination, 
   TNL_ASSERT_TRUE( source, "Attempted to compare data through a nullptr." );



   auto fetch = [=] __device_callable__ ( Index i ) -> bool { return destination[ i ] == source[ i ]; };

   return Reduction< Devices::Hip >::reduce( size, std::logical_and<>{}, fetch, true );

   return Reduction< Devices::Hip >::reduce( ( Index ) 0, size, fetch, std::logical_and<>{}, true );

}



template< typename Element,
@@ -148,7 +148,7 @@ containsValue( const Element* data, 
   TNL_ASSERT_GE( size, (Index) 0, "" );



   auto fetch = [=] __device_callable__ ( Index i ) -> bool { return data[ i ] == value; };

   return Reduction< Devices::Hip >::reduce( size, std::logical_or<>{}, fetch, false );

   return Reduction< Devices::Hip >::reduce( ( Index ) 0, size, fetch, std::logical_or<>{}, false );

}



template< typename Element,
@@ -164,7 +164,7 @@ containsOnlyValue( const Element* data, 
   TNL_ASSERT_GE( size, 0, "" );



   auto fetch = [=] __device_callable__ ( Index i ) -> bool { return data[ i ] == value; };

   return Reduction< Devices::Hip >::reduce( size, std::logical_and<>{}, fetch, true );

   return Reduction< Devices::Hip >::reduce( ( Index ) 0, size, fetch, std::logical_and<>{}, true );

}



} // namespace Algorithms

src/TNL/Algorithms/Segments/BiEllpack.hpp

+1 −0
Original line number Diff line number Diff line
@@ -113,6 +113,7 @@ template< typename Device, 
          typename IndexAllocator,

          ElementsOrganization Organization,

          int WarpSize >

__device_callable__

auto BiEllpack< Device, Index, IndexAllocator, Organization, WarpSize >::

getSegmentsCount() const -> IndexType

{

src/TNL/Assert.h

+8 −4
Original line number Diff line number Diff line
@@ -186,10 +186,7 @@ printDiagnosticsHost( const char* assertion, 
}

#endif // TNL_THROW_ASSERTION_ERROR



//__device_callable__

#ifdef HAVE_HIP

__device__

#endif

__device_callable__

inline void

printDiagnosticsCuda( const char* assertion,

                      const char* message,
@@ -215,6 +212,8 @@ fatalFailure() 
   // https://devtalk.nvidia.com/default/topic/509584/how-to-cancel-a-running-cuda-kernel-/

   // TODO: it is reported as "illegal instruction", but that leads to an abort as well...

   asm("trap;");

#elif defined __HIP_DEVICE_COMPILE__

   //asm("s_trap;"); // TODO HIP: Find how to cancel HIP kernel

#else

   throw EXIT_FAILURE;

#endif
@@ -261,6 +260,11 @@ cmpHelperOpFailure( const char* assertion, 
   // to construct the dynamic error message

   printDiagnosticsCuda( assertion, message, file, function, line,

                         "Not supported in CUDA kernels." );

#elif defined __HIP_DEVICE_COMPILE__

   // diagnostics is not supported - we don't have the machinery

   // to construct the dynamic error message

   printDiagnosticsCuda( assertion, message, file, function, line,

                         "Not supported in HIP kernels." );

#else

   const std::string formatted_lhs_value = Formatter< T1 >::printToString( lhs_value );

   const std::string formatted_rhs_value = Formatter< T2 >::printToString( rhs_value );