medical image analysis ocltools

// This file is part of fluid3d - a software for image registration
//
// Copyright (c) Madrid 2013 Gert Wollny
//
// fluid3d is free software; you can redistribute it and/or modify
// it under the terms of the GNU General Public License as published by
// the Free Software Foundation; either version 3 of the License, or
// (at your option) any later version.
//
// This program 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 General Public License for more details.
//
// You should have received a copy of the GNU General Public License
// along with fluid3d; if not, see <http://www.gnu.org/licenses/>.


#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; 

__kernel void downscale3d( __write_only image3d_t  output,  
			   __read_only image3d_t  input, 
			   __constant float *kern, int4 hgw, 
			   float4 scales)
{
    // this version runs the filter in a non-separable way 
    int4 outpos = {get_global_id(0), get_global_id(1), get_global_id(2), 0}; 
    
    float4 opf = convert_float4(outpos); 
    
    int4 inpos = convert_int4_rte(opf * scales);
    int4 delta = {0, 0, 0, 0};
    int ki = 0; 
    float4 out_value = 0.0f; 
    for (delta.z = - hgw.z ; delta.z <= hgw.z; ++delta.z) 
        for (delta.y = - hgw.y ; delta.y <= hgw.y; ++delta.y) 
            for (delta.x = - hgw.x ; delta.x <= hgw.x; ++delta.x, ++ki) {
		    out_value += kern[ki] * read_imagef(input, inpos + delta); 
            }
    write_imagef(output, outpos, out_value); 
}