image_gpu.h

#pragma once

#include "image.h"
#include "image_allocator_gpu.h"
#include <thrust/device_ptr.h>

template<typename type, int dims> class ndarray_ref;

namespace iu {


template<typename PixelType, bool zero_reminder>
struct use_PixelType{
      static int stride(int pitch_){
          return pitch_/sizeof(PixelType);
      }
};

template<typename PixelType>
struct use_PixelType<PixelType,false>{
      static int __attribute__((deprecated("this function may return wrong result in case pitch_ is not divisible by sizeof(PixelType) (typically the case when using _C3 images)"))) stride(int pitch_){
          if(pitch_ % sizeof(PixelType) != 0){
              throw std::runtime_error("bad! your are using wrong stride, because pitch is not divisible (are you using _C3 images? You shouldn't)");
          };
          return pitch_/sizeof(PixelType);
      }
};


template<typename PixelType, class Allocator>
class ImageGpu : public Image
{
public:
  typedef PixelType pixel_type;

  ImageGpu() :
    Image(),
    data_(0), pitch_(0), ext_data_pointer_(false), texture_(0)
  {
  }

  virtual ~ImageGpu()
  {
    if(!ext_data_pointer_)
    {
      // do not delete externally handeled data pointers.
      Allocator::free(data_);
      data_ = 0;
    }
    pitch_ = 0;

    if (texture_)
      IU_CUDA_SAFE_CALL(cudaDestroyTextureObject(texture_));
  }

  ImageGpu(unsigned int _width, unsigned int _height) :
      Image(_width, _height), data_(0), pitch_(0),
      ext_data_pointer_(false), texture_(0)
  {
    data_ = Allocator::alloc(this->size(), &pitch_);
  }

  ImageGpu(const iu::Size<2>& size) :
      Image(size), data_(0), pitch_(0),
      ext_data_pointer_(false), texture_(0)
  {
    data_ = Allocator::alloc(size, &pitch_);
  }

  ImageGpu(PixelType* _data, unsigned int _width, unsigned int _height,
           size_t _pitch, bool ext_data_pointer = false) :
      Image(_width, _height), data_(0), pitch_(0), ext_data_pointer_(ext_data_pointer), texture_(0)
  {
    if(ext_data_pointer_)
    {
      // This uses the external data pointer as internal data pointer.
      data_ = _data;
      pitch_ = _pitch;
    }
    else
    {
      // allocates an internal data pointer and copies the external data onto it.
      if(_data == 0)
        return;

      data_ = Allocator::alloc(this->size(), &pitch_);
      Allocator::copy(_data, _pitch, data_, pitch_, this->size());
    }
  }

  PixelType getPixel(unsigned int x, unsigned int y)
  {
    PixelType value;
    IU_CUDA_SAFE_CALL(cudaMemcpy2D(&value, sizeof(PixelType), &data_[y*stride()+x], pitch_,
                 sizeof(PixelType), 1, cudaMemcpyDeviceToHost));
    return value;
  }

  virtual size_t bytes() const
  {
    return height()*pitch_;
  }

  virtual size_t pitch() const
  {
    return pitch_;
  }

  virtual size_t  stride() const
  {
    return use_PixelType<PixelType, sizeof(PixelType) % 32 == 0 ||sizeof(PixelType)==1 || sizeof(PixelType) == 2 || sizeof(PixelType) == 4 || sizeof(PixelType) == 8 || sizeof(PixelType) == 16 >::stride(pitch_);
    //return pitch_/sizeof(PixelType);
  }

  virtual unsigned int bitDepth() const
  {
    return 8*sizeof(PixelType);
  }

  virtual bool onDevice() const
  {
    return true;
  }

  PixelType* data(int ox = 0, int oy = 0)
  {
    //return &data_[oy * stride() + ox];
      return (PixelType*)(  (char*)data_ + pitch_*oy + ox*sizeof(PixelType) );
  }

  const PixelType* data(int ox = 0, int oy = 0) const
  {
      return const_cast<ImageGpu*>(this)->data(ox,oy);
    //return reinterpret_cast<const PixelType*>(
     //   &data_[oy * stride() + ox]);
  }

  thrust::device_ptr<PixelType> begin(void)
  {
      return thrust::device_ptr<PixelType>(data());
  }

  thrust::device_ptr<PixelType> end(void)
  {
      //return thrust::device_ptr<PixelType>(data()+stride()*height());
      return thrust::device_ptr<PixelType>(data(width()-1,height()-1) + 1);
  }

  void prepareTexture(cudaTextureReadMode readMode = cudaReadModeElementType,
                      cudaTextureFilterMode filterMode = cudaFilterModeLinear,
                      cudaTextureAddressMode addressMode = cudaAddressModeClamp)
  {
      if (texture_)             // delete if already exists
        IU_CUDA_SAFE_CALL(cudaDestroyTextureObject(texture_));

      cudaResourceDesc resDesc;
      memset(&resDesc, 0, sizeof(resDesc));

      cudaTextureDesc texDesc;
      memset(&texDesc, 0, sizeof(texDesc));

      resDesc.resType = cudaResourceTypePitch2D;
      resDesc.res.pitch2D.devPtr = data();
      resDesc.res.pitch2D.pitchInBytes = pitch();
      resDesc.res.pitch2D.width = width();
      resDesc.res.pitch2D.height = height();
      resDesc.res.pitch2D.desc = cudaCreateChannelDesc<PixelType>();

      texDesc.readMode = readMode;
      texDesc.normalizedCoords = (addressMode == cudaAddressModeClamp) ? false : true;
      texDesc.addressMode[0] = addressMode;
      texDesc.addressMode[1] = addressMode;
      texDesc.filterMode = filterMode;

      IU_CUDA_SAFE_CALL(cudaCreateTextureObject(&texture_, &resDesc, &texDesc, NULL));
  }

  cudaTextureObject_t getTexture()
  {
      if (!texture_)        // create texture object implicitly
          this->prepareTexture();

      return texture_;
  }

  inline cudaTextureObject_t getTexture() const
  {
      if (!texture_)
          throw IuException("Warning: getTexture() on const image requires explicit call to prepareTexture(),"
                            " returned cudaTextureObject will be invalid\n", __FILE__, __FUNCTION__, __LINE__);

      return texture_;
  }

  struct KernelData
  {
      PixelType* data_;
      int width_;
      int height_;
      int stride_;

      __device__ PixelType& operator()(int x, int y)
      {
          return data_[y*stride_ + x];
      }

      __host__ KernelData(const ImageGpu<PixelType, Allocator> &im)
          : data_(const_cast<PixelType*>(im.data())), width_(im.width()), height_(im.height()),
            stride_(im.stride())
      { }
  };

  ndarray_ref<PixelType,2> ref() const;

   ImageGpu(const ndarray_ref<PixelType,2> &x);

protected:
  PixelType* data_;
  size_t pitch_;
  bool ext_data_pointer_;
  cudaTextureObject_t texture_;

private:
  ImageGpu(const ImageGpu&);
  ImageGpu& operator=(const ImageGpu&);
};

} // namespace iu