This seems to be solid enough to allow me to launch a render of the Gooseberry benchmark scene on my 4GB GPU with TDR turned off. However, some tiles render extremely slow (counting seconds per sample instead of samples per second!), I haven't found out yet what crazy things happening in them.

This seems incomplete? The patch in T48651 contained more changes.

Yes. I'm still trying to figure out how this system and git patches work. It looks like when I try to upload my diff file with multiple commits in it, it takes only the first one.

A couple of ideas for improvement:

• It may be possible to do 3D texture allocation through cuMemAlloc()/cuMemAllocHost() too.
• I'd love to have a better heuristic for finding out the actual memory requirements of our kernel launch. Memory fragmentation may be important too when estimating the required amount.
• When multiple CUDA devices are in use, host allocations should be shared between devices.
• Eventually, one could support Pascal/CUDA 8's virtual memory by small changes to mem_alloc_internal()
• Some high end cards (Telsa only?) support Peer-To-Peer memory access, multi-device renders could then share allocations in VRAM.
• There could be user preferences for enabling/disabling pinned memory and setting the max amount.

Do we have to pin memory ahead of a time and doomed to never change that? Basically, what i mean is: can't we pin more and more memory once our requirement grows (similar to what we do in ccl::array already)?

Other idea here is to do delayed device allocation. This way we'll know overall device memory requirement and can calculate how much of RAM we'll have to pin.

I'm not a strong fan of user preferences for that. They might be acceptable for local operation on desktop, but once you have to deal with renderfarms need to tweak each individual node is quite annoying. Not as if CUDA renderfarms are common at this moment, but still. Also, even for local operation, you shouldn't be in a situations like simple scene using too much of your RAM or not being able to render just because you've set low memory limit. IMO, the way to go here would be to pin RequiredDeviceMemory-VRAM megs of RAM.

It only pins the amount of memory required, no more. The user preference would only set an upper bound to that.

Well, that's good only required memory is pinned. But then i'm not sure why to have limit? We don't limit CPU memory and added code to gracefully (as much as we can) to abort render when allocation fails due to bad_alloc. Why not to use same for host memory with GPU rendering? Basically, allow to pin up to whole RAM available but stop rendering when even that is not enough?

From the CUDA documentation it sounds like pinning all available memory may not be a good idea:

"Allocating excessive amounts of memory with cudaMallocHost() may degrade system performance, since it reduces the amount of memory available to the system for paging. As a result, this function is best used sparingly to allocate staging areas for data"

When pinning too much, it will affect overall system performance, not just the rendering process.

For the lowest memory usage and least copying we can make device_vector point to pinned memory directly. This would require some refactoring of device_vector so all memory allocations and frees go through the device, but it's not a huge change I think. If pinned memory is being shared between multiple CUDA devices it seems logical to do it this way.

@Stefan Werner (swerner): I wouldn't worry about that at all. When you have a small RAM and try to render a large scene on CPU, your system memory gets filled too, swapping begins and the system becomes unresponsive / slow. I don't see why we should treat GPU host allocations any different. I am not a fan of adding a user pref option either.

It might also make sense to use pinned memory for all device_vector allocations even if we want to copy memory to the GPU. That way we can use async copies which should be faster, particularly for multi GPU.

For limiting memory, I find that once you start swapping performance is unacceptably slow already, often to the point that the computer is almost unusable. So it would be interesting to see what happens when you degrade system performance even worse than that :). In theory it would be really bad if all memory except e.g. 1 MB is pinned and all memory need to swapped into that tiny amount of space, but I wonder if it can really get much worse than typical memory swapping in practice.

If we need to do this then a reasonable heuristic might be to allow pinning total system memory minus 2 GB, to leave enough space for the OS and other applications to run, but it's all quite arbitrary.

Would it be worthwhile having a user control to specify how much ram to use on the GPU? because of the 970 3.5gb/4gb issue?

That last 512 MB on the GTX 970 is still faster than using host memory, so I don't think it would be useful.

Different operating systems may handle this differently, but for what it's worth, I ran a simple test on OS X. cuMemAllocHost() in a loop, allocating 200MB chunks, 100 times - which should try to allocate ~20GB in total. My machine has 16GB of physical memory. At just under 10GB allocated, the machine froze completely. That is, not even the mouse pointer will move and the machine does not respond any more to network requests. After about one or two two minutes of being frozen, the machine rebooted, the "crash" log pointing to the Watchdog task.

So at least on OS X, allocating excessive pinned memory is different from excessive regular memory and a hard limit on the maximum amount is in my opinion very much preferred over a "one more texture and your render will crash the entire machine" user experience.

In D2056#47274, @Brecht Van Lommel (brecht) wrote:

It might also make sense to use pinned memory for all device_vector allocations even if we want to copy memory to the GPU. That way we can use async copies which should be faster, particularly for multi GPU.

Yes, that would be an improvement that we can split out into a separate followup task. One must take care though of pitch and alignment, the CPU side would need to be aware of the memory layout for textures.

@Stefan Werner (swerner), that's a bummer it freezes so completely, but here are some points:

