cu_affine_matrix_function.cu

Go to the documentation of this file.

/*
 *  SPDX-FileCopyrightText: Copyright 2021, Siavash Ameli <sameli@berkeley.edu>
 *  SPDX-License-Identifier: BSD-3-Clause
 *  SPDX-FileType: SOURCE
 *
 *  This program is free software: you can redistribute it and/or modify it
 *  under the terms of the license found in the LICENSE.txt file in the root
 *  directory of this source tree.
 */
 
 
// =======
// Headers
// =======
 
#include "./cu_affine_matrix_function.h"
#include <cassert>  // assert
#include "../_cu_definitions/cu_types.h" // __nv_fp8_e5m2, __nv_fp8_e4m3,
                                         // __half, __nv_bfloat16
#include "../_definitions/debugging.h"  // ASSERT
#include "../_cu_basic_algebra/cu_vector_operations.h"  // cuVectorOperations
#include "../_cuda_utilities/cuda_api.h"  // CudaAPI
#include "../_cu_arithmetics/cu_arithmetics.h"  // cu_arithmetics
 
 
// ===========
// constructor
// ===========
 
 
template <typename DataType>
cuAffineMatrixFunction<DataType>::cuAffineMatrixFunction():
    B_is_identity(false)
{
    // This class has one parameter that is t in A+tB
    this->num_parameters = 1;
}
 
 
// ==========
// destructor
// ==========
 
 
template <typename DataType>
cuAffineMatrixFunction<DataType>::~cuAffineMatrixFunction()
{
}
 
 
// ==============
// get eigenvalue
// ==============
 
 
template <typename DataType>
DataType cuAffineMatrixFunction<DataType>::get_eigenvalue(
        const DataType* known_parameters,
        const DataType known_eigenvalue,
        const DataType* inquiry_parameters) const
{
    ASSERT((this->eigenvalue_relation_known == 1),
            "An eigenvalue relation is not known. This function should be "
            "called only when the matrix B is a scalar multiple of the "
            "identity matrix");
 
    // Shift the eigenvalue by the parameter
    DataType inquiry_eigenvalue = \
        cu_arithmetics::add<DataType>(
            cu_arithmetics::sub<DataType>(
                known_eigenvalue,
                known_parameters[0]),
            inquiry_parameters[0]
        );
 
    return inquiry_eigenvalue;
}
 
 
// =================
// add scaled vector
// =================
 
 
template <typename DataType>
void cuAffineMatrixFunction<DataType>::_add_scaled_vector(
        const DataType* input_vector,
        const LongIndexType vector_size,
        const DataType scale,
        DataType* output_vector) const
{
    // Get device id
    int device_id = CudaAPI<DataType>::get_device();
 
    // Negative of scale (negate in double type as some types like
    // __nv_fp8_exmx does not yet support such operation yet)
    DataType neg_scale = cu_arithmetics::cast<double, DataType>(
            -cu_arithmetics::cast<DataType, double>(scale));
 
    // Subtracting two vectors with minus scale sign, which is adding.
    cuVectorOperations<DataType>::subtract_scaled_vector(
            this->cublas_handle[device_id], input_vector, vector_size,
            neg_scale, output_vector);
}
 
 
// ===============================
// Explicit template instantiation
// ===============================
 
#if defined(USE_CUDA_FP8_E5M2) && (USE_CUDA_FP8_E5M2 == 1)
    template class cuAffineMatrixFunction<__nv_fp8_e5m2>;
#endif
 
#if defined(USE_CUDA_FP8_E4M3) && (USE_CUDA_FP8_E4M3 == 1)
    template class cuAffineMatrixFunction<__nv_fp8_e4m3>;
#endif
 
#if defined(USE_CUDA_FP16) && (USE_CUDA_FP16 == 1)
    template class cuAffineMatrixFunction<__half>;
#endif
 
#if defined(USE_CUDA_BF16) && (USE_CUDA_BF16 == 1)
    template class cuAffineMatrixFunction<__nv_bfloat16>;
#endif
 
#if defined(USE_CUDA_FP32) && (USE_CUDA_FP32 == 1)
    template class cuAffineMatrixFunction<float>;
#endif
 
#if defined(USE_CUDA_FP64) && (USE_CUDA_FP64 == 1)
    template class cuAffineMatrixFunction<double>;
#endif