[BUG]: single_device_tls_caching_allocator doesn't free memory on failed allocation

[gevtushenko]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Runtime Error

Component

Thrust

Describe the bug

Users have to workaround an issue in thrust::detail::single_device_tls_caching_allocator by reaching into detail namespace to release unused memory upon failed allocation. Without a try-catch block, the following code fails:

#include <thrust/detail/caching_allocator.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/sort.h>

#include <iostream>

int main(int argc, char const *argv[])
{
  size_t mode = 0;
  if (argc > 1)
  {
    mode = std::stoi(argv[1]);
  }

  size_t nbSteps   = 100000;
  size_t deltaStep = 5;

  thrust::device_vector<size_t> target(nbSteps * deltaStep, 1);

  for (size_t i = 0; i < nbSteps; ++i)
  {
    std::cerr << "Steps " << i << "\r";

    if (mode == 0)
    {
      thrust::sort(thrust::cuda::par(thrust::detail::single_device_tls_caching_allocator()),
                   target.begin(),
                   target.end());
    }
    else
    {
      try
      {
        thrust::sort(thrust::cuda::par(thrust::detail::single_device_tls_caching_allocator()),
                     target.begin(),
                     target.begin() + (deltaStep * (i + 1)));
      }
      catch (...)
      {
        thrust::detail::single_device_tls_caching_allocator().resource()->release();
        thrust::sort(thrust::cuda::par(thrust::detail::single_device_tls_caching_allocator()),
                     target.begin(),
                     target.begin() + (deltaStep * (i + 1)));
      }
    }
  }
}

How to Reproduce

Remove try catch in the code above.

Expected behavior

Unsuccessful do_allocate in disjoint_pool frees unused memory and attemps allocation again.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[gevtushenko]

CC @dkolsen-pgi