
#define VALTYPE float

int ToGlobalRow( int gidRow, int lszRow, int lidRow );
int ToGlobalCol( int gidCol, int lszCol, int lidCol );
int ToFlatHaloedIdx( int row, int col, int rowPitch ); 
int ToFlatIdx( int row, int col, int pitch );

__kernel
void
CopyRect( __global VALTYPE* dest,
            int doffset,
            int dpitch,
            __global VALTYPE* src,
            int soffset,
            int spitch,
            int width,
            int height )
{
    int gid = get_group_id(0);
    int lid = get_local_id(0);
    int gsz = get_global_size(0);
    int lsz = get_local_size(0);
    int grow = gid * lsz + lid;

    if( grow < height )
    {
        for( int c = 0; c < width; c++ )
        {
            (dest + doffset)[ToFlatIdx(grow,c,dpitch)] = (src + soffset)[ToFlatIdx(grow,c,spitch)];
        }
    }
}
