title: CUDA allocate memory in __device__ function
body: <p>Is there a way in CUDA to allocate memory dynamically in device-side functions ?
  I could not find any examples of doing this.</p>
  <p>From the CUDA C Programming manual:</p>
  <blockquote>
  <p>B.15 Dynamic Global Memory Allocation</p>
  <pre><code>void* malloc(size_t size); 
  void free(void* ptr); 
  </code></pre>
  <p>allocate and free memory dynamically from a fixed-size heap in global memory.</p>
  <p>The CUDA in-kernel <code>malloc()</code> function allocates at least size bytes from the device heap and returns a pointer to the allocated memory or NULL if insufficient memory exists to fulfill the request. The returned pointer is guaranteed to be aligned to a 16-byte boundary.</p>
  <p>The CUDA in-kernel <code>free()</code> function deallocates the memory pointed to by <code>ptr</code>, which must have been returned by a previous call to <code>malloc()</code>. If <code>ptr</code> is <code>NULL</code>, the call to free() is ignored. Repeated calls to free() with the same ptr has undefined behavior.</p>
  <p>The memory allocated by a given CUDA thread via <code>malloc()</code> remains allocated for the lifetime of the CUDA context, or until it is explicitly released by a call to <code>free()</code>. It can be used by any other CUDA threads even from subsequent kernel launches. Any CUDA thread may free memory allocated by another thread, but care should be taken to ensure that the same pointer is not freed more than once.</p>
  </blockquote>