/* * Each thread calculates a pixel component(rgba), by applying a filter * on group of 8 neighbouring pixels in both x and y directions. */ // // DECLARATION OF CONSTANTS // TODO : pass all the arguments with clBuildProgram(..) // like subSamplingDivider ("divider" here) // see page 240 of openCL 1.1 specification __constant sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; __constant int divider = 4; // sub-sampling decimation factor __constant int kernelRadius = 2; // radius of the kernel for all convolution filters // value of pixel data (from image2d_t) __constant int DATA_MIN = 0; __constant int DATA_MAX = 255; // the width for searching the disparity between 2 blocks of pixels (see report smart OSD insertion section 2.5) __constant int DISP_LOOKUP_WIDTH = 16; __constant int GRADIANT_THRESHOLD = 5; __constant uint4 GRADIANT_UNDERFLOW_TAG = (uint4)(0, 255, 0, 1); // RGBA __constant int SAD_THRESHOLD = 200; __constant uint4 SAD_OVERFLOW_TAG = (uint4)(255, 0, 255, 1); // RGBA //////////////////////////////////////////////////////////////////////// // HELPER countingSort: // Sort by using programming tips meaning that it's a bucket sort but // the elements of the array MUST be unsigned and the range of values known // see : http://www.algorithmist.com/index.php/Counting_sort // TODO : parallelize the code ? // //////////////////////////////////////////////////////////////////////// void countingSort(uint * a, uint a_size) { // there's 256 possible elements in 8bit uint count[256] = {0}; // Scan the array and count the number of times each value appears for( uint l = 0; l < a_size; l++ ) count[ a[ l ] ]++; // Fill the array with the number of 0's that were counted, followed by the number of 1's, and then 2's and so on uint n = 0; for( uint i = 0; i < 256; i++ ) for( uint j = 0; j < count[ i ]; j++ ) a[ n++ ] = i; } //////////////////////////////////////////////////////////////////////// // KERNEL medianfilter: // we apply a vertical median filter to remove noise from image // by performing a sort and searching the median value (numberOfElements / 2) // //////////////////////////////////////////////////////////////////////// __kernel void medianfilter(__read_only image2d_t src, __write_only image2d_t dest) { int2 coord = (int2)(get_global_id(0), get_global_id(1)); int radius = kernelRadius; // MODIFY the num. of elements in lumValues if you change this // because OpenCL does not allow to declare variable size... uint numValues = (radius * 2) + 1; // we need to take into account the full range (-radius; radius) uint lumValues[(2 * 2) + 1] = {0}; uint4 pixel = (uint4)0; uint i = 0; // we take the first component (Y) of every neighbours for (int dy=-radius; dy<=radius; dy++) { pixel = read_imageui(src, imageSampler, (int2)(coord.x, coord.y + dy)); lumValues[i] = pixel.x; ++i; } // then, we sort the elements by ascending order countingSort(lumValues, numValues); // finally, we take the (numValues / 2)th element of the sorted array uint final = lumValues[numValues/2]; write_imageui(dest, coord, (uint4)(final, final, final, 1)); // testing purpose ! //write_imageui(dest, coord, read_imageui(src, imageSampler, coord.xy)); } //////////////////////////////////////////////////////////////////////// // KERNEL lowestdisparity: // // //////////////////////////////////////////////////////////////////////// __kernel void lowestdisparity(__global uint * src, __global uint *lowestDispValue) { int2 coord = (int2)(get_global_id(0), get_global_id(1)); } //////////////////////////////////////////////////////////////////////// // KERNEL insertOSD: // // //////////////////////////////////////////////////////////////////////// __kernel void insertOSD(__write_only image2d_t srcLeft, __write_only image2d_t srcRight, __read_only image2d_t logo, uint lowestDisparity, uint xOffset, uint yOffset, uint width, uint height) { int2 coord = (int2)(get_global_id(0), get_global_id(1)); }