fix: HIP build

src/coreComponents/common/LifoStorageCommon.hpp

@@ -164,7 +164,7 @@ class LifoStorageCommon
     GEOS_ERROR( "To use LifoStorage, both _SC_AVPHYS_PAGES and _SC_PAGESIZE must be defined." );
 #endif
     int numberOfBuffersToStoreOnHost = std::max( 1, std::min( ( int )( 0.01 * percent * free / bufferSize ), maxNumberOfBuffers - numberOfBuffersToStoreOnDevice ) );
-    double freeGB = ( ( double ) free ) / ( 1024.0 * 1024.0 * 1024.0 ) / MpiWrapper::nodeCommSize();
+    GEOS_MAYBE_UNUSED double freeGB = ( ( double ) free ) / ( 1024.0 * 1024.0 * 1024.0 ) / MpiWrapper::nodeCommSize();
     LIFO_LOG_RANK( " LIFO : available memory on host " << freeGB << " GB" );
     return numberOfBuffersToStoreOnHost;
   }

src/coreComponents/common/TypeDispatch.hpp

@@ -376,7 +376,7 @@ bool dispatch( LIST const combinations,
   if( !success )
   {
     auto typePrinter = []( auto t ){ return LvArray::system::demangle( typeid( typename decltype(t)::type ).name() ); };
-    auto typeListPrinter = [typePrinter]( auto tlist ){ return internal::listToString( typename decltype( tlist )::type{}, "\n  ", "", typePrinter ); };
+    GEOS_MAYBE_UNUSED auto typeListPrinter = [typePrinter]( auto tlist ){ return internal::listToString( typename decltype( tlist )::type{}, "\n  ", "", typePrinter ); };
 
     GEOS_ERROR( "Types were not dispatched to the lambda of type\n"
                 << LvArray::system::demangleType< LAMBDA >() << "\n"

src/coreComponents/common/format/table/TableFormatter.cpp

@@ -98,6 +98,8 @@ void toStream( std::ostream & outputStream, string_view content, string_view str
   {
     GEOS_WARNING( GEOS_FMT( "Error while writing to '{}':\n{}", streamName, msgs ) );
   }
+  // TODO: remove after implementing GEOS_ERROR for HIP
+  GEOS_UNUSED_VAR(streamName);
 }
 
 void TableFormatter::toStreamImpl( std::ostream & outputStream, string_view content ) const

src/coreComponents/common/logger/Logger.hpp

@@ -201,6 +201,27 @@
       asm ( "trap;" ); \
     } \
   } while( false )
+#elif __HIP_DEVICE_COMPILE__
+#define GEOS_ERROR_IF_CAUSE( COND, CAUSE_MESSAGE, ... ) \
+  do \
+  { \
+    if( COND ) \
+    { \
+      constexpr char const * formatString = "***** ERROR\n" \
+                                            "***** LOCATION" LOCATION "\n" \
+                                            "***** BLOCK:  [%u, %u, %u]\n" \
+                                            "***** THREAD: [%u, %u, %u]\n" \
+                                            "***** %s\n" \
+                                            "***** %s\n\n"; \
+      printf( formatString, \
+              blockIdx.x, blockIdx.y, blockIdx.z, \
+              threadIdx.x, threadIdx.y, threadIdx.z, \
+              STRINGIZE( CAUSE_MESSAGE ), \
+              STRINGIZE( GEOS_DETAIL_FIRST_ARG( __VA_ARGS__ ) ) \
+            ); \
+      asm volatile ( "s_trap 2" ); \
+    } \
+  } while( false )
 #endif
 
 /**
@@ -283,6 +304,27 @@
       asm ( "trap;" ); \
     } \
   } while( false )
+#elif __HIP_DEVICE_COMPILE__
+#define GEOS_THROW_IF_CAUSE( COND, CAUSE_MESSAGE, MSG, ... ) \
+  do \
+  { \
+    if( COND ) \
+    { \
+      static char const formatString[] = "***** ERROR\n" \
+                                         "***** LOCATION" LOCATION "\n" \
+                                         "***** BLOCK:  [%u, %u, %u]\n" \
+                                         "***** THREAD: [%u, %u, %u]\n" \
+                                         "***** %s\n" \
+                                         "***** %s\n\n"; \
+      printf( formatString, \
+              blockIdx.x, blockIdx.y, blockIdx.z, \
+              threadIdx.x, threadIdx.y, threadIdx.z, \
+              STRINGIZE( CAUSE_MESSAGE ), \
+              STRINGIZE( GEOS_DETAIL_FIRST_ARG( __VA_ARGS__ ) ) \
+            ); \
+      asm volatile ( "s_trap 2" ); \
+    } \
+  } while( false )
 #endif
 
 /**
@@ -360,6 +402,27 @@
       asm ( "trap;" ); \
     } \
   } while( false )
+#elif __HIP_DEVICE_COMPILE__
+#define GEOS_WARNING_IF_CAUSE( COND, CAUSE_MESSAGE, ... ) \
+  do \
+  { \
+    if( COND ) \
+    { \
+      static char const formatString[] = "***** WARNING\n" \
+                                         "***** LOCATION" LOCATION "\n" \
+                                         "***** BLOCK:  [%u, %u, %u]\n" \
+                                         "***** THREAD: [%u, %u, %u]\n" \
+                                         "***** %s\n" \
+                                         "***** %s\n\n"; \
+      printf( formatString, \
+              blockIdx.x, blockIdx.y, blockIdx.z, \
+              threadIdx.x, threadIdx.y, threadIdx.z, \
+              STRINGIZE( CAUSE_MESSAGE ), \
+              STRINGIZE( GEOS_DETAIL_FIRST_ARG( __VA_ARGS__ ) ) \
+            ); \
+      asm volatile ( "s_trap 2" ); \
+    } \
+  } while( false )
 #endif
 
 /**

src/coreComponents/constitutive/fluid/multifluid/blackOil/BlackOilFluid.cpp

@@ -480,6 +480,9 @@ void BlackOilFluid::checkTableConsistency() const
                      InputError, getDataContext() );
     }
   }
+
+  // Needed to remove compilation warnings in the HIP build
+  GEOS_UNUSED_VAR(PT::OIL);
 }
 
 BlackOilFluid::KernelWrapper::

src/coreComponents/constitutiveDrivers/fluid/multiFluid/reactive/ReactiveFluidDriver.hpp

@@ -92,7 +92,7 @@ class ReactiveFluidDriver : public TaskBase
   };
 
   integer m_numSteps;      ///< Number of load steps
-  integer m_numColumns;    ///< Number of columns in data table (depends on number of fluid phases)
+  GEOS_MAYBE_UNUSED integer m_numColumns;    ///< Number of columns in data table (depends on number of fluid phases)
   integer m_numPhases;     ///< Number of fluid phases
 
   integer m_numPrimarySpecies; ///< Number of primary species (components)

src/coreComponents/dataRepository/ObjectCatalog.hpp

@@ -181,7 +181,7 @@ class CatalogInterface
    */
   //START_SPHINX_2
   static std::unique_ptr< BASETYPE > factory( std::string const & objectTypeName,
-                                              DataContext const & context,
+                                              GEOS_MAYBE_UNUSED DataContext const & context,
                                               ARGS... args )
   {
     // We stop the simulation if the type to create is not found

src/coreComponents/mesh/FaceElementSubRegion.cpp

@@ -452,10 +452,11 @@ localIndex FaceElementSubRegion::unpackToFaceRelation( buffer_unit_type const *
  * @param[in] elem2dToElems3d A mapping.
  * @param[in,out] elem2dToFaces This mapping will be corrected if needed to match @p elem2dToElems3d.
  */
-void fixNeighborMappingsInconsistency( string const & fractureName,
+void fixNeighborMappingsInconsistency( GEOS_MAYBE_UNUSED string const & fractureName,
                                        FixedToManyElementRelation const & elem2dToElems3d,
                                        FaceElementSubRegion::FaceMapType & elem2dToFaces )
 {
+  GEOS_MAYBE_UNUSED static constexpr std::string_view mappingInconsistency= "Mapping neighbor inconsistency detected for fracture {}.";
   {
     localIndex const num2dElems = elem2dToFaces.size( 0 );
     for( int e2d = 0; e2d < num2dElems; ++e2d )
@@ -489,9 +490,9 @@ void fixNeighborMappingsInconsistency( string const & fractureName,
         {
           std::swap( elem2dToFaces[e2d][0], elem2dToFaces[e2d][1] );
         }
-        else if( !matchStraight )
+        else
         {
-          GEOS_ERROR( "Mapping neighbor inconsistency detected for fracture " << fractureName );
+          GEOS_ERROR_IF( !matchStraight, GEOS_FMT( mappingInconsistency, fractureName ) );
         }
       }
     }
@@ -611,8 +612,7 @@ buildCollocatedEdgeBuckets( stdMap< globalIndex, globalIndex > const & reference
   stdMap< std::pair< globalIndex, globalIndex >, std::set< localIndex > > collocatedEdgeBuckets;
   for( auto const & p: edgesIds )
   {
-    static constexpr auto nodeNotFound = "Internal error when trying to access the reference collocated node for global node {}.";
-    GEOS_UNUSED_VAR( nodeNotFound ); // Not used in GPU builds.
+    GEOS_MAYBE_UNUSED static constexpr auto nodeNotFound = "Internal error when trying to access the reference collocated node for global node {}.";
 
     std::pair< globalIndex, globalIndex > const & nodes = p.first;
     localIndex const & edge = p.second;

src/coreComponents/mesh/PerforationData.cpp

@@ -143,7 +143,7 @@ void PerforationData::computeWellTransmissibility( MeshLevel const & mesh,
       if( close_perf )
       {
         m_localPerfStatus[iperf] = PerforationStatus::CLOSED;
-        WellElementRegion const & wellRegion = dynamicCast< WellElementRegion const & >( wellElemSubRegion.getParent().getParent() );
+        GEOS_MAYBE_UNUSED WellElementRegion const & wellRegion = dynamicCast< WellElementRegion const & >( wellElemSubRegion.getParent().getParent() );
         GEOS_LOG_RANK( "\n \nWarning! Perforation " << wellRegion.getWellGeneratorName() <<
                        " is defined with a zero transmissibility.\n" <<
                        "The simulation is going to proceed with this zero transmissibility,\n" <<

src/coreComponents/mesh/generators/ExternalMeshGeneratorBase.cpp

@@ -62,8 +62,8 @@ ExternalMeshGeneratorBase::ExternalMeshGeneratorBase( string const & name,
 
 void ExternalMeshGeneratorBase::postInputInitialization()
 {
-  auto const checkSizes = [this]( string_array const & from, string_array const & to,
-                                  string const & fromKey, string const & toKey )
+  auto const checkSizes = [&]( string_array const & from, string_array const & to,
+                               string const & fromKey, string const & toKey )
   {
     GEOS_THROW_IF_NE_MSG( from.size(), to.size(),
                           getWrapperDataContext( fromKey ) <<
@@ -74,7 +74,7 @@ void ExternalMeshGeneratorBase::postInputInitialization()
   checkSizes( m_volumicFieldsToImport, m_volumicFieldsInGEOS, viewKeyStruct::volumicFieldsToImportString(), viewKeyStruct::volumicFieldsInGEOSString() );
   checkSizes( m_surfacicFieldsToImport, m_surfacicFieldsInGEOS, viewKeyStruct::surfacicFieldsToImportString(), viewKeyStruct::surfacicFieldsInGEOSString() );
 
-  auto const checkDuplicates = [this]( string_array const & v, string const & key )
+  auto const checkDuplicates = [&]( string_array const & v, string const & key )
   {
     std::set< string > const tmp{ v.begin(), v.end() };
     bool const hasDuplicates = tmp.size() != LvArray::integerConversion< std::size_t >( v.size() );

src/coreComponents/mesh/generators/ParMETISInterface.hpp

@@ -26,13 +26,8 @@
 namespace geos
 {
 
-#if defined(GEOS_USE_HIP) // still need int32 hypre for the current hip-capable build
-/// Typedef to allow us to specify required parmetis integer type.
-using pmet_idx_t = int32_t;
-#else
 /// Typedef to allow us to specify required parmetis integer type.
 using pmet_idx_t = int64_t;
-#endif
 
 namespace parmetis
 {

src/coreComponents/mesh/generators/WellGeneratorBase.cpp

@@ -78,7 +78,7 @@ WellGeneratorBase::WellGeneratorBase( string const & name, Group * const parent
 Group * WellGeneratorBase::createChild( string const & childKey, string const & childName )
 {
   GEOS_LOG_RANK_0( GEOS_FMT( "{}: adding {} {}", getName(), childKey, childName ) );
-  const auto childTypes = { viewKeyStruct::perforationString() };
+  GEOS_MAYBE_UNUSED const auto childTypes = { viewKeyStruct::perforationString() };
   GEOS_ERROR_IF( childKey != viewKeyStruct::perforationString(),
                  CatalogInterface::unknownTypeError( childKey, getDataContext(), childTypes ),
                  getDataContext() );

src/coreComponents/mesh/mpiCommunications/SpatialPartition.cpp

@@ -215,7 +215,7 @@ void SpatialPartition::setSizes( real64 const ( &min )[ 3 ],
 
     //check to make sure our dimensions agree
     {
-      string_view partitionsLogMessage =
+      GEOS_MAYBE_UNUSED string_view partitionsLogMessage =
         "The total number of processes = {} does not correspond to the total number of partitions = {}.\n"
         "The number of cells in an axis cannot be lower that the partition count of this axis\n";