CudaMemory.cuh

/*
 * CudaMemory.cu
 *
 *  Created on: Aug 17, 2014
 *      Author: Pietro Incardona
 */

#ifndef CUDA_MEMORY_CUH_
#define CUDA_MEMORY_CUH_

#define EXCEPT_MC noexcept

#include "config.h"
#include "memory.hpp"
#include <iostream>

#include "util/cuda_util.hpp"

extern size_t TotCudaMemoryAllocated;

__device__ inline size_t align_number_device(size_t al, size_t number)
{
    return number + ((number % al) != 0)*(al - number % al);
}

static __device__ unsigned char global_cuda_error_array[256];

class CudaMemory : public memory
{
    bool is_hm_sync;
    
    size_t sz;
    
    void * dm;
    
    mutable void * hm;

    size_t ref_cnt;
    
    void allocate_host(size_t sz) const;
    
    bool copyFromPointer(const void * ptr);
    
public:
    
    bool copyDeviceToDevice(const CudaMemory & m);

    virtual bool flush();
    virtual bool allocate(size_t sz);
    virtual void destroy();
    virtual bool copy(const memory & m);
    virtual size_t size() const;
    virtual bool resize(size_t sz);
    virtual void * getPointer();
    
    virtual const void * getPointer() const;
    
    virtual void * getDevicePointer();

    virtual void hostToDevice();

    virtual void deviceToHost();

    virtual void deviceToHost(size_t start, size_t stop);

    virtual void hostToDevice(size_t start, size_t top);

    void hostToDevice(CudaMemory & mem);

    void deviceToHost(CudaMemory & mem);

    virtual void fill(unsigned char c);

    void isNotSync() {is_hm_sync = false;}
    
    public:
    
    virtual void incRef()
    {
        ref_cnt++;
    }

    virtual void decRef()
    {ref_cnt--;}
    
    virtual long int ref()
    {
        return ref_cnt;
    }

    bool isInitialized()
    {
        return false;
    }
    
    // Copy the memory (device and host)
    CudaMemory & operator=(const CudaMemory & mem)
    {
        copy(mem);
        return *this;
    }

    // Copy the Cuda memory
    CudaMemory(const CudaMemory & mem)
    :CudaMemory()
    {
        allocate(mem.size());
        copy(mem);
    }

    CudaMemory(CudaMemory && mem) EXCEPT_MC
    {
        is_hm_sync = mem.is_hm_sync;
        sz = mem.sz;
        dm = mem.dm;
        hm = mem.hm;
        ref_cnt = mem.ref_cnt;

        // reset mem
        mem.is_hm_sync = false;
        mem.sz = 0;
        mem.dm = NULL;
        mem.hm = NULL;
        mem.ref_cnt = 0;
    }
    
    CudaMemory():is_hm_sync(true),sz(0),dm(0),hm(0),ref_cnt(0) {};
    
    CudaMemory(size_t sz):is_hm_sync(true),sz(0),dm(0),hm(0),ref_cnt(0)
    {
        allocate(sz);
    };

    ~CudaMemory()   
    {
        if(ref_cnt == 0)
            destroy();
        else
            std::cerr << "Error: " << __FILE__ << " " << __LINE__ << " destroying a live object" << "\n"; 
    };

    void deviceToDevice(void * ptr, size_t start, size_t stop, size_t offset);

    void swap(CudaMemory & mem);

    constexpr static bool isDeviceHostSame()
    {
        return false;
    }

    void * toKernel()
    {
        return getDevicePointer();
    }
};


#endif