CudaMemory.cu

#include "config.h"
#include <cstddef>
#include "CudaMemory.cuh"
#include "cuda_macro.h"
#include "util/cudify/cudify.hpp"
#include <cstring>

#define CUDA_EVENT 0x1201

bool CudaMemory::flush()
{
    if (hm != NULL && dm != NULL)
    {

        #ifdef __HIP__
        CUDA_SAFE_CALL(hipMemcpy(dm,hm,sz,hipMemcpyHostToDevice));
        #else
        CUDA_SAFE_CALL(cudaMemcpy(dm,hm,sz,cudaMemcpyHostToDevice));
        #endif
    }
    
    return true;
}

bool CudaMemory::allocate(size_t sz)
{
    if (dm == NULL)
    {
        #ifdef __HIP__
        CUDA_SAFE_CALL(hipMalloc(&dm,sz));
        #elif defined(CUDIFY_USE_CUDA)
        CUDA_SAFE_CALL(cudaMalloc(&dm,sz));
        #else
        if (sz != 0)
        {
            dm = new unsigned char[sz];
            #ifdef GARBAGE_INJECTOR
            memset(dm,0xFF,sz);
            #endif
        }
        #endif
    }
    else
    {
        if (sz != this->sz)
        {
            std::cout << __FILE__ << ":" << __LINE__ << " error FATAL: using allocate to resize the memory, please use resize." << std::endl;
            return false;
        }
    }

    this->sz = sz;

#if defined(GARBAGE_INJECTOR) && !defined(CUDA_ON_CPU)
    CUDA_SAFE_CALL(cudaMemset(dm,-1,sz))
#endif

    return true;
}

void CudaMemory::destroy()
{
    if (dm != NULL)
    {
        #ifdef __HIP__
        CUDA_SAFE_CALL(hipFree(dm));
        #elif defined(CUDIFY_USE_CUDA)
        CUDA_SAFE_CALL(cudaFree(dm));
        #else
        delete [] (unsigned char *)dm;
        #endif
        dm = NULL;
    }

    if (hm != NULL)
    {
        #ifdef __HIP__
        CUDA_SAFE_CALL(hipHostFree(hm));
        #elif defined(CUDIFY_USE_CUDA)
        CUDA_SAFE_CALL(cudaFreeHost(hm));
        #else
        delete [] (unsigned char *)hm;
        #endif
        hm = NULL;
    }
    
    sz = 0;
}

void CudaMemory::deviceToDevice(void * ptr, size_t start, size_t stop, size_t offset)
{
    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)dm)+offset,((unsigned char *)ptr)+start,(stop-start),hipMemcpyDeviceToDevice));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)dm)+offset,((unsigned char *)ptr)+start,(stop-start),cudaMemcpyDeviceToDevice));
    #endif
}

void CudaMemory::allocate_host(size_t sz) const
{
    if (hm == NULL)
    {
        #ifdef __HIP__
        CUDA_SAFE_CALL(hipHostMalloc(&hm,sz,hipHostMallocMapped))
        #elif defined(CUDIFY_USE_CUDA)
        CUDA_SAFE_CALL(cudaHostAlloc(&hm,sz,cudaHostAllocMapped))
        #else
        hm = new unsigned char[sz];
        #ifdef GARBAGE_INJECTOR
        memset(hm,0xFF,sz);
        #endif
        #endif
    }
}

bool CudaMemory::copyFromPointer(const void * ptr)
{
    // check if we have a host buffer, if not allocate it

    allocate_host(sz);

    // get the device pointer

    void * dvp;
    #ifdef __HIP__
    CUDA_SAFE_CALL(hipHostGetDevicePointer(&dvp,hm,0));
    // memory copy
    memcpy(dvp,ptr,sz);
    #else
    CUDA_SAFE_CALL(cudaHostGetDevicePointer(&dvp,hm,0));
    // memory copy
    memcpy(dvp,ptr,sz);
    #endif

    return true;
}

bool CudaMemory::copyDeviceToDevice(const CudaMemory & m)
{

    if (m.sz > sz)
    {
        std::cerr << "Error " << __LINE__ << __FILE__ << ": source buffer is too big to copy";
        return false;
    }

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(dm,m.dm,m.sz,hipMemcpyDeviceToDevice));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(dm,m.dm,m.sz,cudaMemcpyDeviceToDevice));
    #endif

    return true;
}

bool CudaMemory::copy(const memory & m)
{
    const CudaMemory * ofpm = dynamic_cast<const CudaMemory *>(&m);


    if (ofpm == NULL)
    {
        // copy the memory from device to host and from host to device

        return copyFromPointer(m.getPointer());
    }
    else
    {
        // they are the same memory type, use cuda/thrust buffer copy

        return copyDeviceToDevice(*ofpm);
    }
}

size_t CudaMemory::size() const
{
    return sz;
}


