device.h

Go to the documentation of this file.

#ifndef MODULE_DEVICE_H_
#define MODULE_DEVICE_H_
 
#include "types.h"
#include <fstream>
 
#ifdef __MPI
#include "mpi.h"
#endif
 
namespace base_device
{
 
template <typename Device>
base_device::AbacusDevice_t get_device_type(const Device* dev);
 
template <typename T>
std::string get_current_precision(const T* var);
 
namespace information
{
 
std::string get_device_name(std::string device_flag);
 
int get_device_num(std::string device_flag);
 
void output_device_info(std::ostream& output);
 
int get_device_kpar(const int& kpar, const int& bndpar);
 
bool probe_gpu_availability();
 
std::string get_device_flag(const std::string& device,
                            const std::string& basis_type);
 
#if __MPI
int get_node_rank();
int get_node_rank_with_mpi_shared(const MPI_Comm mpi_comm = MPI_COMM_WORLD);
int stringCmp(const void* a, const void* b);
 
#ifdef __CUDA
int set_device_by_rank(const MPI_Comm mpi_comm = MPI_COMM_WORLD);
#endif
 
#endif
 
template <typename Device>
void print_device_info(const Device* dev, std::ofstream& ofs_device)
{
    return;
}
 
template <typename Device>
void record_device_memory(const Device* dev, std::ofstream& ofs_device, std::string str, size_t size)
{
    return;
}
 
#if defined(__CUDA) || defined(__ROCM)
template <>
void print_device_info<base_device::DEVICE_GPU>(const base_device::DEVICE_GPU *ctx, std::ofstream &ofs_device);
 
template <>
void record_device_memory<base_device::DEVICE_GPU>(const base_device::DEVICE_GPU* dev, std::ofstream& ofs_device, std::string str, size_t size);
#endif
 
} // end of namespace information
} // end of namespace base_device
 
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600 && !defined(__CUDA_ON_DCU)
static __inline__ __device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull = (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do
    {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN !=
        // NaN) } while (assumed != old);
    } while (assumed != old);
    return __longlong_as_double(old);
}
#endif
 
#endif // MODULE_DEVICE_H_