Beignet Haswell Bug

beignet-bug.c

#include <CL/cl.h>
#include <stdio.h>

int
main()
{
    cl_int ret;

    cl_platform_id platform_id = NULL;
    cl_uint num_platforms = 0;
    clGetPlatformIDs(1, &platform_id, &num_platforms);

    cl_device_id device_id = NULL;
    cl_uint num_devices = 0;
    clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1,
                   &device_id, &num_devices);
    cl_context context = clCreateContext(NULL, 1, &device_id,
                                         NULL, NULL, &ret);
    cl_command_queue queue = clCreateCommandQueue(context, device_id, 0, &ret);

    cl_mem buf = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                                sizeof(float) * 100, NULL, &ret);
    const char *kernel_src = ""
        "__kernel void fill_one(__global float *out__base, long out__offset, long n)\n"
        "{\n"
        "    int lid = get_local_id(0);\n"
        "    int gsize = get_global_size(0);\n"
        "    int work_group_start = get_local_size(0) * get_group_id(0);\n"
        "    long i;\n"
        "    __global float *out = (__global float *) ((__global char *) out__base + out__offset);\n"
        "    for (i = work_group_start + lid; i < n; i += gsize) {\n"
        "        out[i] = 1;\n"
        "    }\n"
        "}\n";
    cl_program program =
        clCreateProgramWithSource(context, 1, &kernel_src, NULL, &ret);
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    cl_kernel kernel = clCreateKernel(program, "fill_one", &ret);
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf);
    cl_long out__offset = 0;
    clSetKernelArg(kernel, 1, sizeof(cl_long), &out__offset);
    cl_long n = 100;
    clSetKernelArg(kernel, 2, sizeof(cl_long), &n);
    size_t global = 128;
    size_t local = 32;
    size_t offset = 0;
    ret = clEnqueueNDRangeKernel(queue, kernel, 1, &offset, &global,
                                 &local, 0, NULL, NULL);
    printf("%d\n", ret);
    clFlush(queue);

    float host_buf[100];
    clEnqueueReadBuffer(queue, buf, CL_TRUE, 0, sizeof(host_buf), host_buf,
                        0, NULL, NULL);
    for (int i = 0;i < 100;i++) {
        printf("%f, ", host_buf[i]);
    }
    printf("\n");

    clFinish(queue);

    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseMemObject(buf);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);

    return 0;
}