device_mem_pointer_matrix_cu.h

Go to the documentation of this file.

#ifndef DEVICE_MEM_POINTER_MATRIX_CU_H
#define DEVICE_MEM_POINTER_MATRIX_CU_H

#ifdef SKEPU_CUDA

#include <iostream>
#include <cuda.h>


#include "device_cu.h"

namespace skepu
{

template <typename T>
class DeviceMemPointer_Matrix_CU
{

public:
   DeviceMemPointer_Matrix_CU(T* start, int rows, int cols, Device_CU *device, bool usePitch=false);
   ~DeviceMemPointer_Matrix_CU();

   void copyHostToDevice(int rows=-1, int cols=-1) const;
   void copyDeviceToHost(int rows=-1, int cols=-1) const;
   T* getDeviceDataPointer() const;
   int getDeviceID() const;
   void changeDeviceData();

   size_t m_pitch;
private:
   T* m_hostDataPointer;
   T* m_deviceDataPointer;
   int m_rows;
   int m_cols;
   int m_deviceID;
   Device_CU *m_dev;

   bool m_usePitch;

   inline int size()
   {
      return m_rows * m_cols;
   }

   mutable bool deviceDataHasChanged;
};

template <typename T>
DeviceMemPointer_Matrix_CU<T>::DeviceMemPointer_Matrix_CU(T* start, int rows, int cols, Device_CU *device, bool usePitch) : m_hostDataPointer(start), m_rows(rows), m_cols(cols), m_dev(device)
{
   cudaError_t err;
   size_t sizeVec = rows*cols*sizeof(T);

   DEBUG_TEXT_LEVEL2("Alloc: " << rows*cols <<"\n")

   m_deviceID = m_dev->getDeviceID();

   cudaSetDevice(m_deviceID);

   m_usePitch = usePitch;
   if(usePitch)
   {
      err = cudaMallocPitch((void**)&m_deviceDataPointer, &m_pitch, cols * sizeof(T), rows);
      m_pitch = (m_pitch)/sizeof(T);
   }
   else
   {
      err = cudaMalloc((void**)&m_deviceDataPointer, sizeVec);
      m_pitch = cols;
   }
   if(err != cudaSuccess)
   {
      std::cerr<<"Error allocating memory on device\n";
   }

   deviceDataHasChanged = false;
}




template <typename T>
DeviceMemPointer_Matrix_CU<T>::~DeviceMemPointer_Matrix_CU()
{
   DEBUG_TEXT_LEVEL2("DeAlloc: "  <<"\n")

   cudaSetDevice(m_deviceID);

   cudaFree(m_deviceDataPointer);
}

template <typename T>
void DeviceMemPointer_Matrix_CU<T>::copyHostToDevice(int rows, int cols) const
{
   if(m_hostDataPointer != NULL)
   {
      DEBUG_TEXT_LEVEL2("HOST_TO_DEVICE!!!\n")

      cudaError_t err;
      size_t sizeVec;
      int _rows, _cols;

      _rows = ((rows == -1) ? m_rows : rows);
      _cols = ((cols == -1) ? m_cols : cols);

      sizeVec = _rows * _cols * sizeof(T);

      cudaSetDevice(m_deviceID);

#ifdef USE_PINNED_MEMORY
      if(m_usePitch)
         err = cudaMemcpy2DAsync(m_deviceDataPointer,m_pitch*sizeof(T),m_hostDataPointer,_cols*sizeof(T), _cols*sizeof(T), _rows, cudaMemcpyHostToDevice, m_dev->stream);
      else
         err = cudaMemcpyAsync(m_deviceDataPointer, m_hostDataPointer, sizeVec, cudaMemcpyHostToDevice, m_dev->stream);
#else
      if(m_usePitch)
         err = cudaMemcpy2D(m_deviceDataPointer,m_pitch*sizeof(T),m_hostDataPointer,_cols*sizeof(T), _cols*sizeof(T), _rows, cudaMemcpyHostToDevice);
      else
         err = cudaMemcpy(m_deviceDataPointer, m_hostDataPointer, sizeVec, cudaMemcpyHostToDevice);
#endif

      if(err != cudaSuccess)
      {
         std::cerr<<"Error copying data to device\n" <<cudaGetErrorString(err) <<"\n";
      }

      deviceDataHasChanged = false;
   }
}

template <typename T>
void DeviceMemPointer_Matrix_CU<T>::copyDeviceToHost(int rows, int cols) const
{
   if(deviceDataHasChanged && m_hostDataPointer != NULL)
   {
      DEBUG_TEXT_LEVEL2("DEVICE_TO_HOST!!!\n")

      cudaError_t err;
      size_t sizeVec;

      int _rows, _cols;

      _rows = ((rows == -1) ? m_rows : rows);
      _cols = ((cols == -1) ? m_cols : cols);

      sizeVec = _rows * _cols * sizeof(T);

      cudaSetDevice(m_deviceID);

#ifdef USE_PINNED_MEMORY
      if(m_usePitch)
         err = cudaMemcpy2DAsync(m_hostDataPointer,_cols*sizeof(T),m_deviceDataPointer,m_pitch*sizeof(T), _cols*sizeof(T), _rows, cudaMemcpyDeviceToHost, m_dev->stream);
      else
         err = cudaMemcpyAsync(m_hostDataPointer, m_deviceDataPointer, sizeVec, cudaMemcpyDeviceToHost, m_dev->stream);
      cudaStreamSynchronize(m_dev->stream);
#else
      if(m_usePitch)
         err = cudaMemcpy2D(m_hostDataPointer,_cols*sizeof(T),m_deviceDataPointer,m_pitch*sizeof(T), _cols*sizeof(T), _rows, cudaMemcpyDeviceToHost);
      else
         err = cudaMemcpy(m_hostDataPointer, m_deviceDataPointer, sizeVec, cudaMemcpyDeviceToHost);
#endif

      if(err != cudaSuccess)
      {
         std::cerr<<"Error copying data from device: " <<cudaGetErrorString(err) <<"\n";
      }

      deviceDataHasChanged = false;
   }
}

template <typename T>
T* DeviceMemPointer_Matrix_CU<T>::getDeviceDataPointer() const
{
   return m_deviceDataPointer;
}

template <typename T>
int DeviceMemPointer_Matrix_CU<T>::getDeviceID() const
{
   return m_deviceID;
}

template <typename T>
void DeviceMemPointer_Matrix_CU<T>::changeDeviceData()
{
   deviceDataHasChanged = true;
}

}

#endif

#endif