_cu_epsilon.h

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.
 */
 
#ifndef _CU_ARITHMETICS_CU_EPSILON_H_
#define _CU_ARITHMETICS_CU_EPSILON_H_
 
// =======
// Headers
// =======
 
#include "../_cu_definitions/cu_types.h" // __nv_fp8_e5m2, __nv_fp8_e4m3,
                                         // __half, __nv_bfloat16
 
 
// =============
// cu arithmetic
// =============
 
 
namespace cu_arithmetics
{
    // =======
    // epsilon
    // =======
 
 
    template <typename DataType>
    inline __host__ __device__ DataType epsilon();
 
 
    // =======
    // epsilon (__nv_fp8_e5m2)
    // =======
 
    
    #if defined(USE_CUDA_FP8_E5M2) && (USE_CUDA_FP8_E5M2 == 1)
    template<>
    inline __host__ __device__ __nv_fp8_e5m2 epsilon<__nv_fp8_e5m2>()
    {
        // This is 2^{-10}, as __nv_fp8_e5m2 type has 2 digits for mantissa. 
        return __nv_fp8_e5m2(0.25f);
    }
    #endif
 
 
    // =======
    // epsilon (__nv_fp8_e4m3)
    // =======
 
    
    #if defined(USE_CUDA_FP8_E4M3) && (USE_CUDA_FP8_E4M3 == 1)
    template<>
    inline __host__ __device__ __nv_fp8_e4m3 epsilon<__nv_fp8_e4m3>()
    {
        // This is 2^{-10}, as __nv_fp8_e4m3 type has 3 digits for mantissa. 
        return __nv_fp8_e4m3(0.125f);
    }
    #endif
 
 
    // =======
    // epsilon (__half)
    // =======
 
    
    #if defined(USE_CUDA_FP16) && (USE_CUDA_FP16 == 1)
    template<>
    inline __host__ __device__ __half epsilon<__half>()
    {
        // This is 2^{-10}, as __half type has 10 digits for mantissa. 
        return __float2half(0.00097656f);
    }
    #endif
 
 
    // =======
    // epsilon (__nv_bfloat16)
    // =======
 
    
#if defined(USE_CUDA_BF16) && (USE_CUDA_BF16 == 1)
    template<>
    inline __host__ __device__ __nv_bfloat16 epsilon<__nv_bfloat16>()
    {
        // This is 2^{-7}, as __half type has 10 digits for mantissa. 
        return __float2bfloat16(0.0078125f);
    }
    #endif
 
 
    // =======
    // epsilon (float)
    // =======
 
    
    #if defined(USE_CUDA_FP32) && (USE_CUDA_FP32 == 1)
    template<>
    inline __host__ __device__ float epsilon<float>()
    {
        // Instead of the hard-coded number below, one may use
        // the std::numeric_limits::epsilon(), but, then this function cannot
        // be called as a __device__ code. Hence, the value of 2^(-23) is hard
        // coded below.
        return 1.1920929e-7f;
    }
    #endif
 
 
    // =======
    // epsilon (float)
    // =======
 
    
    #if defined(USE_CUDA_FP64) && (USE_CUDA_FP64 == 1)
    template<>
    inline __host__ __device__ double epsilon<double>()
    {
        // Instead of the hard-coded number below, one may use
        // the std::numeric_limits::epsilon(), but, then this function cannot
        // be called as a __device__ code. Hence, the value of 2^(-52) is hard
        // coded below.
        return 2.220446049250313e-16;
    }
    #endif
 
}  // namespace cu_arithmetics
 
#endif  // _CU_ARITHMETICS_CU_EPSILON_H_