• As you've mentioned, it might depend on OS/driver and even their version. I don't think we should commit ourselves to maintenance of all those per-platform heuristics, they never work reliably enough anyway. We've got that happening for AMD OpenCL, and it's really rather pathetic. Heuristic created for one driver does not work for another one, making me almost thinking of dropping all of them and demand companies to deliver proper drivers. We are too small to keep adding workarounds for each new driver.
• Drivers might be just buggy. They definitely shouldn't crash entire OS kernel, and if they do -- it totally worth submitting report to NVidia IMO.
• For any setting you need a default setting. Now, what is the value in this case? Is it same value for all platforms? is it something user would need to set up himself (then it kinda defeats whole idea of avoid user frustration because he will first run into the issue before discovering that option)?
• Is this limit dictated by a memory size after all or by bus throughput and power of GPU? As in, imagine you've increased your machine memory to 32gig keeping everything else the same. Will crash system freeze at 20gig now? Or because of some "heavy" load crash still will happen at 10gig?

Also, is this some known issue and everyone is "solving" it by adding hard limits to host memory?

Octane exposes what they call "out-of-core" textures, which I assume is textures in pinned memory, as a feature that the user has to enable and to pick how much memory it is supposed to use: https://docs.otoy.com/Standalone_2_0/?page_id=3216

In Poser we ship our version of Cycles with a hardcoded 1/2 physical limit.

Stefan Werner, Do you have in your OS the virtual memory located on a 7200 rpm disk?
Sorry if I misunderstood, but your system freeze reminds me of those I have in Linux when the system begins to fill the Swap intensively on a slow disk (as one of 7200rpm). That always happens in Linux when intensively fills the swap, swap management is not good on Linux (you look in google how a lot of people try to avoid using the swap by configuring swappiness).
What I mean is that if the freeze on your OS was due to the intensive use of the swap, then maybe in a system with more RAM the problem does not occur.

Edit: sorry, apparently I have subscribed to the wrong person when I tried to quote Stefan Werner. I have to learn how to do this properly.

There is no swapping with pinned memory. That's the whole point about pinned memory - it always remains fixed in physical memory, at the same address, no matter how severe the OS' memory pressure is. That said, it was an SSD.

My computer just finished rendering the Victor/Gooseberry benchmark scene on two GPUs, coming in at about 1h:36m (conservative tile size of 64x64). So I think my patch is pretty solid.

I did have to enable SHADOW_RECORD_ALL for CUDA with an intersection buffer size of 10 to prevent timeouts, though - the incremental stepping through transparent shadows absolutely kills GPU performance. I may try to tackle that next.

Update: Bumped the tile size to 128x128, now a single GPU (K5000, 4GB) renders this scene in 1h:53m. This beats my CPU which takes 3h:15m. MemPeak is at 6681MB, so even with PCI bus transfers, GPU rendering can be faster than CPU rendering (if you have a CPU as weak as mine, I guess...).

@Stefan Werner (swerner), that's not really comparing apples to apples.. Also, we can't that easily enable __SHADOW_RECORD_ALL__ on GPU because:

1. With transparent bounces larger than 64 it'll use malloc() which isn't fast (read as: really slow)
2. It increases ctask usage by 1.5KB per thread (which is 5% of overall current stack usage as far as i remember)

That being said, some code feedback.

Also seems dealing with multi-device is not in there yet?

And still not sure having memory limit setting is indeed a way to go, there are various ways it'll fail on various levels:

• it's yet-another-magic-setting which users would have to adjust based on (a) particular scene (b) particular system and there's no goo default value
• It is based on physical memory, not free one. Meaning, cases when Blender is not an exclusive running application (or there are multiple Blender instances running) current approach of "safety" will lead to a machine death yet again
• Documentation doesn't mention any issues using more than N% of host memory and would expect everything to be good as long as there's enough free RAM. IMO it totally worth checking if the computer hungs you're experiencing is a specific of behavior or is a bug of some sort. at least form reading documentation it's unclear why pinning memory but keeping enough memory for running processes will cause any troubles.

One more thing, documentation mentions context is to be created with CU_CTX_MAP_HOST flag, which i don't see in this patch.1

Did quick test with pinning memory using OS API (namely mlock since i'm on Linux). I had 3+Gb used by running processes and i've successfully pinned 10Gb out of 12Gb of RAM without any noticeable regression. Pinning 11Gb of RAM caused some hiccups in running youtube, but that's kinda expected. In any case, the system survived this just fine.

Wouldn't expect CUDA API doing something much more special than using kernel's calls to pin pages anyway.

Here's a test code i've used

The change to SHADOW_RECORD_ALL wasn't meant as a permanent solution, but to compare how GPU performance would be in the largest Blender scene available to me (will happily render any other scenes). When SHADOW_RECORD_ALL is different for CPU and GPU, we're not tracing the same number of rays, and benchmarks with high levels of transparent shadows will always perform better on the CPU.

If you are in favour of not having limits on pinned memory, sure, I can change it to that. Your sample code allocating pinned memory worked without hiccups on my OS X machine (14 GB out of 16 GB just fine), so I'm not sure if the CUDA Runtime does something additional there or if my machine had a bad day when I tested it last time.

I will add the CU_CTX_MAP_HOST, thanks for pointing that out.

Looks like cuMemAllocHost() is not the same as malloc() followed by mlock(). While malloc/mlock easily allows me to get 14 out of 16GB on my machine, cuMemAllocHost() freezes the entire machine at 12GB.

Great idea for a patch! What blender version does it currently apply to?
It doesn't seem to work for me on 2.79, 2.77 nor 2.80 release candidate.

(Host OS: Linux Ubuntu 16.04.3 64bit)

Thanks for your work, this is looking good! Especially the addition for sharing memory with CPU devices is good.

I've added code to dynamically move texture memory from device to host during a render, and now it seems to be robust even when there are big denoising buffers that need to be allocated during a render.

Committed and release notes updated.