Results 1 to 4 of 4

Thread: How to read elements from a float16 (or other) using index?

  1. #1
    Junior Member
    Join Date
    Dec 2009
    Posts
    16

    How to read elements from a float16 (or other) using index?

    I'm wanting to read an arbitrary element from a float16. The kernel code below using array subscript syntax "weights[i]" works on Apple's OpenCL implementation, however it errors on Nvidia's Linux implementation saying "subscripted value is not an array, pointer, or vector"
    Not sure if this is valid OpenCL syntax, or if Apple just happens to support it (?)
    Code :
    __kernel void convolve_x16(__global const float *source,
                   __global float *dest,
                   const int2 image_dims,
                   const float16 weights) {
      int x = get_global_id(0);
      int y = get_global_id(1);
      if ((x < image_dims.x) && (y < image_dims.y)) {
        int index = y * image_dims.x + x;
        float d = .0f;
        float wsum = .0f;
        float w;
        int imax = min(16, image_dims.x - x);
        for (int i = 0; i < imax; i++) {
          w = weights[i];  // <---- this is the syntax in question
          d += w * source[index + i];
          wsum = w;
        }
        dest[index] = d / w;
      }
    }

    My question: is there another way to index into a float16 that is correct OpenCL syntax? (I don't want to use weights.s0 etc and unroll the loop.)

    If not, is there another way to pass an arbitrary sized chunk of constant memory as a kernel argument (such as an array 16 floats) that can be indexed like this? I know I could create a read-only global memory buffer, fill it using clEnqueueWriteImage, and then pass that in as a float* kernel argument, but I'm wondering if I'm missing a simpler way (more like the float16) ?

    Thanks!

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

    Re: How to read elements from a float16 (or other) using ind

    My question: is there another way to index into a float16 that is correct OpenCL syntax? (I don't want to use weights.s0 etc and unroll the loop.)
    I thought that accessing vector elements using array indexing notation as your example was legal since CL 1.1. but my memory must be failing because I don't find that in the spec. Can a language lawyer clarify this?

    There are ways to work around that, though. You can use either a cast or a union. I would use a cast for convenience:

    Code :
    float *sweights = (float*)&weights;
    sweights[i];
    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

    Re: How to read elements from a float16 (or other) using ind

    When writing ordinary host code, at least with gcc, then accessing data of some type through a pointer of another type gives warnings about type-punned pointers and strict aliasing rules. Is the same true for OpenCL kernel code? Or in other worlds, could
    Code :
    float *sweights = (float*)&weights;
    // Do a bunch of stuff to float16 weights.
    // Dereference float* sweights.
    produce unexpected results if weights wasn't const and could be written to?

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

    Re: How to read elements from a float16 (or other) using ind

    Sorry, I didn't notice that it was const. Also, you are correct about strict aliasing. My bad!

    There's yet another workaround: use a union like this:

    Code :
    union
    {
        float     s[16];
        float16 v;
    } uweights;
     
    ...
     
    uweights.v = weights;
     
    // Here you can read from uweights.s[i] safely thanks to section 6.2.4.1
    // Notice that this is a nice property of OpenCL C that does not exist in C99

    That said, I will ask around because it feels a bit silly that we don't allow reading from vector elements using array subscripting syntax.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

Similar Threads

  1. [XF] Float16 vs 16 float
    By xfaure in forum OpenCL
    Replies: 5
    Last Post: 09-07-2011, 05:47 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
  •