title: Calling operator new in a CUDA kernel
body: <p>In the following, I'm trying to allocate device memory in kernel <code>create_A()</code> for a trivial object <code>A</code>, then delete it from another kernel <code>delete_A()</code>. It seems the allocation of <code>ptr_</code> in <code>create_A</code> does not persist in device memory after <code>create_A</code> returns.</p>
  <h1>main1.cu</h1>
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;cuda.h&gt;
  
  class A {
  public:
      int a;
  
      __device__ A(int x)
          : a(x) { 
          printf(&quot;\nA()&quot;); 
      }
      
      __device__ ~A() { printf(&quot;\n~A()&quot;); }
  };
  
  __global__ void create_A(A* ptr_) {
  
      printf(&quot;\nCreating A...&quot;);
  
      ptr_ = new A(7);
  
      return;
  }
  
  __global__ void delete_A(A* ptr_) {
  
      printf(&quot;\nA.a = %d \nDeleting A...&quot;, ptr_-&gt;a);
  
      delete ptr_;
       
      return;
  }
  
  int main(void) {
      A* ptr_ = NULL;
      
      create_A&lt;&lt;&lt;1,1&gt;&gt;&gt;(ptr_);
  
      delete_A&lt;&lt;&lt;1,1&gt;&gt;&gt;(ptr_);
  
      cudaDeviceSynchronize();
  
      cudaError_t cudaerr = cudaDeviceSynchronize();
      if (cudaerr != CUDA_SUCCESS)
          printf(&quot;kernel launch failed with error \&quot;%s\&quot;.\n&quot;,
                 cudaGetErrorString(cudaerr));
  
      return 0;
  }
  </code></pre>
  <p>Output:</p>
  <pre><code>Creating A...
  A()
  kernel launch failed with error &quot;an illegal memory access was encountered&quot;.
  </code></pre>
  <p>On the other hand, if I make <code>ptr_</code> a global variable as in the following, it works.</p>
  <h1>main2.cu</h1>
  <pre><code>class A {
  public:
      int a;
  
      __device__ A(int x) 
          : a(x) { 
          
          printf(&quot;\nA()&quot;); 
      }
      
      __device__ ~A() { printf(&quot;\n~A()&quot;); }
  };
  
  __device__ A* ptr_;
  
  __global__ void create_A() {
  
      printf(&quot;\nCreating A...&quot;);
  
      ptr_ = new A(7);
  
      return;
  }
  
  __global__ void delete_A() {
          
      printf(&quot;A.a is: %d&quot;, ptr_-&gt;a);
  
      printf(&quot;\nDeleting A...&quot;);
  
      delete ptr_;
       
      return;
  }
  
  int main(void) {
      
      create_A&lt;&lt;&lt;1,1&gt;&gt;&gt;();
  
      delete_A&lt;&lt;&lt;1,1&gt;&gt;&gt;();
  
      cudaDeviceSynchronize();
  
      cudaError_t cudaerr = cudaDeviceSynchronize();
      if (cudaerr != CUDA_SUCCESS)
          printf(&quot;kernel launch failed with error \&quot;%s\&quot;.\n&quot;,
                 cudaGetErrorString(cudaerr));
  
      return 0;
  }
  </code></pre>
  <p>Output:</p>
  <pre><code>Creating A...
  A()
  A.a is: 7
  Deleting A...
  ~A()
  </code></pre>
  <p>Why does the allocation of <code>ptr_</code> persist on the device only in the second case?</p>