src/PSkelArray.h

Go to the documentation of this file.

//-------------------------------------------------------------------------------
// Copyright (c) 2015, ICEI - PUC Minas
// All rights reserved.
// 
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// 
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
// 
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
// 
// 3. Neither the name of the copyright holder nor the names of its contributors
// may be used to endorse or promote products derived from this software without
// specific prior written permission.
// 
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//-------------------------------------------------------------------------------

#ifndef PSKEL_ARRAY_H
#define PSKEL_ARRAY_H

#include <cuda.h>
#include "PSkelDefs.h"

namespace PSkel{

template<typename T>
class ArrayBase{
private:
        //variables that hold the real boundaries (total allocated data.)
        size_t realWidth, realHeight, realDepth;
        //variables that hold the "virtual" array boundaries (important for sliced arrays.)
        size_t width, height,depth;
        //offsets for the sliced array.
        size_t widthOffset, heightOffset, depthOffset;
        //host and device (GPU memory) pointers
        T *hostArray, *deviceArray;
protected:
        __device__ __forceinline__ T & deviceGet(size_t h,size_t w,size_t d) const ;

        __host__ __forceinline__ T & hostGet(size_t h,size_t w,size_t d) const ;

        ArrayBase(size_t width, size_t height, size_t depth);
public: 
        void deviceAlloc();

        void deviceFree();

        void hostAlloc(size_t width, size_t height, size_t depth);

        void hostAlloc();

        void hostFree();
        
        __device__ __host__ size_t getWidth() const;
        
        __device__ __host__ size_t getHeight() const;
        
        __device__ __host__ size_t getDepth() const;
        
        __device__ __host__ size_t memSize() const;

        __device__ __host__ size_t size() const;

        __device__ __host__ size_t realSize() const;

        template<typename Arrays>
        void hostSlice(Arrays array, size_t widthOffset, size_t heightOffset, size_t depthOffset, size_t width, size_t height, size_t depth);

        template<typename Arrays>
        void hostClone(Arrays array);
        
        template<typename Arrays>
        void hostMemCopy(Arrays array);

        void copyToDevice();

        template<typename Arrays>
        void copyFromDevice(Arrays array);

        void copyToHost();

        __device__ __host__ operator bool() const ;
};

//*******************************************************************************************
// Array 3D
//*******************************************************************************************

template <typename T>
class Array3D: public ArrayBase<T>{
public:
        Array3D();
        
        /*
        //TODO O kernel cuda não aceita structs com destrutores. Corrigir nas próximas versões
        ~Array3D(){
                free(hostArray);
                cudaFree(deviceArray);
        }*/
        
        Array3D(size_t width, size_t height, size_t depth);
        
        __attribute__((always_inline)) __forceinline__ __device__ __host__ T & operator()(size_t h,size_t w,size_t d) const ;
        
};

//*******************************************************************************************
// Array 2D
//*******************************************************************************************

template<typename T>
class Array2D: public ArrayBase<T>{
public: 
        Array2D();
        
        Array2D(size_t width, size_t height);
        
        __attribute__((always_inline)) __forceinline__ __device__ __host__ T & operator()(size_t h, size_t w) const ;
        
};

//*******************************************************************************************
// Array 1D
//*******************************************************************************************

template<typename T>
class Array: public ArrayBase<T>{
public:
        Array();
        
        Array(size_t size);

        __attribute__((always_inline)) __forceinline__ __device__ __host__ T & operator()(size_t w) const;      
};

}//end namespace

#include "PSkelArray.hpp"

#endif