kernel_enums.hpp

/*
  Copyright 2024 Equinor ASA

  This file is part of the Open Porous Media project (OPM).

  OPM is free software: you can redistribute it and/or modify
  it under the terms of the GNU General Public License as published by
  the Free Software Foundation, either version 3 of the License, or
  (at your option) any later version.

  OPM is distributed in the hope that it will be useful,
  but WITHOUT ANY WARRANTY; without even the implied warranty of
  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  GNU General Public License for more details.

  You should have received a copy of the GNU General Public License
  along with OPM.  If not, see <http://www.gnu.org/licenses/>.
*/

#ifndef OPM_GPUISTL_KERNEL_ENUMS_HPP
#define OPM_GPUISTL_KERNEL_ENUMS_HPP

#include <opm/common/ErrorMacros.hpp>

#include <cuda_runtime.h>

#if CUDA_VERSION >= 12100
#include <fmt/core.h>
#else
#include <sstream>
#endif

/*
    This file organizes a growing amount of different mixed precision options for the preconditioners.
*/

namespace Opm::gpuistl {
    // Mixed precision schemes used for storing the matrix in GPU memory
    enum class MatrixStorageMPScheme {
        DOUBLE_DIAG_DOUBLE_OFFDIAG = 0, // full precision should be default
        FLOAT_DIAG_FLOAT_OFFDIAG = 1,
        DOUBLE_DIAG_FLOAT_OFFDIAG = 2
    };

    namespace detail {
        bool isValidMatrixStorageMPScheme(int scheme);
    }

    inline MatrixStorageMPScheme makeMatrixStorageMPScheme(int scheme) {
        if (!detail::isValidMatrixStorageMPScheme(scheme)) {
#if CUDA_VERSION >= 12100
            OPM_THROW(std::invalid_argument,
                      fmt::format(fmt::runtime("Invalid matrix storage mixed precision scheme chosen: {}.\n"
                                  "Valid Schemes:\n"
                                  "\t0: DOUBLE_DIAG_DOUBLE_OFFDIAG\n"
                                  "\t1: FLOAT_DIAG_FLOAT_OFFDIAG\n"
                                  "\t2: DOUBLE_DIAG_FLOAT_OFFDIAG"),
                                  scheme));
#else
            std::stringstream str;
            str << "Invalid matrix storage mixed precision scheme chosen: " << scheme << ".\n"
                << "Valid Schemes:\n"
                   "\t0: DOUBLE_DIAG_DOUBLE_OFFDIAG\n"
                   "\t1: FLOAT_DIAG_FLOAT_OFFDIAG\n"
                   "\t2: DOUBLE_DIAG_FLOAT_OFFDIAG";
            OPM_THROW(std::invalid_argument, str.str());
#endif
        }
        return static_cast<MatrixStorageMPScheme>(scheme);
    }

    namespace detail {

        __host__ __device__ constexpr bool storeDiagonalAsFloat(MatrixStorageMPScheme scheme) {
            switch (scheme) {
                case MatrixStorageMPScheme::DOUBLE_DIAG_DOUBLE_OFFDIAG:
                    return false;
                case MatrixStorageMPScheme::FLOAT_DIAG_FLOAT_OFFDIAG:
                    return true;
                case MatrixStorageMPScheme::DOUBLE_DIAG_FLOAT_OFFDIAG:
                    return false;
                default:
                    return false;
            }
        }

        __host__ __device__ constexpr bool storeOffDiagonalAsFloat(MatrixStorageMPScheme scheme) {
            switch (scheme) {
                case MatrixStorageMPScheme::DOUBLE_DIAG_DOUBLE_OFFDIAG:
                    return false;
                case MatrixStorageMPScheme::FLOAT_DIAG_FLOAT_OFFDIAG:
                    return true;
                case MatrixStorageMPScheme::DOUBLE_DIAG_FLOAT_OFFDIAG:
                    return true;
                default:
                    return false;
            }
        }

        // returns true if we use anything else that the the default double precision for everything
        __host__ __device__ constexpr bool usingMixedPrecision(MatrixStorageMPScheme scheme) {
            switch (scheme) {
                case MatrixStorageMPScheme::DOUBLE_DIAG_DOUBLE_OFFDIAG:
                    return false;
                case MatrixStorageMPScheme::FLOAT_DIAG_FLOAT_OFFDIAG:
                    return true;
                case MatrixStorageMPScheme::DOUBLE_DIAG_FLOAT_OFFDIAG:
                    return true;
                default:
                    return false;
            }
        }

        inline bool isValidMatrixStorageMPScheme(int scheme) {
            switch (static_cast<MatrixStorageMPScheme>(scheme)) {
                case MatrixStorageMPScheme::DOUBLE_DIAG_DOUBLE_OFFDIAG:
                case MatrixStorageMPScheme::FLOAT_DIAG_FLOAT_OFFDIAG:
                case MatrixStorageMPScheme::DOUBLE_DIAG_FLOAT_OFFDIAG:
                    return true;
                default:
                    return false;
            }
        }
    }
}

#endif // OPM_GPUISTL_KERNEL_ENUMS_HPP