Repeated calls to OpenACC-using function leak VRAM

Accelerated Computing HPC Compilers nvc, nvc++ and nvfortran
1

The below test executable demonstrates the problem. With the ever-growing range-size (N), the VRAM is never released and keeps accumulating… Using nvc++ version 2025.11

Please, advise.

/*
 * Minimal reproducer: scoped #pragma acc data copyin leaks VRAM
 * when allocation sizes vary across calls.
 *
 * Build:
 *   nvc++ -acc -gpu=cc100 -O2 -o vram_leak vram_leak_test4_min.cpp -ldl
 *
 * Expected: delta stays ~0 MB (memory freed at end of scoped block).
 * Actual:   delta grows monotonically, never stabilizes.
 */

#include <cstdio>
#include <cstdlib>
#include <dlfcn.h>
#include <openacc.h>

static size_t
gpu_free_mb(void)
{
	size_t free_bytes, total_bytes;
	typedef int (*cuMemGetInfo_t)(size_t *, size_t *);
	static cuMemGetInfo_t fn;

	if (!fn) {
		void *h = dlopen("libcuda.so.1", RTLD_NOW | RTLD_GLOBAL);
		if (h)
			fn = (cuMemGetInfo_t)dlsym(h, "cuMemGetInfo_v2");
	}
	if (fn && fn(&free_bytes, &total_bytes) == 0)
		return (free_bytes / (1024 * 1024));
	return (0);
}

static void
run(size_t N)
{
	double *buf = new double[N];

	for (size_t i = 0; i < N; ++i)
		buf[i] = 1.0;

	#pragma acc data copyin(buf[0:N])
	{
		#pragma acc parallel loop
		for (size_t i = 0; i < N; ++i)
			buf[i] *= 2.0;
	}
	/* buf[0:N] should be freed here -- but it is not. */

	delete[] buf;
}

int
main(void)
{
	/* Prime-number element counts so every allocation is unique. */
	static const size_t primes[] = {
		 500009,  520013,  540007,  560017,  580027,
		 600011,  620003,  640007,  660013,  680003,
		 700001,  720007,  740011,  760007,  780029,
		 800011,  820007,  840023,  860009,  880001,
		 900001,  920011,  940003,  960017,  980027,
	};
	const int niters = sizeof(primes) / sizeof(primes[0]);

	run(1000);  /* warm up the OpenACC runtime */

	size_t base = gpu_free_mb();
	size_t cum = 0;

	printf("baseline: %zu MB free\n\n", base);

	for (int i = 0; i < niters; i++) {
		size_t N = primes[i];

		cum += N * sizeof(double);
		run(N);

		size_t now = gpu_free_mb();
		printf("iter %2d  N=%7zu (%2zu MB)  free=%5zu MB"
		    "  delta=%4zd MB  cumalloc=%3zu MB\n",
		    i, N, N * 8 / (1024 * 1024), now,
		    (long)base - (long)now, cum / (1024 * 1024));
	}

	printf("\ntotal leaked: %zd MB  (cumulative allocs: %zu MB)\n",
	    (long)base - (long)gpu_free_mb(), cum / (1024 * 1024));
	return (0);
}
2

When running with the NV_ACC_MEM_MANAGE=0 the reported leak goes to zero, which is a relief.

I feel like this should be the default – and a lot more prominently documented…

3

It is done for performance reasons. Allocating memory on GPU is slow, it is better to keep a pool of memory around for reuse.

4

Multiple malloc-library implementations have been written over the decades for the same reason – the problem is not GPU-specific.

Still, a developer making use of acc-pragmas has a right to expect, all of the memory used by a kernel is freed after the kernel ends.

The usage of pools may make sense for one process accesses the GPU. But multiple workers accessing it in parallel – each with its own collection of allocated and never really freed blocks – blows up VRAM for no fault of the programmer…

5

Compared to CPU malloc, cudaMalloc is commonly 100x to 10,000x slower for small allocation.
So, for your very specific use case, NV_ACC_MEM_MANAGE=0 is the way to go. For general use, a memory pool is the right choice since it is performant.