title: OpenCL - copying a struct from global to local memory
body: <p>I have a struct Primitive, which has the following definition:</p>
  
  <pre><code>typedef struct Primitive {
      float m[12];
      float invm[12];
      enum PrimitiveType type;
      int rayDensity;
      float util1;             
      float util2;                
  } Primitive;
  </code></pre>
  
  <p>I pass an array of these structs to my kernel in a constant memory buffer:</p>
  
  <pre><code>__constant Primitive *objects;
  </code></pre>
  
  <p>As part of an optimization exercise I want to look at loading the structs into local memory, so my kernel has code to the likes of this:</p>
  
  <pre><code>__kernel void test(int n_objects, __constant Primitives *objects) {
      local Primitive pFrom, pTo;
  
      for(int i = 0; i &lt; n_objects; i++) {
          pFrom = objects[i]; 
      }
  
  }
  </code></pre>
  
  <p>When I run this I get a compilation error saying :</p>
  
  <pre><code>ptxas application ptx input, line 42; error: State space mismatch between instruction and address in instruction 'ld'
  </code></pre>
  
  <p>As an experiment I have tried first copying the struct to a private variable and then to the local variable as follows :</p>
  
  <pre><code>__kernel void test(int n_objects, __constant Primitives *objects) {
      Primitive pF, Pt;
      local Primitive pFrom, pTo;
  
      for(int i = 0; i &lt; n_objects; i++) {
          pF = objects[i]
          pFrom = pF; 
      }
  
  }
  </code></pre>
  
  <p>Which now compiles and runs however it seems like the object is not deeply copied into the local variable pFrom.</p>
  
  <p>Please note that my code samples are purely samples and I have removed everything for the sake of brevity. Also my code works fine when I use the primitive structs directly from the constant global memory.</p>
  
  <p>Does anyone know what I am missing here, surely its some basic fundamental to deep copying or OpenCL address spaces.</p>