Attachment 127966 Details for Bug 98717 – Minimal example to reproduce

Minimal example to reproduce

packbug.c (text/x-c), 3.73 KB, created by Kai on 2016-11-14 12:15:16 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: Kai

Created: 2016-11-14 12:15:16 UTC

Size: 3.73 KB

patch

obsolete

>#include <CL/cl.h>
>#include <stdio.h>
>#include <assert.h>
>#include <string.h>
>
>#define NELEMS(x) (sizeof(x) / sizeof(*x))
>
>#define PLATFORM 0
>#define DEVICE 0
>
>const char src[] =
>    "// Works if not 'packed'\n"
>    "struct Node {\n"
>    "    unsigned int right_edge;\n"
>    "    unsigned int children[2];\n"
>    "} __attribute__((packed));\n"
>    "\n"
>    "kernel void generate(global struct Node * const nodes) {\n"
>    "    const unsigned int idx = get_global_id(0);\n"
>    "    nodes[idx].right_edge = 0;\n"
>    "    nodes[idx].children[0] = get_global_size(0);\n"
>    "    nodes[idx].children[1] = get_global_size(0);\n"
>    "    //Works if right_edge is set here instead\n"
>    "    //nodes[idx].right_edge = 0;\n"
>    "}";
>
>struct Node {
>    cl_uint right_edge;
>    cl_uint children[2];
>} __attribute__((packed));
>
>void printNode(struct Node node) {
>    puts("Node {");
>    printf("\tright_edge = %u,\n", node.right_edge);
>    printf("\tchildren = {%u, %u},\n", node.children[0], node.children[1]);
>    puts("}");
>}
>
>const size_t src_len = sizeof(src);
>const char * const kernels[] = {"generate"};
>const size_t n = 1;
>
>int main(void) {
>    cl_int err = -1;
>
>    cl_uint num_platforms;
>    clGetPlatformIDs(0, NULL, &num_platforms);
>    assert(num_platforms > PLATFORM);
>
>    cl_platform_id  * platforms = malloc(sizeof(*platforms) * num_platforms);
>    clGetPlatformIDs(num_platforms, platforms, NULL);
>
>    cl_uint num_devices;
>    clGetDeviceIDs(platforms[PLATFORM], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);
>    assert(num_devices >= DEVICE);
>
>    cl_device_id * devices = malloc(sizeof(*devices) * num_devices);
>    clGetDeviceIDs(platforms[PLATFORM], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL);
>
>    cl_context_properties context_properties[] = {
>        CL_CONTEXT_PLATFORM, (cl_context_properties) platforms[PLATFORM], 0
>    };
>
>    cl_context context = clCreateContext(context_properties, 1, &devices[DEVICE], NULL, NULL, &err);
>    assert(err == CL_SUCCESS);
>
>#pragma GCC diagnostic push
>#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
>    cl_command_queue queue = clCreateCommandQueue(context, devices[DEVICE], 0, &err);
>#pragma GCC diagnostic pop
>    assert(err == CL_SUCCESS);
>
>    cl_program program;
>    {
>        // Crashes if directly using src[]
>        char * source = malloc(src_len);
>        memcpy(source, src, src_len);
>        program = clCreateProgramWithSource(context, 1, (const char **) &source, &src_len, &err);
>        assert(err == CL_SUCCESS);
>        free(source);
>    }
>
>    err = clBuildProgram(program, 1, &devices[DEVICE], "", NULL, NULL);
>    assert(err == CL_SUCCESS);
>
>    cl_mem nodes_buf = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(struct Node), NULL, &err);
>    assert(err == CL_SUCCESS);
>
>    {
>        cl_kernel bvh_kernel = clCreateKernel(program, kernels[0], &err);
>        assert(err == CL_SUCCESS);
>
>        err = clSetKernelArg(bvh_kernel, 0, sizeof(cl_mem), &nodes_buf);
>        assert(err == CL_SUCCESS);
>
>        cl_event e_bvh;
>        err = clEnqueueNDRangeKernel(queue, bvh_kernel, 1, NULL, &n, NULL, 0, NULL, &e_bvh);
>        assert(err == CL_SUCCESS);
>        err = clWaitForEvents(1, &e_bvh);
>        assert(err == CL_SUCCESS);
>
>        clReleaseEvent(e_bvh);
>        clReleaseKernel(bvh_kernel);
>    }
>
>    {
>        struct Node * data = clEnqueueMapBuffer(queue, nodes_buf, CL_TRUE, CL_MAP_READ, 0, n * sizeof(struct Node), 0, NULL, NULL, &err);
>        assert(err == CL_SUCCESS);
>
>        printNode(data[0]);
>
>        err = clEnqueueUnmapMemObject(queue, nodes_buf, data, 0, NULL, NULL);
>        assert(err == CL_SUCCESS);
>    }
>
>    clReleaseMemObject(nodes_buf);
>    clReleaseProgram(program);
>    clReleaseCommandQueue(queue);
>    clReleaseContext(context);
>
>    free(devices);
>    free(platforms);
>
>    return 0;
>}

Actions: View

Attachments on bug 98717: 127966

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.