Macros.hpp

Go to the documentation of this file.

/*
 * Copyright (c) 2021, Lawrence Livermore National Security, LLC and LvArray contributors.
 * All rights reserved.
 * See the LICENSE file for details.
 * SPDX-License-Identifier: (BSD-3-Clause)
 */

#pragma once

// Source includes
#include "LvArrayConfig.hpp"
#include "system.hpp"

// System includes
#include <fstream>
#include <sstream>
#include <iostream>
#include <type_traits>


#if defined(LVARRAY_USE_CUDA) || defined(LVARRAY_USE_HIP)
#define LVARRAY_USE_DEVICE
#endif

#if defined(LVARRAY_USE_CUDA)
#define LVARRAY_DEFAULT_DEVICE_SPACE MemorySpace::cuda
#elif defined(LVARRAY_USE_HIP)
#define LVARRAY_DEFAULT_DEVICE_SPACE MemorySpace::hip
#endif

#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
#define LVARRAY_DEVICE_COMPILE
#define LVARRAY_FORCE_INLINE __forceinline__
#else
#define LVARRAY_FORCE_INLINE inline
#endif

#if defined(__CUDACC__) || defined(__HIPCC__)
// Denotes whether to define decorator macros later in this file.
#define LVARRAY_DECORATE
#endif


//#if !defined(NDEBUG) && defined(LVARRAY_DEVICE_COMPILE)
  #include <cassert>
//#endif

#define STRINGIZE_NX( A ) #A

#define STRINGIZE( A ) STRINGIZE_NX( A )

#define LVARRAY_UNUSED_ARG( X )

#define LVARRAY_UNUSED_VARIABLE( X ) ( ( void ) X )

#define LVARRAY_DEBUG_VAR( X ) LVARRAY_UNUSED_VARIABLE( X )

#define LOCATION __FILE__ ":" STRINGIZE( __LINE__ )

#define TYPEOFPTR( X ) std::remove_pointer_t< decltype( X ) >

#define TYPEOFREF( X ) std::remove_reference_t< decltype( X ) >

#define LVARRAY_LOG( ... ) std::cout << __VA_ARGS__ << std::endl

#define LVARRAY_LOG_VAR( ... ) LVARRAY_LOG( STRINGIZE( __VA_ARGS__ ) << " = " << __VA_ARGS__ )

#if defined(LVARRAY_DEVICE_COMPILE)
//   #if defined(__HIP_DEVICE_COMPILE__)
// // empty impl to avoid the possibility of printfs in device code
// //   on AMD, which can cause performance degradation just by being present
// #define LVARRAY_ERROR_IF( EXP, MSG )
  #if (!defined(NDEBUG)) || defined(__HIP_DEVICE_COMPILE__)
#define LVARRAY_ERROR_IF( EXP, MSG ) \
  do \
  { \
    if( EXP ) \
    { \
      assert( false && "EXP = " STRINGIZE( EXP ) "MSG = " STRINGIZE( MSG ) ); \
    } \
  } while( false )
  #else
#define LVARRAY_ERROR_IF( EXP, MSG ) \
  do \
  { \
    if( EXP ) \
    { \
      constexpr char const * formatString = "***** ERROR\n" \
                                            "***** LOCATION: " LOCATION "\n" \
                                                                        "***** Block: [%u, %u, %u]\n" \
                                                                        "***** Thread: [%u, %u, %u]\n" \
                                                                        "***** Controlling expression (should be false): " STRINGIZE( EXP ) "\n" \
                                                                                                                                            "***** MSG: " STRINGIZE( MSG ) "\n\n"; \
      printf( formatString, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z ); \
      asm ( "trap;" ); \
    } \
  } while( false )
  #endif
#else
#define LVARRAY_ERROR_IF( EXP, MSG ) \
  do \
  { \
    if( EXP ) \
    { \
      std::ostringstream __oss; \
      __oss << "***** ERROR\n"; \
      __oss << "***** LOCATION: " LOCATION "\n"; \
      __oss << "***** Controlling expression (should be false): " STRINGIZE( EXP ) "\n"; \
      __oss << MSG << "\n"; \
      __oss << LvArray::system::stackTrace( true ); \
      std::cout << __oss.str() << std::endl; \
      LvArray::system::callErrorHandler(); \
    } \
  } while( false )
#endif

#define LVARRAY_ERROR( MSG ) LVARRAY_ERROR_IF( true, MSG )

#if !defined(NDEBUG)
#define LVARRAY_ASSERT_MSG( EXP, MSG ) LVARRAY_ERROR_IF( !(EXP), MSG )
#else
#define LVARRAY_ASSERT_MSG( EXP, MSG ) ((void) 0)
#endif

#define LVARRAY_THROW_IF( EXP, MSG, TYPE ) \
  do \
  { \
    if( EXP ) \
    { \
      std::ostringstream __oss; \
      __oss << "\n"; \
      __oss << "***** LOCATION: " LOCATION "\n"; \
      __oss << "***** Controlling expression (should be false): " STRINGIZE( EXP ) "\n"; \
      __oss << MSG << "\n"; \
      __oss << LvArray::system::stackTrace( true ); \
      throw TYPE( __oss.str() ); \
    } \
  } while( false )

