title: memory allocation inside a CUDA kernel
body: <p>I have the following (snippet) of a kernel.</p>
  
  <pre><code>__global__ void plain(int* geneVec, float* probs, int* nComponents, float* randomNumbers,int *nGenes)
  {
  
      int xid = threadIdx.x + (blockDim.x * blockIdx.x);
  
      float* currentProbs= (float*)malloc(sizeof(float)*tmp);
  
           .....
           .....
  
      currentProbs[0] = probs[start];
      for (k=1;k&lt;nComponents[0]; k++)
      {
          currentProbs[k] = currentProbs[k-1] + prob;
      }
  
         ...
         ...
        free(currentProbs);
  
  }
  </code></pre>
  
  <p>When it's static (even the same sizes) it's very fast, but when CurrentProbs is dynamically allocated (as above) performance is awful.</p>
  
  <p>This question said I could do this inside a kernel: <a href="https://stackoverflow.com/questions/4715238/cuda-allocate-memory-in-device-function">CUDA allocate memory in __device__ function</a></p>
  
  <p>Here is a related question: <a href="https://stackoverflow.com/questions/7476560/efficiency-of-malloc-function-in-cuda">Efficiency of Malloc function in CUDA</a></p>
  
  <p>I was wondering if any other methods have solved this other than the one proposed in the paper?
  It seems ridiculous that one cannot malloc/free inside a kernel without this sort of penalty.</p>