cudaMemGetInfo() doesn't work correctly when explicitly setting heap size for device dynamic mem

Hello there is a problem with cudaMemGetInfo reporting the wrong amount of free memory.

The following code should reproduce what I’m talking about.

__global__ void test()

{

	char* ptr = (char*)malloc(1*1024*1024);	//allocate 1 MB

}

int main()

{

	size_t free_mem, total_mem;

	cudaMemGetInfo(&free_mem, &total_mem);

	printf("Free/Total %lu/%lu %u %%\n",free_mem, total_mem, (unsigned int) ((free_mem*100.0)/total_mem));

	// Set heap limit to 80% of free memory

        // COMMENTING THIS OUT WORKS

	cudaError error = cudaThreadSetLimit(cudaLimitMallocHeapSize, (size_t) (0.8 * free_mem));

		

	test<<<1,1>>>();

	cudaDeviceSynchronize();

	cudaMemGetInfo(&free_mem, &total_mem);

	printf("Free/Total %lu/%lu %u %%\n",free_mem, total_mem, (unsigned int) ((free_mem*100.0)/total_mem));

	return 0;

}
Output:

Free/Total 1080008704/1275658240 84 %

Free/Total 107847680/1275658240 8 %

My kernels dynamically allocate and free memory as they need it so I set the device heap limit before any kernels are called to a certain percentage of available memory.

The problem is that after setting the heap limit, cuda runtime excludes this amount of memory from being reported as free. Device side allocation of more and more memory of course doesn’t fail because the memory set aside for the heap is also free memory.

The following is what cudaMemGetInfo should do:

Returns in *free and *total respectively, the free and total amount of memory available for allocation by the device in bytes

but instead it (perhaps) only reports the amount of free memory available to be allocated from the host.

So, am I wrong to expect cudaMemGetInfo to report the total amount of memory available to be allocated regardless of what portion of it is available to be allocated from the host and what portion is available to be allocated from the device.

Or is this intended to work this way and is just a matter of a vague definition.

Thank you.

Anyone?

Is there a way to find out how much memory has been allocated on the heap without doing any book-keeping?

Same issue here (and lots more!!!)…
If you got some answers i will appreciate to know.