// Copyright (C) 2011 Dmitri Nikulin // Copyright (C) 2011 Monash University // // Permission is hereby granted, free of charge, to any person // obtaining a copy of this software and associated documentation // files (the "Software"), to deal in the Software without // restriction, including without limitation the rights to use, // copy, modify, merge, publish, distribute, sublicense, and/or sell // copies of the Software, and to permit persons to whom the // Software is furnished to do so, subject to the following // conditions: // // The above copyright notice and this permission notice shall be // included in all copies or substantial portions of the Software. // // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, // EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES // OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND // NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT // HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, // WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING // FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR // OTHER DEALINGS IN THE SOFTWARE. kernel void fast_best( read_only image2d_t scores, global int2 * corners, global int2 * filtered, global int * icorner ) { // Prepare a suitable OpenCL image sampler. sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; // Use global work item as 1D offset into corners. int const ic = get_global_id(0); int2 const xy = corners[ic]; // Read the candidate score. int const p00 = read_imageui(scores, sampler, xy).x; // Read other scores in a tight square around the candidate score. int const p01 = read_imageui(scores, sampler, xy + (int2)(-1, -1)).x; int const p02 = read_imageui(scores, sampler, xy + (int2)( 0, -1)).x; int const p03 = read_imageui(scores, sampler, xy + (int2)( 1, -1)).x; int const p04 = read_imageui(scores, sampler, xy + (int2)(-1, 0)).x; int const p05 = read_imageui(scores, sampler, xy + (int2)( 0, 0)).x; int const p06 = read_imageui(scores, sampler, xy + (int2)( 1, 0)).x; int const p07 = read_imageui(scores, sampler, xy + (int2)(-1, 1)).x; int const p08 = read_imageui(scores, sampler, xy + (int2)( 0, 1)).x; int const p09 = read_imageui(scores, sampler, xy + (int2)( 1, 1)).x; // Select the maximum score. int sco = p00; sco = max(sco, p01); sco = max(sco, p02); sco = max(sco, p03); sco = max(sco, p04); sco = max(sco, p05); sco = max(sco, p06); sco = max(sco, p07); sco = max(sco, p08); sco = max(sco, p09); // Keep this score if it is as good as the maximum. if (p00 >= sco) { // Atomically append to filtered corner buffer. filtered[atom_inc(icorner)] = xy; } }