title: Creating an object in device code
body: <p>I want to create an object on the device and allocate it to a pointer available on the host. Is there something I'm doing wrong in here?</p>
  
  <pre><code>__global__ void createAProduction(DeviceProduction* production) {
      production = new AProduction();
  }
  
  DeviceProduction * devAProduction = NULL;
  cudaMalloc(&amp;devAProduction, sizeof(AProduction));
  createAProduction&lt;&lt;&lt;1, 1&gt;&gt;&gt;(devAProduction);
  deviceProductions["A"] = devAProduction;
  </code></pre>
  
  <p>Somewhere further in the code I'd like to do sth. like:</p>
  
  <pre><code>BatchOperation ** devBatchOperations;
      cudaMalloc((void **) &amp;devBatchOperations, sizeof(BatchOperation *) * operationCount);
  </code></pre>
  
  <p>Then I populate that pointer array with that:</p>
  
  <pre><code>void DeviceBatchExecutor::execute(vector&lt;BatchOperation&gt; operationsToPerform) {
      BatchOperation ** devBatchOperations;
      cudaMalloc((void **) &amp;devBatchOperations, sizeof(BatchOperation *) * operationsToPerform.size());
      int i = 0;
      for(batchOperationIt it = operationsToPerform.begin(); it != operationsToPerform.end(); ++it) {
          BatchOperation * devBatchOperation;
          cudaMalloc(&amp;devBatchOperation, sizeof(BatchOperation));
          cudaMemcpy(&amp;devBatchOperation, &amp;it, sizeof(BatchOperation), cudaMemcpyHostToDevice);
          Vertex * devInputNode = it-&gt;inputNode-&gt;allocateToDevice();
          cudaMemcpy(&amp;(devBatchOperation-&gt;inputNode), &amp;devInputNode, sizeof(Vertex *), cudaMemcpyDeviceToDevice);
          cudaMemcpy(&amp;(devBatchOperation-&gt;production), &amp;(it-&gt;production), sizeof(Production *), cudaMemcpyDeviceToDevice);
          cudaMemcpy(&amp;devBatchOperations[i], &amp;devBatchOperation, sizeof(BatchOperation *), cudaMemcpyDeviceToDevice);
          i++;
      }
      int operationCount = operationsToPerform.size();
      executeOperations&lt;&lt;&lt;operationCount, 1&gt;&gt;&gt;(devBatchOperations);
  }
  </code></pre>
  
  <p>where production is a pointer to the device memory holding that created object AProduction. Then I finally invoke processing via</p>
  
  <pre><code>executeOperations&lt;&lt;&lt;operationCount, 1&gt;&gt;&gt;(devBatchOperations);
  </code></pre>
  
  <p>So I'm relying on virtual method calls. As those DeviceProduction objects were created on the device, there is also a virtual pointer table so it should work. See example <a href="http://w3foverflow.com/question/any-known-work-arounds-to-cudas-lack-of-virtual-functions/" rel="nofollow">here</a>. But it doesn't since the received batch operations seem random... crashes on invocation.</p>
  
  <pre><code>__global__ void executeOperations(BatchOperation ** operation) {    
      operation[blockIdx.x]-&gt;production-&gt;apply(operation[blockIdx.x]-&gt;inputNode);
  }
  </code></pre>
  
  <p>Batch operation is a struct holding the production to be executed.</p>
  
  <pre><code>struct BatchOperation {
      Production * production;
      Vertex * inputNode;
      Vertex * outputNode;
  };
  </code></pre>