When calling cudaMallocAsync, the memory pointer is returned immediately, so why is it still named an asynchronous function?
From my understanding, asynchronous operations typically return control to the caller immediately while the operation completes in the background, often with a future or callback mechanism. However, cudaMallocAsync returns the pointer right away, which makes me wonder what aspect of it is actually asynchronous.
Is the asynchronicity related to the actual allocation being deferred or occurring on a separate stream? Or does it refer to something else entirely, such as memory reuse or synchronization behavior with streams?
cudaMallocAsync provides stream-ordered allocation. The pointer is returned immediately. However, unless you explicitly synchronize the allocating stream afterwards, the allocation is only valid to use in the allocating stream.
cudaMalloc(&ptr);
//can use ptr immediately in all streams
kernel1<<<, stream1>>>(ptr); //ok
kernel2<<<, stream2>>>(ptr); //ok
cudaMallocAsync(&ptr, stream);
cudaStreamSynchronize(stream);
//can use ptr immediately in all streams
kernel1<<<, stream1>>>(ptr); //ok
kernel2<<<, stream2>>>(ptr); //ok
cudaMallocAsync(&ptr, stream);
kernel<<<stream>>>(ptr); // ok
kernel1<<<, stream1>>>(ptr); //not ok
kernel2<<<, stream2>>>(ptr); //not ok
Memory allocated via cudaMallocAsync comes from a memory-pool that enables fast re-use of free allocations from the same stream or from a different stream.
There is a whole section in the CUDA Programming Guide dedicated to stream-ordered allocations which might be interesting to you: 4.3. Stream-Ordered Memory Allocator — CUDA Programming Guide
Thank you for your detailed explanation — it’s very helpful.
I think I now understand the stream-ordered semantics better. However, I still have one follow-up question regarding the underlying mechanism:
If cudaMallocAsync returns a valid virtual address (VA) immediately, and the memory comes from a pre-allocated pool that already has physical backing ¶, then why can’t other streams safely use that pointer until the allocating stream is synchronized?
In other words, is the “not ready” state due to:
-
Some internal bookkeeping or ownership metadata that hasn’t been updated yet?
-
The pool reusing memory that is still in flight (e.g., being used by a previous kernel in the allocating stream)?
-
Or is it simply a design choice to enforce stream-ordered semantics and prevent races, even though the physical memory itself is technically already allocated?
I previously thought that the memory pool was just a user-space cache to reduce kernel call overhead, and that once the pointer is returned, the memory is physically ready for use by anyone. It seems my understanding may be incomplete here.
Could you please clarify what exactly makes the memory “not usable” by other streams before synchronization?
I believe the main reason is 2. Deallocation is also stream-ordered. Consider the following workflow
cudaMallocAsync(&ptr, 1024, stream);
kernel<<<stream>>>(ptr);
cudaFreeAsync(ptr, stream);
cudaMallocAsync(&ptr2, 1024, stream);
//ptr2 == ptr is possible.
As you have noticed, the allocated ptr2 is returned immediately. However due to stream-ordering, the backing memory could still be in use by the kernel. Using it in a different stream without synchronization would constitute a race-condition.
Thank you, I’ll write some demos to verify this.