/*
 * Copyright (C) 2011-2014 Karlsruhe Institute of Technology
 *
 * This file is part of Ufo.
 *
 * This library is free software: you can redistribute it and/or
 * modify it under the terms of the GNU Lesser General Public
 * License as published by the Free Software Foundation, either
 * version 3 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Lesser General Public License for more details.
 *
 * You should have received a copy of the GNU Lesser General Public
 * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
 */

 __kernel void backproject (global float *volume,
                            __read_only image2d_t projection,
                            const sampler_t sampler,
                            const int2 x_region,
                            const int2 y_region,
                            const int2 z_region,
                            const float8 tmatrix,
                            const int2 offset,
                            const uint cumulate)
{
     int3 id = (int3) (get_global_id (0), get_global_id (1), get_global_id (2));
     float2 pixel;
     float3 voxel;
 
     voxel.x = x_region.x + id.x * x_region.y;
     voxel.y = y_region.x + id.y * y_region.y;
     voxel.z = z_region.x + id.z * z_region.y;
 
     pixel.x = voxel.x * tmatrix.s0 + voxel.y * tmatrix.s1 + tmatrix.s3 - offset.x;
     pixel.y = voxel.x * tmatrix.s4 + voxel.y * tmatrix.s5 + voxel.z * tmatrix.s6 + tmatrix.s7 - offset.y;
 
    if (cumulate) {
         volume[id.z * get_global_size (0) * get_global_size (1) +
                id.y * get_global_size (0) + id.x] += read_imagef (projection, sampler, pixel).x;
    }
    else {
         volume[id.z * get_global_size (0) * get_global_size (1) +
                id.y * get_global_size (0) + id.x] = read_imagef (projection, sampler, pixel).x;
    }
 }