Results 1 to 5 of 5

Thread: address space qualifiers and pointers to arrays

  1. #1
    Junior Member
    Join Date
    Sep 2012
    Posts
    5

    address space qualifiers and pointers to arrays

    The following code works on Nvidia, but the AMD and Intel OpenCL compilers complain about the absence of a valid address space qualifier:

    Code :
    __kernel void foo(__global float (*pointer_to_matrix)[3][4])
    {
    }

    I think this code is legal. Any opinion?

  2. #2
    Senior Member
    Join Date
    Oct 2012
    Posts
    103

    Re: address space qualifiers and pointers to arrays

    According to OpenCL specification: "arguments to kernel functions in a program cannot be declared as a pointer to a pointer"

  3. #3
    Junior Member
    Join Date
    Sep 2012
    Posts
    5

    Re: address space qualifiers and pointers to arrays

    But this is not a pointer to a pointer, this is a pointer to a 3x4 matrix (i.e., a pointer to a memory area holding 12 consecutive floats).

  4. #4
    Senior Member
    Join Date
    Oct 2012
    Posts
    103

    Re: address space qualifiers and pointers to arrays

    When you declare "float matrix[3][4]", matrix is a pointer to an array of 12 floats.
    So "float (*pointer_to_matrix)[3][4]" is obviously a pointer to a pointer to floats.

    If you intend to use a pointer to an array of 12 floats, you should declare your kernel as

    __kernel void foo(__global float matrix[3][4])

  5. #5
    Junior Member
    Join Date
    Sep 2012
    Posts
    5

    Re: address space qualifiers and pointers to arrays

    Both declarations declare a pointer to an array of 12 floats, even though the syntax for dereferencing the pointer is different.

    Consider the five examples below; these are fully equivalent and could/should lead to the same generated binary code (remove the __kernel and __global, and it is legal C code; you will see that an optimizing C compiler will generate exactly the same instructions for all five cases).

    Code :
    __kernel void foo1(__global float *ptr) // ok on NVIDIA, AMD, and Intel
    {
      ptr[6] = 42;
    }
     
    __kernel void foo2(__global float ptr[12]) // gives error on Intel
    {
      ptr[6] = 42;
    }
     
    __kernel void foo3(__global float (*ptr)[12]) // gives error on AMD and Intel
    {
      (*ptr)[6] = 42;
    }
     
    __kernel void foo4(__global float ptr[3][4]) // gives error on AMD and Intel
    {
      ptr[1][2] = 42;
    }
     
    __kernel void foo5(__global float (*ptr)[3][4]) // gives error on AMD and Intel
    {
      (*ptr)[1][2] = 42;
    }

    As far as I can see, none of these constructs is forbidden by the OpenCL specification, so I think that all these five functions should be accepted by the compiler (only NVIDIA does so).

Similar Threads

  1. Replies: 7
    Last Post: 03-09-2010, 11:34 AM
  2. address space on multiple chips
    By wilbur in forum OpenCL
    Replies: 2
    Last Post: 07-03-2009, 12:09 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
  •