opencl pencil kernel using local memory in planes

pencil_kernel_with_local_mem.cl

// <file: kernel_0.cl>
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#define encode514_514_514(x0, x1, x2) (264196 * (long) (x0) + 514 * (long) (x1) + 1 * (long) (x2))
#define encode514_514(x0, x1) (514 * (long) (x0) + 1 * (long) (x1))
#define encode10_18(x0, x1) (18 * (long) (x0) + 1 * (long) (x1))
__kernel void kernel_0(__global double* out, __global double* mesh) {

    __local double local_buf_0[180];
    __local double local_buf_1[180];
    __local double local_buf_2[180];
    __local double* plane_0 = local_buf_0;
    __local double* plane_1 = local_buf_1;
    __local double* plane_2 = local_buf_2;
    __local double* temp_plane;
    size_t tile_id_1 = get_group_id(0);
    size_t tile_id_2 = get_group_id(1);
    size_t packed_global_id_1 = get_global_id(0);
    size_t packed_global_id_2 = get_global_id(1);
    size_t packed_local_id_1 = get_local_id(0);
    size_t packed_local_id_2 = get_local_id(1);
    size_t thread_id = packed_local_id_1 * 16 + packed_local_id_2;
    size_t group_id_0 = get_group_id(0);
    size_t group_id_1 = get_group_id(1);
    size_t index_0;
    size_t index_1;
    size_t index_2;
    size_t local_index_0;
    size_t local_index_1;
    size_t local_index_2;
    size_t dim_0_offsets[] = {1};
    size_t dim_0_strides[] = {1};
    size_t dim_1_offsets[] = {1};
    size_t dim_1_strides[] = {1};
    index_1 = (packed_global_id_1 % 512) * dim_1_strides[(packed_global_id_1 / 512)] + dim_1_offsets[(packed_global_id_1 / 512)];
    local_index_1 = (packed_local_id_1 % 512) * dim_1_strides[(packed_global_id_1 / 512)] + dim_1_offsets[(packed_global_id_1 / 512)];
    size_t dim_2_offsets[] = {1};
    size_t dim_2_strides[] = {1};
    index_2 = (packed_global_id_2 % 512) * dim_2_strides[(packed_global_id_2 / 512)] + dim_2_offsets[(packed_global_id_2 / 512)];
    local_index_2 = (packed_local_id_2 % 512) * dim_2_strides[(packed_global_id_2 / 512)] + dim_2_offsets[(packed_global_id_2 / 512)];
    //
    // Fill the first local memory planes
    //
    plane_1[(thread_id + 0)] = mesh[encode514_514_514(0, tile_id_1 * 8 + ((thread_id + 0) / 18), tile_id_2 * 16 + ((thread_id + 0) % 18))];
    if (thread_id + 128 < 180) {
        plane_1[(thread_id + 128)] = mesh[encode514_514_514(0, tile_id_1 * 8 + ((thread_id + 128) / 18), tile_id_2 * 16 + ((thread_id + 128) % 18))];
    };
    barrier(CLK_LOCAL_MEM_FENCE);
    plane_2[(thread_id + 0)] = mesh[encode514_514_514(1, tile_id_1 * 8 + ((thread_id + 0) / 18), tile_id_2 * 16 + ((thread_id + 0) % 18))];
    if (thread_id + 128 < 180) {
        plane_2[(thread_id + 128)] = mesh[encode514_514_514(1, tile_id_1 * 8 + ((thread_id + 128) / 18), tile_id_2 * 16 + ((thread_id + 128) % 18))];
    };
    barrier(CLK_LOCAL_MEM_FENCE);
    for (index_0 = 1; index_0 < 513; index_0 ++) {
        temp_plane = plane_0;
        plane_0 = plane_1;
        plane_1 = plane_2;
        plane_2 = temp_plane;
        plane_2[(thread_id + 0)] = mesh[encode514_514_514(index_0 + 1, tile_id_1 * 8 + ((thread_id + 0) / 18), tile_id_2 * 16 + ((thread_id + 0) % 18))];
        if (thread_id + 128 < 180) {
            plane_2[(thread_id + 128)] = mesh[encode514_514_514(index_0 + 1, tile_id_1 * 8 + ((thread_id + 128) / 18), tile_id_2 * 16 + ((thread_id + 128) % 18))];
        };
        barrier(CLK_LOCAL_MEM_FENCE);
        out[encode514_514_514(index_0, index_1, index_2)] = plane_0[encode10_18(local_index_1 + 0, local_index_2 + 0)] * 0.831077496446 + plane_1[encode10_18(local_index_1 + -1, local_index_2 + 0)] * 0.799991162981 + plane_1[encode10_18(local_index_1 + 0, local_index_2 + -1)] * 0.934757141219 + plane_1[encode10_18(local_index_1 + 0, local_index_2 + 0)] * 0.820808768568 + plane_1[encode10_18(local_index_1 + 0, local_index_2 + 1)] * 0.463035373069 + plane_1[encode10_18(local_index_1 + 1, local_index_2 + 0)] * 0.862325968572 + plane_2[encode10_18(local_index_1 + 0, local_index_2 + 0)] * 0.544520144954;
        barrier(CLK_LOCAL_MEM_FENCE);
    };

};