The kernel that is generated for this test: https://github.com/ddemidov/vexcl/blob/a3d940b97/tests/stencil.cpp#L161

user_defined_stencil.cl

#if defined(cl_khr_fp64)
#  pragma OPENCL EXTENSION cl_khr_fp64: enable
#elif defined(cl_amd_fp64)
#  pragma OPENCL EXTENSION cl_amd_fp64: enable
#endif
typedef double real;
real read_x(
    long g_id,
    ulong n,
    char has_left, char has_right,
    int lhalo, int rhalo,
    global const real *xloc,
    global const real *xrem
    )
{
    if (g_id >= 0 && g_id < n) {
        return xloc[g_id];
    } else if (g_id < 0) {
        if (has_left)
            return (lhalo + g_id >= 0) ? xrem[lhalo + g_id] : 0;
        else
            return xloc[0];
    } else {
        if (has_right)
            return (g_id < n + rhalo) ? xrem[lhalo + g_id - n] : 0;
        else
            return xloc[n - 1];
    }
}
real stencil_oper(local real *X) {
return sin(X[1] - X[0]) + sin(X[0] - X[-1]);
}
kernel void convolve(
    ulong n,
    char has_left,
    char has_right,
    int lhalo, int rhalo,
    global const real *xloc,
    global const real *xrem,
    global real *y,
    real alpha, real beta,
    local real *X
    )
{
    int l_id       = get_local_id(0);
    int block_size = get_local_size(0);
    long g_id      = get_global_id(0);
    for(int i = 0, j = g_id - lhalo; i < 1 + lhalo + rhalo; i++, j++)
        X[i] = read_x(j, n, has_left, has_right, lhalo, rhalo, xloc, xrem);
    if (g_id < n) {
        real sum = stencil_oper(X + lhalo);
        if (alpha)
            y[g_id] = alpha * y[g_id] + beta * sum;
        else
            y[g_id] = beta * sum;
    }
}