title: Allocating and freeing memory in a single kernel slower than Allocating in one then freeing in another
body: <p><strong>Background</strong></p>
  
  <p>I had CUDA application in which I wanted to reproduce how I had improved performance.  I ended up writing the entire code over but could not reproduce the performance gains.  Then I noticed that I had a memory leak in the original code.  So I added the required delete, then my performance plummeted.</p>
  
  <p>Now I understand memory allocation and deletion should have some overhead.  But I then did a test and I found out that if I allocate in one kernel and then perform the memory de-allocation in another I don't get the performance penalty.</p>
  
  <p>I have some example code and output of the performance.  From the example you get on average a 2.64 speedup (ignoring the really strange first one).  </p>
  
  <p>In my real world example though I see a speedup of around 10.  For instance I have one that takes 45 minutes when I use two kernels, and 8 hours with one kernel.</p>
  
  <p><strong>Question</strong></p>
  
  <p>Does anyone have an idea what is going on?</p>
  
  <p>The only thing I can think of are these two cases:</p>
  
  <ol>
  <li>There is a bug</li>
  <li>delete does take a really long time
  <ul>
  <li>In the first kernel it waits for it to complete</li>
  <li>In the second kernel it calls delete and returns right away.</li>
  </ul></li>
  </ol>
  
  <p><strong>Setup</strong></p>
  
  <p><em>OS :</em></p>
  
  <p>Windows 7 64-bit</p>
  
  <p><em>nvcc --version :</em> </p>
  
  <p>Built on Fri_Mar_14_19:30:01_PDT_2014
  Cuda compilation tools, release 6.0, V6.0.1</p>
  
  <p><em>Compiler Options :</em></p>
  
  <p>nvcc -O3 -arch=sm_35 forty_hours_debugging.cu</p>
  
  <p><em>---EDIT---</em></p>
  
  <p><em>GPU :</em></p>
  
  <p>GTX 780</p>
  
  <p><strong>Code</strong></p>
  
  <pre><code>#include &lt;iostream&gt;
  #include &lt;string&gt;
  
  __device__ int *D_DATA[1000];
  
  __global__ void MyAllocate(int size,bool perform_delete) {
      int *data;
      __shared__ int *t_data;
  
      if(threadIdx.x==0) {
          t_data = new int[size];
      }__syncthreads();
  
      if(t_data==NULL) return;
  
      D_DATA[blockIdx.x] = data= t_data;  
  
      memset(data,0,size);
      __syncthreads();
  
      for(int i = threadIdx.x;i &lt; size; i+= blockDim.x) {
              data[i] = i * i + perform_delete * i;
      }__syncthreads();
  
      // If we should delete then do so, otherwise another kernel (hopefully) will
      if(threadIdx.x==0 &amp;&amp; perform_delete) {
          delete data;
      }
  }
  
  __global__ void MyDelete() {
      delete D_DATA[blockIdx.x];
  }
  
  int main(int argc,char **argv) {    
  
      cudaEvent_t start,stop;
      float time;
      const std::string pre[2] = {"One Kernel ","Two Kernels "};  
  
      for(int size = 1000000; size &lt; 1024000000; size *= 2) {
          for(int i = 0; i &lt; 2; ++i) {
  
              cudaEventCreate(&amp;start);
              cudaEventCreate(&amp;stop);
  
              cudaEventRecord(start,0);
              MyAllocate&lt;&lt;&lt;1000,128&gt;&gt;&gt;(size,i==0);
              if(i!=0) {
                  MyDelete&lt;&lt;&lt;1000,1&gt;&gt;&gt;();
              }
              cudaEventRecord(stop,0);
              cudaEventSynchronize(stop);
              cudaEventElapsedTime(&amp;time, start, stop);
              std::cout &lt;&lt; pre[i] &lt;&lt; "Time : " &lt;&lt; (time) &lt;&lt; "ms" &lt;&lt; std::endl;
  
              cudaEventDestroy(start);
              cudaEventDestroy(stop);
          }
          cudaDeviceReset();
          std::cout &lt;&lt; std::endl;
      }
      return 0;
  }
  </code></pre>
  
  <p><strong>Results</strong></p>
  
  <pre><code>One Kernel Time : 88.3073ms
  Two Kernels Time : 88.3073ms
  
  One Kernel Time : 0.109024ms
  Two Kernels Time : 0.042912ms
  
  One Kernel Time : 0.11184ms
  Two Kernels Time : 0.042272ms
  
  One Kernel Time : 0.1072ms
  Two Kernels Time : 0.042688ms
  
  One Kernel Time : 0.098464ms
  Two Kernels Time : 0.042208ms
  
  One Kernel Time : 0.103776ms
  Two Kernels Time : 0.0432ms
  
  One Kernel Time : 0.111776ms
  Two Kernels Time : 0.04256ms
  
  One Kernel Time : 0.114592ms
  Two Kernels Time : 0.0424ms
  
  One Kernel Time : 0.109888ms
  Two Kernels Time : 0.042656ms
  
  One Kernel Time : 0.119456ms
  Two Kernels Time : 0.042336ms
  </code></pre>