bool CudaMemory::resize(size_t sz)
{
    // if the allocated memory is enough, do not resize
    if (sz <= CudaMemory::size())
    {return true;}


    if (CudaMemory::size() == 0)
    {return allocate(sz);}

    void * thm = NULL;

    void * tdm = NULL;

    if (dm != NULL)
    {
        if (this->sz < sz)
        {
            #ifdef __HIP__
            CUDA_SAFE_CALL(hipMalloc(&tdm,sz));
            #elif defined(CUDIFY_USE_CUDA)
            CUDA_SAFE_CALL(cudaMalloc(&tdm,sz));
            #else
            tdm = new unsigned char [sz];
            #ifdef GARBAGE_INJECTOR
            memset(tdm,0xFF,sz);
            #endif
            #endif

#ifdef GARBAGE_INJECTOR
            #ifdef __HIP__
            CUDA_SAFE_CALL(hipMemset(tdm,-1,sz));
            #elif defined(CUDIFY_USE_CUDA)
            CUDA_SAFE_CALL(cudaMemset(tdm,-1,sz));
            #endif
#endif
        }

        #ifdef __HIP__
        CUDA_SAFE_CALL(hipMemcpy(tdm,dm,CudaMemory::size(),hipMemcpyDeviceToDevice));
        #else
        CUDA_SAFE_CALL(cudaMemcpy(tdm,dm,CudaMemory::size(),cudaMemcpyDeviceToDevice));
        #endif
    }

    if (hm != NULL)
    {
        if (this->sz < sz)
        {
            #ifdef __HIP__
            CUDA_SAFE_CALL(hipHostMalloc(&thm,sz,hipHostMallocMapped));
            #elif defined(CUDIFY_USE_CUDA)
            CUDA_SAFE_CALL(cudaHostAlloc(&thm,sz,cudaHostAllocMapped));
            #else
            thm = new unsigned char [sz];
            #ifdef GARBAGE_INJECTOR
            memset(thm,0xFF,sz);
            #endif
            #endif
        }

        #ifdef __HIP__
        CUDA_SAFE_CALL(hipMemcpy(thm,hm,CudaMemory::size(),hipMemcpyHostToHost));
        #else
        CUDA_SAFE_CALL(cudaMemcpy(thm,hm,CudaMemory::size(),cudaMemcpyHostToHost));
        #endif
    }


    destroy();

    dm = tdm;
    hm = thm;


    this->sz = sz;

    return true;
}

void * CudaMemory::getPointer()
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    return hm;
}

void CudaMemory::deviceToHost()
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(hm,dm,sz,hipMemcpyDeviceToHost));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(hm,dm,sz,cudaMemcpyDeviceToHost));
    #endif
}

void CudaMemory::deviceToHost(CudaMemory & mem)
{
    // allocate an host memory if not allocated
    if (mem.hm == NULL)
        mem.allocate_host(sz);

    if (mem.sz > sz)
    {resize(mem.sz);}

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(mem.hm,dm,mem.sz,hipMemcpyDeviceToHost));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(mem.hm,dm,mem.sz,cudaMemcpyDeviceToHost));
    #endif
}

void CudaMemory::hostToDevice(CudaMemory & mem)
{
    // allocate an host memory if not allocated
    if (mem.hm == NULL)
        mem.allocate_host(sz);

    if (mem.sz > sz)
    {resize(mem.sz);}

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(dm,mem.hm,mem.sz,hipMemcpyHostToDevice));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(dm,mem.hm,mem.sz,cudaMemcpyHostToDevice));
    #endif
}

void CudaMemory::hostToDevice(size_t start, size_t stop)
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)dm)+start,((unsigned char *)hm)+start,(stop-start),hipMemcpyHostToDevice));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)dm)+start,((unsigned char *)hm)+start,(stop-start),cudaMemcpyHostToDevice));
    #endif
}

void CudaMemory::deviceToHost(size_t start, size_t stop)
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)hm)+start,((unsigned char *)dm)+start,(stop-start),hipMemcpyDeviceToHost));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)hm)+start,((unsigned char *)dm)+start,(stop-start),cudaMemcpyDeviceToHost));
    #endif
}



const void * CudaMemory::getPointer() const
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    return hm;
}

void CudaMemory::fill(unsigned char c)
{
    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemset(dm,c,size()));
    #elif defined(CUDIFY_USE_CUDA)
    CUDA_SAFE_CALL(cudaMemset(dm,c,size()));
    #else
    memset(dm,c,size());
    #endif
    if (hm != NULL)
    {memset(hm,c,size());}
}

void * CudaMemory::getDevicePointer()
{
    return dm;
}

void CudaMemory::hostToDevice()
{
    // allocate an host memory if not allocated
    if (hm == NULL)
        allocate_host(sz);

    #ifdef __HIP__
    CUDA_SAFE_CALL(hipMemcpy(dm,hm,sz,hipMemcpyHostToDevice));
    #else
    CUDA_SAFE_CALL(cudaMemcpy(dm,hm,sz,cudaMemcpyHostToDevice));
    #endif
}


void CudaMemory::swap(CudaMemory & mem)
{
    size_t sz_tmp;
    void * dm_tmp;
//  long int ref_cnt_tmp;
    bool is_hm_sync_tmp;
    void * hm_tmp;

    hm_tmp = hm;
    is_hm_sync_tmp = is_hm_sync;
    sz_tmp = sz;
    dm_tmp = dm;
//  ref_cnt_tmp = ref_cnt;

    hm = mem.hm;
    is_hm_sync = mem.is_hm_sync;
    sz = mem.sz;
    dm = mem.dm;
    ref_cnt = mem.ref_cnt;

    mem.hm = hm_tmp;
    mem.is_hm_sync = is_hm_sync_tmp;
    mem.sz = sz_tmp;
    mem.dm = dm_tmp;
//  mem.ref_cnt = ref_cnt_tmp;
}