#define LVARRAY_THROW( MSG, TYPE ) LVARRAY_THROW_IF( true, MSG, TYPE )

#define LVARRAY_ASSERT( EXP ) LVARRAY_ASSERT_MSG( EXP, "" )

#define LVARRAY_WARNING_IF( EXP, MSG ) \
  do \
  { \
    if( EXP ) \
    { \
      std::ostringstream __oss; \
      __oss << "***** WARNING\n"; \
      __oss << "***** LOCATION: " LOCATION "\n"; \
      __oss << "***** Controlling expression (should be false): " STRINGIZE( EXP ) "\n"; \
      __oss << MSG; \
      std::cout << __oss.str() << std::endl; \
    } \
  } while( false )

#define LVARRAY_WARNING( MSG ) LVARRAY_WARNING_IF( true, MSG )

#define LVARRAY_INFO_IF( EXP, MSG ) \
  do \
  { \
    if( EXP ) \
    { \
      std::ostringstream __oss; \
      __oss << "***** INFO\n"; \
      __oss << "***** LOCATION: " LOCATION "\n"; \
      __oss << "***** Controlling expression: " STRINGIZE( EXP ) "\n"; \
      __oss << MSG; \
      std::cout << __oss.str() << std::endl; \
    } \
  } while( false )

#define LVARRAY_INFO( msg ) LVARRAY_INFO_IF( true, msg )

