One of the attributes of CUDA memory pools is CU_MEMPOOL_ATTR_REUSE_ALLOW_OPPORTUNISTIC, described in the doxygen as follows:
Allow reuse of already completed frees when there is no dependency between the free and allocation.
If a free (a cuFreeAsync() I presume) depends on an allocation - how can that free be completed when the allocation needs to happen? Or - am I misunderstanding what this attribute allows?
This flag is explained in the CUDA programming guide.
11.9.2. cudaMemPoolReuseAllowOpportunistic
According to the cudaMemPoolReuseAllowOpportunistic policy, the allocator examines freed allocations to see if the free’s stream order semantic has been met (such as the stream has passed the point of execution indicated by the free). When this is disabled, the allocator will still reuse memory made available when a stream is synchronized with the CPU. Disabling this policy does not stop the cudaMemPoolReuseFollowEventDependencies from applying.
cudaMallocAsync(&ptr, size, originalStream);
kernel<<<..., originalStream>>>(ptr, ...);
cudaFreeAsync(ptr, originalStream);
// after some time, the kernel finishes running
wait(10);
// When cudaMemPoolReuseAllowOpportunistic is enabled this allocation request
// can be fulfilled with the prior allocation based on the progress of originalStream.
cudaMallocAsync(&ptr2, size, otherStream);
