title: CUDA device function pointers in structure without static pointers or symbol copies
body: <p>My intended program flow would look like the following if it were possible:</p>
  
  <pre><code>typedef struct structure_t
  {
    [...]
    /* device function pointer. */
    __device__ float (*function_pointer)(float, float, float[]);
    [...]
  } structure;
  
  [...]
  
  /* function to be assigned. */
  __device__ float
  my_function (float a, float b, float c[])
  {
    /* do some stuff on the device. */
    [...]
  }
  
  void
  some_structure_initialization_function (structure *st)
  {
    /* assign. */
    st-&gt;function_pointer = my_function;
    [...]
  }
  </code></pre>
  
  <p>This is not possible, and ends in a familiar error during compilation regarding the placement of __device__ in the structure.</p>
  
  <pre><code> error: attribute "device" does not apply here
  </code></pre>
  
  <p>There are some examples of similar types of problems here on stackoverflow, but they all involve the use of static pointers outside the structure. Examples are <a href="https://stackoverflow.com/questions/11102945/device-function-pointers-as-struct-members">device function pointers as struct members</a> and <a href="https://stackoverflow.com/questions/9000388/device-function-pointers">device function pointers</a>. I've taken a similar approach with success previously in other codes where it's easy for me to use static device pointers and define them outside of any structures. Currently though this is a problem. It's written as an API of sorts and the user may define one or two or dozens of structures which need to include a device function pointer. So, defining static device pointers outside of the structure is a major problem.</p>
  
  <p>I'm fairly certain the answer exists within the posts I have linked above, through use symbol copies, but I've not been able to put them to successful use.</p>