#define LVARRAY_ERROR_IF_OP_MSG( lhs, OP, NOP, rhs, msg ) \
  LVARRAY_ERROR_IF( lhs OP rhs, \
                    msg << "\n" << \
                    "Expected " << #lhs << " " << #NOP << " " << #rhs << "\n" << \
                    "  " << #lhs << " = " << lhs << "\n" << \
                    "  " << #rhs << " = " << rhs << "\n" )

#define LVARRAY_THROW_IF_OP_MSG( lhs, OP, NOP, rhs, msg, TYPE ) \
  LVARRAY_THROW_IF( lhs OP rhs, \
                    msg << "\n" << \
                    "Expected " << #lhs << " " << #NOP << " " << #rhs << "\n" << \
                    "  " << #lhs << " = " << lhs << "\n" << \
                    "  " << #rhs << " = " << rhs << "\n", TYPE )

#define LVARRAY_ERROR_IF_EQ_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, ==, !=, rhs, msg )

#define LVARRAY_THROW_IF_EQ_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, ==, !=, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_EQ( lhs, rhs ) LVARRAY_ERROR_IF_EQ_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_EQ( lhs, rhs, TYPE ) LVARRAY_THROW_IF_EQ_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ERROR_IF_NE_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, !=, ==, rhs, msg )

#define LVARRAY_THROW_IF_NE_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, !=, ==, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_NE( lhs, rhs ) LVARRAY_ERROR_IF_NE_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_NE( lhs, rhs, TYPE ) LVARRAY_THROW_IF_NE_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ERROR_IF_GT_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, >, <=, rhs, msg )

#define LVARRAY_THROW_IF_GT_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, >, <=, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_GT( lhs, rhs ) LVARRAY_ERROR_IF_GT_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_GT( lhs, rhs, TYPE ) LVARRAY_THROW_IF_GT_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ERROR_IF_GE_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, >=, <, rhs, msg )

#define LVARRAY_THROW_IF_GE_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, >=, <, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_GE( lhs, rhs ) LVARRAY_ERROR_IF_GE_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_GE( lhs, rhs, TYPE ) LVARRAY_THROW_IF_GE_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ERROR_IF_LT_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, <, >=, rhs, msg )

#define LVARRAY_THROW_IF_LT_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, <, >=, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_LT( lhs, rhs ) LVARRAY_ERROR_IF_LT_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_LT( lhs, rhs, TYPE ) LVARRAY_THROW_IF_LT_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ERROR_IF_LE_MSG( lhs, rhs, msg ) LVARRAY_ERROR_IF_OP_MSG( lhs, <=, >, rhs, msg )

#define LVARRAY_THROW_IF_LE_MSG( lhs, rhs, msg, TYPE ) LVARRAY_THROW_IF_OP_MSG( lhs, <=, >, rhs, msg, TYPE )

#define LVARRAY_ERROR_IF_LE( lhs, rhs ) LVARRAY_ERROR_IF_GE_MSG( lhs, rhs, "" )

#define LVARRAY_THROW_IF_LE( lhs, rhs, TYPE ) LVARRAY_THROW_IF_GE_MSG( lhs, rhs, "", TYPE )

#define LVARRAY_ASSERT_OP_MSG( lhs, OP, rhs, msg ) \
  LVARRAY_ASSERT_MSG( lhs OP rhs, \
                      msg << "\n" << \
                      "  " << #lhs << " = " << lhs << "\n" << \
                      "  " << #rhs << " = " << rhs << "\n" )

#define LVARRAY_ASSERT_EQ_MSG( lhs, rhs, msg ) LVARRAY_ASSERT_OP_MSG( lhs, ==, rhs, msg )

#define LVARRAY_ASSERT_EQ( lhs, rhs ) LVARRAY_ASSERT_EQ_MSG( lhs, rhs, "" )

#define LVARRAY_ASSERT_NE_MSG( lhs, rhs, msg ) LVARRAY_ASSERT_OP_MSG( lhs, !=, rhs, msg )

#define LVARRAY_ASSERT_NE( lhs, rhs ) LVARRAY_ASSERT_NE_MSG( lhs, rhs, "" )

#define LVARRAY_ASSERT_GT_MSG( lhs, rhs, msg ) LVARRAY_ASSERT_OP_MSG( lhs, >, rhs, msg )

#define LVARRAY_ASSERT_GT( lhs, rhs ) LVARRAY_ASSERT_GT_MSG( lhs, rhs, "" )

#define LVARRAY_ASSERT_GE_MSG( lhs, rhs, msg ) LVARRAY_ASSERT_OP_MSG( lhs, >=, rhs, msg )

#define LVARRAY_ASSERT_GE( lhs, rhs ) LVARRAY_ASSERT_GE_MSG( lhs, rhs, "" )

#if defined(LVARRAY_DECORATE)
#define LVARRAY_HOST_DEVICE __host__ __device__

#if defined( LVARRAY_USE_HIP )
#define LVARRAY_HOST_DEVICE_HIP __host__ __device__
#else
#define LVARRAY_HOST_DEVICE_HIP
#endif

#define LVARRAY_DEVICE __device__

#if defined(LVARRAY_USE_CUDA)
#define DISABLE_HD_WARNING _Pragma("hd_warning_disable")
#else
#define DISABLE_HD_WARNING
#endif
#else
#define LVARRAY_HOST_DEVICE
#define LVARRAY_HOST_DEVICE_HIP

#define LVARRAY_DEVICE

#define DISABLE_HD_WARNING
#endif


#if defined(__clang__)
#define LVARRAY_RESTRICT __restrict__
#define LVARRAY_RESTRICT_REF __restrict__
#define LVARRAY_INTEL_CONSTEXPR constexpr
#elif defined(__GNUC__)
  #if defined(__INTEL_COMPILER)
#define LVARRAY_RESTRICT __restrict__
#define LVARRAY_RESTRICT_REF __restrict__
#define LVARRAY_INTEL_CONSTEXPR
  #else
#define LVARRAY_RESTRICT __restrict__
#define LVARRAY_RESTRICT_REF __restrict__
#define LVARRAY_INTEL_CONSTEXPR constexpr
  #endif
#endif

#if !defined(LVARRAY_BOUNDS_CHECK)

#define CONSTEXPR_WITHOUT_BOUNDS_CHECK constexpr
#else

#define CONSTEXPR_WITHOUT_BOUNDS_CHECK
#endif

#if defined(NDEBUG)

#define CONSTEXPR_WITH_NDEBUG constexpr
#else

#define CONSTEXPR_WITH_NDEBUG
#endif

#if !defined(LVARRAY_BOUNDS_CHECK)

#define CONSTEXPR_WITHOUT_BOUNDS_CHECK constexpr
#else

#define CONSTEXPR_WITHOUT_BOUNDS_CHECK
#endif

#if defined(NDEBUG)

#define CONSTEXPR_WITH_NDEBUG constexpr
#else

#define CONSTEXPR_WITH_NDEBUG
#endif

// TPL includes
#include <RAJA/RAJA.hpp>

template< typename >
struct RAJAHelper
{};

using serialPolicy = RAJA::loop_exec;

template<>
struct RAJAHelper< serialPolicy >
{
  using ReducePolicy = RAJA::seq_reduce;
  using AtomicPolicy = RAJA::seq_atomic;
};

#if defined(RAJA_ENABLE_OPENMP)

using parallelHostPolicy = RAJA::omp_parallel_for_exec;

template<>
struct RAJAHelper< parallelHostPolicy >
{
  using ReducePolicy = RAJA::omp_reduce;
  using AtomicPolicy = RAJA::omp_atomic;
};

#endif

#if defined(LVARRAY_USE_CUDA)

template< unsigned long THREADS_PER_BLOCK >
using parallelDevicePolicy = RAJA::cuda_exec< THREADS_PER_BLOCK >;

template< unsigned long N >
struct RAJAHelper< RAJA::cuda_exec< N > >
{
  using ReducePolicy = RAJA::cuda_reduce;
  using AtomicPolicy = RAJA::cuda_atomic;
};

#elif defined(LVARRAY_USE_HIP)

template< unsigned long THREADS_PER_BLOCK >
using parallelDevicePolicy = RAJA::hip_exec< THREADS_PER_BLOCK >;

template< unsigned long N >
struct RAJAHelper< RAJA::hip_exec< N > >
{
  using ReducePolicy = RAJA::hip_reduce;
  using AtomicPolicy = RAJA::hip_atomic;
};

#endif