clEnqueueNDRangeKernel returns -59 on HARPv2

I write a short code in OpenCL. The host fuction is below

    float   h_gaus[3][3] = {{0.0625, 0.125, 0.0625}, {0.1250, 0.250, 0.1250}, {0.0625, 0.125, 0.0625}};
    cl_mem  d_gaus = clCreateBuffer(ocl.clContext, CL_MEM_READ_WRITE, 3 * 3 * sizeof(float), NULL, &clStatus);
    clStatus = clEnqueueWriteBuffer(ocl.clCommandQueue, d_gaus, CL_TRUE, 0, 3 * 3 * sizeof(float), h_gaus, 0, NULL, NULL);
    int rowsc, colsc, in_size;
#ifdef TEST_ROWS
    rowsc = TEST_ROWS;
#endif
 
#ifdef TEST_COLS
    colsc = TEST_COLS;
#endif
        in_size = rowsc * colsc * sizeof(unsigned char);
        // define input and output(local buffer)
        unsigned char *test_frame;
        test_frame = (unsigned char*) alignedMalloc(in_size);
 
        // initialize input frame
        for (int i = 0; i < rowsc; i++)
            for (int j = 0; j < colsc; j++)
            {
                test_frame[i * rowsc + j] = (unsigned char)100;
            }
 
        unsigned char *h_test;
        unsigned char *h_out;
        h_test = (unsigned char *)clSVMAllocAltera(ocl.clContext, 0, in_size, 1024);h_out = (unsigned char *)clSVMAllocAltera(ocl.clContext, 0, in_size, 1024); 
    memcpy(h_test, test_frame, in_size);
    size_t ls[2] = {(size_t)16, (size_t)16};
    size_t gs[2] = {(size_t)rowsc,(size_t)colsc};
    size_t *offset = NULL;
    clSetKernelArgSVMPointerAltera(ocl.clKernel_gauss, 0, (void*)h_test);
    clSetKernelArgSVMPointerAltera(ocl.clKernel_gauss, 1, (void*)h_out);
    clSetKernelArg(ocl.clKernel_gauss, 2, sizeof(int), &rowsc);
    clSetKernelArg(ocl.clKernel_gauss, 3, sizeof(int), &colsc);
    clSetKernelArg(ocl.clKernel_gauss, 4, (L_SIZE + 2) * (L_SIZE + 2) * sizeof(int), NULL);
    clSetKernelArg(ocl.clKernel_gauss, 5, sizeof(cl_mem), &d_gaus);
    clStatus = clEnqueueSVMMap(ocl.clCommandQueue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, (void *)h_test, in_size, 0, NULL, NULL);
                    CL_ERR();
    clStatus = clEnqueueSVMMap(ocl.clCommandQueue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, (void *)h_out, in_size, 0, NULL, NULL);
                    CL_ERR();
    clStatus = clEnqueueNDRangeKernel(ocl.clCommandQueue, ocl.clKernel_gauss, 2, offset, gs, ls, 0, NULL, NULL);

Here I define rowsc and colsc are both 64. The kernel function is below

__kernel void gaussian_kernel(__global unsigned char *data, __global unsigned char *out, int rows, int cols,
                              __local int *l_data, __global float *gaus)
{
    int sum = 0;
    int g_row = get_global_id(0);
    int g_col = get_global_id(1);
    int l_row = get_local_id(0) + 1;
    int l_col = get_local_id(1) + 1;
 
    int pos = g_row * cols + g_col;
    out[pos] = data[pos];
    // copy to local, the position itself
    l_data[l_row * (L_SIZE + 2) + l_col] = data[pos];
 
    // top most row
    if (l_row == 1)
    {
        // top most global row, fill it with 0
        l_data[0 * (L_SIZE + 2) + l_col] = (g_row==0 ? 0 : data[pos - cols]);
 
        // top left
        if (l_col == 1)
            l_data[0 * (L_SIZE + 2) + 0] = (g_row==0 ? 0 : data[pos - cols - 1]);
 
        // top right
        else if (l_col == L_SIZE)
            l_data[0 * (L_SIZE + 2) + L_SIZE + 1] = (g_row==0 ? 0 : data[pos - cols + 1]);
    }
 
    // bottom most row
    else if (l_row == L_SIZE)
    {
        l_data[(L_SIZE + 1) * (L_SIZE + 2) + l_col] = (g_row==rows-1 ? 0 : data[pos + cols]);
 
        // bottom left
        if (l_col == 1)
            l_data[(L_SIZE + 1) * (L_SIZE + 2) + 0] = (g_row==rows-1 ? 0 : data[pos + cols - 1]);
 
        // bottom right
        else if (l_col == L_SIZE)
            l_data[(L_SIZE + 1) * (L_SIZE + 2) + L_SIZE + 1] = (g_row == rows-1 ? 0 : data[pos + cols + 1]);
    }
 
    // left most col
    if (l_col == 1)
        l_data[l_row * (L_SIZE + 2) + 0] = (g_col == 0 ? 0 : data[pos - 1]);
    // right most col
    else if (l_col == L_SIZE)
        l_data[l_row * (L_SIZE + 2) + L_SIZE + 1] = (g_col == cols-1 ? 0 : data[pos + 1]);
 
    barrier(CLK_LOCAL_MEM_FENCE);
 
    // compute convolution
    for (int i = 0; i < 3; i++)
    {
        for (int j = 0; j < 3; j++)
        {
            sum += gaus[i * 3 + j] * l_data[(i + l_row - 1) * (L_SIZE + 2) + j + l_col - 1];
        }
    }
 
    out[pos] = min(255, max(0, sum));
}

I didn't add any restriction to neither max work group size nor max work items. I didn't assign the kernel a "task" attribute, either. But when I try to enqueue a 2D NDRange kernel, it returns -59. I didn't find any introduction about this condition on khronos' group website. So I wonder where is wrong in my code? Or any point I didn't get about HARPv2?