| /* |
| * function: kernel_demosaic |
| * input: image2d_t as read only |
| * output: image2d_t as write only |
| */ |
| |
| "__kernel void kernel_demosaic (__read_only image2d_t input, __write_only image2d_t output) " |
| "{ " |
| " int x = 2 * get_global_id (0); " |
| " int y = 2 * get_global_id (1); " |
| //" sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; " |
| " sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; " |
| |
| " " |
| " int x0 = x - 1; " |
| " int y0 = y - 1; " |
| " float4 p[16]; " |
| " p[0] = read_imagef (input, sampler, (int2)(x0, y0)); " |
| " p[1] = read_imagef (input, sampler, (int2)(x0 + 1, y0)); " |
| " p[2] = read_imagef (input, sampler, (int2)(x0 + 2, y0)); " |
| " p[3] = read_imagef (input, sampler, (int2)(x0 + 3, y0)); " |
| " p[4] = read_imagef (input, sampler, (int2)(x0, y0 + 1)); " |
| " p[5] = read_imagef (input, sampler, (int2)(x0 + 1, y0 + 1)); " |
| " p[6] = read_imagef (input, sampler, (int2)(x0 + 2, y0 + 1)); " |
| " p[7] = read_imagef (input, sampler, (int2)(x0 + 3, y0 + 1)); " |
| " p[8] = read_imagef (input, sampler, (int2)(x0, y0 + 2)); " |
| " p[9] = read_imagef (input, sampler, (int2)(x0 + 1, y0 + 2)); " |
| " p[10] = read_imagef (input, sampler, (int2)(x0 + 2, y0 + 2)); " |
| " p[11] = read_imagef (input, sampler, (int2)(x0 + 3, y0 + 2)); " |
| " p[12] = read_imagef (input, sampler, (int2)(x0, y0 + 3)); " |
| " p[13] = read_imagef (input, sampler, (int2)(x0 + 1, y0 + 3)); " |
| " p[14] = read_imagef (input, sampler, (int2)(x0 + 2, y0 + 3)); " |
| " p[15] = read_imagef (input, sampler, (int2)(x0 + 3, y0 + 3)); " |
| |
| //" for (int i = 0; i < 16; ++i) { " |
| //" p[i] = read_imagef (input, sampler, (int2)(x0 + i % 4, y0 + i / 4)); " |
| //" } " |
| " " |
| " float4 p00, p01, p10, p11; " |
| " p00.x = (p[4].x + p[6].x) / 2.0; " |
| " p00.y = (p[5].x * 4 + p[0].x + p[2].x + p[8].x + p[10].x) / 8.0; " |
| " p00.z = (p[1].x + p[9].x) / 2.0; " |
| " p01.x = p[6].x; " |
| " p01.y = (p[2].x + p[5].x + p[7].x+ p[10].x) / 4.0; " |
| " p01.z = (p[1].x + p[3].x + p[9].x+ p[11].x) / 4.0; " |
| " p10.x = (p[4].x + p[6].x + p[12].x+ p[14].x) / 4.0; " |
| " p10.y = (p[5].x + p[8].x + p[10].x+ p[13].x) / 4.0; " |
| " p10.z = p[9].x; " |
| " p11.x = (p[6].x + p[14].x) / 2.0; " |
| " p11.y = (p[10].x * 4 + p[5].x + p[7].x + p[13].x + p[15].x) / 8.0; " |
| " p11.z = (p[9].x + p[11].x) / 2.0; " |
| " " |
| " write_imagef (output, (int2)(x, y), p00); " |
| " write_imagef (output, (int2)(x + 1, y), p01); " |
| " write_imagef (output, (int2)(x, y + 1), p10); " |
| " write_imagef (output, (int2)(x + 1, y + 1), p11); " |
| "} " |
| |