Results 1 to 6 of 6

Thread: casting char* to struct

  1. #1

    casting char* to struct

    I'm passing in an uchar array to my kernel then trying to cast to a struct. The GPU driver crashes after a few frames. Here's the OpenCL code. I've stripped out everything else to simplify it and it still crashes. What am I missing?

    typedef struct
    {
    uchar r, g, b;
    } Pixel;

    __kernel void main (__global uchar* framebuffer, int width, int height)
    {
    int index = get_global_id(1) * width + get_global_id(0);
    Pixel* p = (Pixel* )&framebuffer[index * 3];
    p->r = 255;
    p->g = 0;
    p->b = 0;
    }

  2. #2
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: casting char* to struct

    Assuming that you have already verified that width and height match the global work-size that you passed to clEnqueueNDRangeKernel() and assuming that you have also verified that the buffer object you passed to the kernel is large enough, my guess would be a problem with alignment.

    There's a good chance that sizeof(Pixel) is 4 bytes if not larger. The spec is a bit vague here, although it makes the following two interesting statements.

    Section 6.1.5:
    A data item declared to be a data type in memory is always aligned to the size of the data type in bytes.
    Section 6.2.5.:
    Pointers to old and new types may be cast back and forth to each other. Casting a pointer to a new type represents an unchecked assertion that the address is correctly aligned.
    In other words, by casting the uchar pointer to Pixel your code is asserting that the pointer is aligned to sizeof(Pixel). If it's not the case, the code may and will fail to run properly.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  3. #3
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: casting char* to struct

    Thinking about it a bit more, I'm not convinced that the alignment of the struct would be an issue from the viewpoint of standard OpenCL. As long as the members of the struct are properly aligned when you dereference the struct pointer, it should be fine.

    The bug reproducer you have is nice and short. Have you tried sending it to your hardware vendor?
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  4. #4
    Member
    Join Date
    Jul 2011
    Location
    Moscow, Russia
    Posts
    41

    Re: casting char* to struct

    The OpenCL extension cl_khr_byte_addressable_store removes certain restrictions on built-in types char, uchar, char2, uchar2, short, and half. An application that wants to be able to write to elements of a pointer (or struct) that are of type char, uchar, char2, uchar2, short, ushort, and half will need to include the #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable directive before any code that performs writes that may not be supported.

    http://www.khronos.org/registry/cl/s...ble_store.html
    Blog (in russian)

  5. #5

    Re: casting char* to struct

    I've already included the cl_khr_byte_addressable_store extension. It still crashes. I've also tried adding the alignment and byte packing attributes to the Pixel struct but to no avail.

    typedef struct
    {
    uchar r __attribute__ ((packed));
    uchar g __attribute__ ((packed));
    uchar b __attribute__ ((packed));
    } Pixel __attribute__ ((aligned (1)));

    On my NVIDIA Geforce 480 it freezes and crashes. On my AMD Radeon 5850 it actually throws a compiler error stating "invalid type conversion: Pixel* p = (Pixel* )&colorBuffer[index*3];"

  6. #6

    Re: casting char* to struct

    **EDIT**

    I mistakenly left out the __global modifier when declaring the Pixel variable. It all works fine now.

    __global Pixel* p = (__global Pixel* )&colorBuffer[index * 3];

Similar Threads

  1. Type casting question
    By dlw in forum OpenCL
    Replies: 1
    Last Post: 07-06-2011, 03:35 PM
  2. Casting int2, float2
    By toneburst in forum OpenCL
    Replies: 1
    Last Post: 12-04-2010, 06:49 PM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •