#ifndef __CUARRAY_IMPL_HPP__
#define __CUARRAY_IMPL_HPP__

namespace amscuda
{

// New Version cuarray<T>
// simpler, less crap going on

template<typename T> __device__ __host__ cuarray<T>::cuarray()
{
    length = 0;
    data = NULL;
}

template<typename T> __device__ __host__ cuarray<T>::~cuarray()
{
    if(data!=NULL)
    {
        delete[] data; data = NULL;
    }
    length = 0;
}

template<typename T> __device__ __host__ int cuarray<T>::resize(const int _length)
{
    int ret = 0;

    T *newbuffer = NULL;

    if(length==_length)
    {
        //do nothing
        ret = 1;
        return ret;
    }
    if(_length<=0)
    {
        if(data!=NULL)
        {
            delete[] data;
            data = NULL;
        }
        length = 0;
        ret = 1;
    }

    newbuffer = new T[_length];
    if(newbuffer==NULL)
    {
        ret = -1; //failed to allocate memory
        return ret;
    }

    int I;
    T def;

    if(data!=NULL)
    {
        for(I=0;I<length&&I<_length;I++)
        {
            newbuffer[I] = data[I];
        }
        for(I=length;I<_length;I++)
        {
            newbuffer[I] = def;
        }
        delete[] data; data=NULL;
    }
    else
    {
        for(I=0;I<_length;I++)
        {
            newbuffer[I] = def;
        }
    }

    data = newbuffer;
    length = _length;
    ret = 1;

    return ret;
}


template<typename T> __host__ int cuarray<T>::device_send(cuarray<T> **dptr)
{
    int ret = 0;
    int dlength;

    if(*dptr==NULL)
    {
        ret = _device_send_overwrite(dptr);
    }
    else
    {
        dlength = device_length(*dptr);
        if(dlength=length)
        {
            ret = _device_send_copy(*dptr);
        }
        else
        {
            ret = _device_send_overwrite(dptr);
        }
    }

    return ret;
}

template<typename T> __host__ int cuarray<T>::_device_send_overwrite(cuarray<T> **dptr)
{
    int ret = 0;
    cuarray<T> dlocal;
    cudaError_t err = cudaSuccess;
    device_free(dptr);
    
    if(length>=0 && data!=NULL)
    {
        err = cudaMalloc(dptr,sizeof(cuarray<T>));
        if(err==cudaSuccess)
        {
            err = cudaMalloc(&(dlocal.data),sizeof(T)*length);
            dlocal.length = length;

            if(err==cudaSuccess)
            {
                cudaMemcpy(*dptr,&dlocal,sizeof(cuarray<T>),cudaMemcpyHostToDevice);
                if(data!=NULL)
                    err = cudaMemcpy(dlocal.data,data,sizeof(T)*length,cudaMemcpyHostToDevice);
                else
                    err = cudaSuccess;
                if(err==cudaSuccess)
                {
                    ret = 1;
                }
                else
                {
                    ret = -3;
                }
            }
            else
            {
                ret = -2;
            }
        }
        else
        {
            ret = -1;
        }
    }
    else
    {
        dlocal.data = NULL;
        dlocal.length = 0;
        err = cudaMalloc(dptr,sizeof(cuarray<T>));
        if(err==cudaSuccess)
        {
            cudaMemcpy(*dptr,&dlocal,sizeof(cuarray<T>),cudaMemcpyHostToDevice);
            ret = 1;
        }
        else
        {
            ret = -4;
        }
    }


    dlocal.data = NULL;
    dlocal.length = -1;

    return ret;
}

template<typename T> __host__ int cuarray<T>::_device_send_copy(cuarray<T> *dptr)
{
    int ret = 0;
    cudaError_t err = cudaSuccess;
    T* ddata = NULL;
    ddata = device_data_ptr(dptr);

    err = cudaMemcpy(ddata,data,sizeof(T)*length,cudaMemcpyHostToDevice);
    if(err==cudaSuccess)
    {
        ret = 1;
    }
    else
    {
        ret = -1;
    }

    return ret;
}

template<typename T> __host__ int cuarray<T>::device_pull(cuarray<T> *dptr)
{
    int ret = 0;
    int dlength;
    T* ddata;
    cudaError_t err;

    if(dptr==NULL) 
    {
        ret = -1; // null d pointer
        return ret;
    }

    dlength = device_length(dptr);
    if(dlength!=length)
    {
        this->resize(dlength);
    }

    ddata = device_data_ptr(dptr);

    if(length>0 && data!=NULL && ddata!=NULL)
    {
        err = cudaMemcpy(data,dptr,length*sizeof(T),cudaMemcpyDeviceToHost);
        if(err==cudaSuccess)
        {
            ret = 1;
        }
        else
        {
            ret = -2;
        }
    }

    return ret;
}

template<typename T> __host__ int cuarray<T>::device_free(cuarray<T> **dptr)
{
    int ret = 0;
    cuarray<T> dlocal;

    if(*dptr!=NULL)
    {
        cudaMemcpy(&dlocal,dptr,sizeof(cuarray<T>),cudaMemcpyDeviceToHost);
        if(dlocal.data!=NULL)
        {
            cudaFree(dlocal.data);
            dlocal.data = NULL;
        }

        cudaFree(*dptr);
        *dptr = NULL;
        ret = 1;
    }

    dlocal.data = NULL;
    dlocal.length = -1;

    return ret;
}

template<typename T> __host__ int cuarray<T>::device_length(cuarray<T> *dptr)
{
    int ret = -1;
    cuarray<T> dlocal;

    if(dptr==NULL)
    {
        return ret;
    }

    cudaMemcpy(&dlocal,dptr,sizeof(cuarray<T>),cudaMemcpyDeviceToHost);
    ret = dlocal.length;

    dlocal.data = NULL;
    dlocal.length = -1;

    return ret;
}

template<typename T> __host__ T* cuarray<T>::device_data_ptr(cuarray<T> *dptr)
{
    T* ret = NULL;
    cuarray<T> dlocal;

    if(dptr==NULL)
    {
        return ret;
    }

    cudaMemcpy(&dlocal,dptr,sizeof(cuarray<T>),cudaMemcpyDeviceToHost);
    ret = dlocal.data;

    dlocal.data = NULL;
    dlocal.length = -1;

    return ret;
}

template<typename T>  __device__ __host__ int cuarray<T>::size() const
{
    return this->length;
}

template<typename T>  __device__ __host__ T& cuarray<T>::at(const int I)
{
    return this->data[I];
}

template<typename T>  __device__ __host__ const T& cuarray<T>::at(const int I) const
{
    return this->data[I];
}

template<typename T> __device__ __host__ T& cuarray<T>::operator[](const int I)
{
    return this->data[I];
}

template<typename T> __device__ __host__ const T& cuarray<T>::operator[](const int I) const
{
    return this->data[I];
}

};


#endif