CUDA memory leak persistence

[ashwin]

Consider a CUDA program that has a memory leak – some device memory allocated using [font=“Courier New”]cudaMalloc()[/font] is not being freed using [font=“Courier New”]cudaFree()[/font].

In this case, is the CUDA device memory properly cleaned up after this leaky program has exited? Or does this leak persist even after the demise of the program and cause problems/slowups for later CUDA programs?

[seibert]

The memory is freed on the GPU when the program exits.

[SPWorley]

I think it’s when the CUDA context is destroyed, likely when the thread that created it is exited.
However I admit I have never tested this!

[ashwin]

I think it’s when the CUDA context is destroyed, likely when the thread that created it is exited.

The reason I ask this is because I’m sometimes executing my code many times, single stepping through the host code and such. And suddenly on some invocation, the CUDA program starts to freeze up and then come back to normal after a few seconds. A few invocations after this, the CUDA program totally freezes the computer leaving me no choice but to reboot.

The only way I can explain this slow transition from normality to freezing up is if either me or the graphics driver is leaking memory on the device somehow over many invocations.

[seibert]

I’ve killed many processes doing CUDA work in a row, and not seen a leak. One thing to check is that the processes are actually dead. If you are building up an accumulation of zombie processes, they might continue to hold onto memory on the device even though they are not doing anything.

[T.B]

I’ve killed many processes doing CUDA work in a row, and not seen a leak. One thing to check is that the processes are actually dead. If you are building up an accumulation of zombie processes, they might continue to hold onto memory on the device even though they are not doing anything.

This doesn’t seem to be quite that clear cut. I’ve seen many instances where my allocated memory didn’t get free’d after a crash (or after I forgot to free it myself). A good general practice is to use cuMemGetInfo at the end of program to see if everything shut down nicely. (You might think “Hey, why not do it first thing when I start up?” Well, turns out that this is only returns the proper amount of free memory after the context is up, which may not be when I think it is, if you’re using the RT)

Coming back to the original problem: When the device almost runs out of memory, it can happen that it simply dies. I’m seeing this in some situations where I still have something like 8MB of VRAM unallocated and as soon as I start my kernel, the computer freezes. This is with valgrind-verified code, BTW.

I used to have other weird "depends on the lunar phase’ kind of bugs with my 8600GT, which I was never able to attribute to anything. Sometimes, I would simply get repeated unspecified launch errors. Rebooting would always help, sometimes running another kernel was enough… It was serious voodoo-stuff. I considered getting some rubber chickens to sacrifice to the damn thing.

FWIW, switching to a 285 removed all those bugs.

[MisterAnderson42]

I used to have other weird "depends on the lunar phase’ kind of bugs with my 8600GT, which I was never able to attribute to anything. Sometimes, I would simply get repeated unspecified launch errors. Rebooting would always help, sometimes running another kernel was enough… It was serious voodoo-stuff. I considered getting some rubber chickens to sacrifice to the damn thing.

FWIW, switching to a 285 removed all those bugs.

We’ve documented this behavior in other threads, and also (through the help of tmurray) found that all hardware released after (and including) Tesla C1060 does not exhibit the problem. It is an extremely rare, but extraordinarily annoying problem.

http://forums.nvidia.com/index.php?showtop…+launch+failure

[E.D_Riedijk]

FWIW, switching to a 285 removed all those bugs.

Could you do a bandwidthtest for me? I am considering to replace all PC’s at my work with 285-equipped boxes. Good experiences are certainly recomforting before placing the order (as are very nice bandwidth numbers ;))

[T.B]

Could you do a bandwidthtest for me? I am considering to replace all PC’s at my work with 285-equipped boxes. Good experiences are certainly recomforting before placing the order (as are very nice bandwidth numbers ;))

Sure. But be aware that I’m not on a PCIE 2.0 board. So Host<->Device is severely gimped. The one thing that is annoying about the card (and all 2xx cards) is that the profiler doesn’t tell you the number of uncoalesced memory accesses any more.

Other than that, it’s a really nice and fast piece of hardware.

[codebox]Running on…

  device 0:GeForce GTX 285

Quick Mode

Host to Device Bandwidth for Pageable memory

.

Transfer Size (Bytes) Bandwidth(MB/s)

33554432 1850.7

Quick Mode

Device to Host Bandwidth for Pageable memory

.

Transfer Size (Bytes) Bandwidth(MB/s)

33554432 1657.4

Quick Mode

Device to Device Bandwidth

.

Transfer Size (Bytes) Bandwidth(MB/s)

33554432 124344.3[/codebox]

[seibert]

For comparison, here’s a GTX 280:

[stan@lurch release]$ ./bandwidthTest --memory=pinned

Using device 0: GeForce GTX 280

Quick Mode

Host to Device Bandwidth for Pinned memory

.

Transfer Size (Bytes)	Bandwidth(MB/s)

 33554432		5290.2

Quick Mode

Device to Host Bandwidth for Pinned memory

.

Transfer Size (Bytes)	Bandwidth(MB/s)

 33554432		5969.5

Quick Mode

Device to Device Bandwidth

.

Transfer Size (Bytes)	Bandwidth(MB/s)

 33554432		114469.7

So imagine these HtoD and DtoH numbers glued to the DtoD number posted by T.B. while you hit the “purchase” button. :)

[E.D_Riedijk]

Sure. But be aware that I’m not on a PCIE 2.0 board. So Host<->Device is severely gimped. The one thing that is annoying about the card (and all 2xx cards) is that the profiler doesn’t tell you the number of uncoalesced memory accesses any more.

Other than that, it’s a really nice and fast piece of hardware.

Thanks, that is some nice device-device number! I know about the profiler thing (have a C1060 and FX4800), but that is because there are no uncoalesced memory accesses anymore ;)

[Austin]

I know about the profiler thing (have a C1060 and FX4800), but that is because there are no uncoalesced memory accesses anymore ;)

On the page 55 -56 of CUDA Programming Guide 2.1, there is a description:

Coalescing on Devices with Compute Capability 1.2 and Higher

             The global memory access by all threads of a half-warp is coalesced into a single

             memory transaction as soon as the words accessed by all threads lie in the same

             segment of size equal to:

                 32 bytes if all threads access 8-bit words,

                 64 bytes if all threads access 16-bit words,

                 128 bytes if all threads access 32-bit or 64-bit words.

             Coalescing is achieved for any pattern of addresses requested by the half-warp,

             including patterns where multiple threads access the same address. This is in

             contrast with devices of lower compute capabilities where threads need to access

             words in sequence.

and

More precisely, the following protocol is used to issue a memory transaction for a

half-warp:

Find the memory segment that contains the address requested by the lowest

numbered active thread. Segment size is 32 bytes for 8-bit data, 64 bytes for

16-bit data, 128 bytes for 32-, 64- and 128-bit data.

Find all other active threads whose requested address lies in the same segment.

Reduce the transaction size, if possible:

   If the transation size is 128 bytes and only the lower or upper half is used,

   reduce the transaction size to 64 bytes;

   If the transation size is 64 bytes and only the lower or upper half is used,

   reduce the transaction sizez to 32 bytes.

Carry out the transaction and mark the serviced threads as inactive.

Repeat until all threads in the half-warp are serviced.

In other words, the cards with a computation capability 1.2 and higher make an optimization an global memory access, it make its great effort to reduce the uncoalescence memory access.

[Austin]

I know about the profiler thing (have a C1060 and FX4800), but that is because there are no uncoalesced memory accesses anymore ;)

On the page 55 -56 of CUDA Programming Guide 2.1, there is a description:

Coalescing on Devices with Compute Capability 1.2 and Higher

             The global memory access by all threads of a half-warp is coalesced into a single

             memory transaction as soon as the words accessed by all threads lie in the same

             segment of size equal to:

                 32 bytes if all threads access 8-bit words,

                 64 bytes if all threads access 16-bit words,

                 128 bytes if all threads access 32-bit or 64-bit words.

             Coalescing is achieved for any pattern of addresses requested by the half-warp,

             including patterns where multiple threads access the same address. This is in

             contrast with devices of lower compute capabilities where threads need to access

             words in sequence.

and

More precisely, the following protocol is used to issue a memory transaction for a

half-warp:

Find the memory segment that contains the address requested by the lowest

numbered active thread. Segment size is 32 bytes for 8-bit data, 64 bytes for

16-bit data, 128 bytes for 32-, 64- and 128-bit data.

Find all other active threads whose requested address lies in the same segment.

Reduce the transaction size, if possible:

   If the transation size is 128 bytes and only the lower or upper half is used,

   reduce the transaction size to 64 bytes;

   If the transation size is 64 bytes and only the lower or upper half is used,

   reduce the transaction sizez to 32 bytes.

Carry out the transaction and mark the serviced threads as inactive.

Repeat until all threads in the half-warp are serviced.

In other words, the cards with a computation capability 1.2 and higher make an optimization an global memory access, it make its great effort to reduce the uncoalescence memory access.

[E.D_Riedijk]

In other words, the cards with a computation capability 1.2 and higher make an optimization an global memory access, it make its great effort to reduce the uncoalescence memory access.

Yes, but where before it was a boolean (coalesced or not) in 1.2 or higher devices you would need to have an indication of how many memory transactions were issued. Because a previously uncoalesced access can in >=1.2 generate between 1 and 16 transactions.