Commit Graph

295 Commits

Author SHA1 Message Date
fb3f32760d Cycles: Add an experimental CUDA kernel.
Now we build 2 .cubins per architecture (e.g. kernel_sm_21.cubin, kernel_experimental_sm_21.cubin).
The experimental kernel can be used by switching to the Experimental Feature Set: http://wiki.blender.org/index.php/Doc:2.6/Manual/Render/Cycles/Experimental_Features

This enables Subsurface Scattering and Correlated Multi Jitter Sampling on GPU, while keeping the stability and performance of the regular kernel.

Differential Revision: https://developer.blender.org/D762
Patch by Sergey and myself.

Developer / Builder Note:
CUDA Toolkit 6.5 is highly recommended for this, also note that building the experimental kernel requires a lot of system memory (~7-8GB).
2014-08-26 17:02:26 +02:00
603348c56e Cycles: Drop support for CUDA 5.0 Toolkit, only 6.0 and 6.5 (recommended) are supported now. 2014-08-21 23:35:20 +02:00
Dalai Felinto
8d3cc431d7 Fix T41471 Cycles Bake: Setting small tile size results in wrong bake with stripes rather than the expected noise pattern
This problem was introduced in 983cbafd18
Basically the issue is that we were not getting a unique index in the
baking routine for the RNG (random number generator).

Reviewers: sergey

Differential Revision: https://developer.blender.org/D749
2014-08-19 11:40:33 +02:00
Dalai Felinto
2c5b6859d9 Revert "Fix T41222 Blender gives weird output when baking (4096*4096) resolution on GPU"
This reverts commit a48b372b04.

Leaving only the part that fix device_multi.cpp
2014-08-15 11:27:42 +02:00
Dalai Felinto
a48b372b04 Fix T41222 Blender gives weird output when baking (4096*4096) resolution on GPU
In collaboration with Sergey Sharybin.

Also thanks to Wolfgang Faehnle (mib2berlin) for help testing the
solutions.

Reviewers: sergey

Differential Revision: https://developer.blender.org/D690
2014-08-05 13:50:50 -03:00
77b7e1fe9a Deduplicate CUDA and OpenCL wranglers
For now it was mainly about OpenCL wrangler being duplicated
between Cycles and Compositor, but with OpenSubdiv work those
wranglers were gonna to be duplicated just once again.

This commit makes it so Cycles and Compositor uses wranglers
from this repositories:

  - https://github.com/CudaWrangler/cuew
  - https://github.com/OpenCLWrangler/clew

This repositories are based on the wranglers we used before
and they'll be likely continued maintaining by us plus some
more players in the market.

Pretty much straightforward change with some tricks in the
CMake/SCons to make this libs being passed to the linker
after all other libraries in order to make OpenSubdiv linked
against those wranglers in the future.

For those who're worrying about Cycles being less standalone,
it's not truth, it's rather more flexible now and in the future
different wranglers might be used in Cycles. For now it'll
just mean those libs would need to be put into Cycles repository
together with some other libs from Blender such as mikkspace.

This is mainly platform maintenance commit, should not be any
changes to the user space.

Reviewers: juicyfruit, dingto, campbellbarton

Reviewed By: juicyfruit, dingto, campbellbarton

Differential Revision: https://developer.blender.org/D707
2014-08-05 13:57:50 +06:00
9c3025cd26 Spelling 2014-08-02 16:53:52 +10:00
Dalai Felinto
fc55c41bba Cycles Bake: show progress bar during bake
Baking progress preview is not possible, in parts due to the way the API
was designed. But at least you get to see the progress bar while baking.

Reviewers: sergey

Differential Revision: https://developer.blender.org/D656
2014-07-25 11:42:53 -03:00
Martijn Berger
bae2b3a688 Switch to Cuda 4.0 style api for kernel invocation. This is a small clean-up that has no functional changes but makes code a bit more readable.
Differential revision: https://developer.blender.org/D659

Reviewed by: Sergey Sharybin, Thomas Dinges
2014-07-25 13:33:19 +02:00
9acabc13de Cleanup: Typo fixes. 2014-07-05 14:25:34 +02:00
5898abe99d Cycles: Update CUDA error messages, based on Toolkit 6.0.
* Removed deprecated erros, and added some new ones, which might help to figure out problems in the future.
2014-07-02 01:50:42 +02:00
4800c52700 Cleanup: Remove unused checks in CUDA device code. 2014-07-02 01:12:13 +02:00
e4e58d4612 Fix T40370: cycles CUDA baking timeout with high number of AA samples.
Now baking does one AA sample at a time, just like final render. There is
also some code for shader antialiasing that solves T40369 but it is disabled
for now because there may be unpredictable side effects.
2014-06-06 15:39:04 +02:00
865dfa8a7e Fix T40228: cycles CUDA multi GPU + world MIS giving error. 2014-06-05 18:10:32 +02:00
69c7522b24 Fix T40379: world MIS causing too much CUDA memory usage.
The kernel for baking the world texture was the same as the one used for
baking. Now that's separate which allows the kernel to reserve much less
memory.
2014-05-27 15:11:32 +02:00
3b53fffb77 Cycles: revert async CUDA changes, these are giving too much trouble still.
Fixes T40027. This means we get more CPU usage again when using multiple CUDA,
but the impact on performance is too big a problem with the current code.
2014-05-19 19:33:09 +02:00
c08c931fb6 Cycles / CUDA: Increase maximum image textures on GPU.
Instead of 95, we can use 145 images now. This only affects Kepler and above (sm30, sm_35 and sm_50).

This can be increased further if needed, but let's first test if this does not come with a performance impact.

Originally developed during my GSoC 2013.
2014-05-11 03:38:39 +02:00
fd26a32aa5 Fix T40119, CUDA Toolkit version mismatch 2014-05-10 01:26:04 +02:00
dc13969e48 Style cleanup: indentation, braces 2014-05-05 02:19:08 +10:00
1618329b00 Code cleanup: style, require ; for cuda_assert, opencl_assert 2014-05-04 03:57:50 +10:00
198f5e506a Cycles: CUDA changes for kernel evaluation cancel 2014-05-02 21:19:10 -03:00
8d16869d83 Code cleanup: Add -Werror=float-conversion to Cycles 2014-05-03 07:31:46 +10:00
741f17f05b Cycles CUDA: make CUDA toolkit 6.0 the official supported version.
This also updates the configurations to build kernels for compute capability
5.0 cards, when using and older CUDA toolkit version this will be skipped.

Also includes tweaks to improve performance with this version:
* Increase max registers on sm_30, sm_35 and sm_50
* No longer use texture storage on sm_30
2014-04-30 16:07:27 +02:00
39bfde674c Cycles CUDA: don't use cuLaunchGridAsync at all for display devices.
As suggested by Martijn, this is slower than cuLaunchGrid.
2014-04-17 12:18:49 +02:00
18da79f471 Cycles CUDA: only do async execution for GPUs not used for display.
Otherwise devices used for display will lock up the UI too much. This means
you might still get 100% CPU for the display device, but for others CPU usage
should be low still.

The check to see if a device is used for display may not be entirely reliable,
it checks if there is a watchdog timeout on the device, but I'm not entirely
sure that always exists for display devices or is disabled for non-display
devices, though some tools like cuda-gdb seem to make the same assumption.

Ref T39559
2014-04-17 12:08:18 +02:00
415e10a0ef Fix another compile error with recent commit on visual studio. 2014-04-16 21:36:19 +02:00
6f1afdbbfc Cycles CUDA: enabled branched path kernel again, with more registers. 2014-04-16 21:05:04 +02:00
2851ed4a55 Cycles code refactor: use __launch_bounds__ instead of -maxrregcount for CUDA.
This makes it easier to have per kernel number of registers. Also, all the
tunable parameters for this are now in kernel.cu, rather than spread over cmake,
scons and device_cuda.cpp.
2014-04-16 21:05:04 +02:00
297a2223b5 Cycles / CUDA: Increase sm_2x registers to 40.
This fixes the ptaxs "ACCESS_VIOLATION" error and should allow our Linux and Windows build bots to compile again.
Unfortunately this comes with a performance penalty on sm_2x cards, so this is only a workaround for now. Branched Path is still globally disabled on GPU.
2014-04-08 23:25:54 +02:00
d923720312 Cycles: Disable Branched Path on all GPUs for now, until we separate the cubins.
SM_20 fails now as well, reported by Zanqdo in IRC.
2014-04-03 22:18:40 +02:00
a2e4ebd36a Cycles code internals: add CPU kernel support for 3D image textures. 2014-03-29 13:03:48 +01:00
859039f732 Cycles: Raise a proper error message when using Branched Path on sm_30, this is currently still disabled. 2014-03-27 10:29:22 +01:00
74518b2826 Fix T39420: Cycles viewport/preview flickers, when moving mouse across editors
Issue was caused by the wrong usage of OCIO GLSL binding API. To make it
work properly on pre-GLSL-1.3 drivers shader is to be enabled after the
texture is binded to the opengl context. Otherwise it wouldn't know the
proper texture size.

This is actually a regression in 2.70 and to be ported to 'a'.
2014-03-26 15:58:53 +06:00
28c1a860e2 Fix T39247
Changes to interpolation break texture allocation on sm35 and greater.
2014-03-19 07:37:18 +01:00
dd2dca2f7e Add support for multiple interpolation modes on cycles image textures
All textures are sampled bi-linear currently with the exception of OSL there texture sampling is fixed and set to smart bi-cubic.

This patch adds user control to this setting.

Added:
- bits to DNA / RNA in the form of an enum for supporting multiple interpolations types
- changes to the image texture node drawing code ( add enum)
- to ImageManager (this needs to know to allocate second texture when interpolation type is different)
- to node compiler (pass on interpolation type)
- to device tex_alloc this also needs to get the concept of multiple interpolation types
- implementation for doing non interpolated lookup for cuda and cpu
- implementation where we pass this along to osl ( this makes OSL also do linear untill I add smartcubic to the interface / DNA/ RNA)

Reviewers: brecht, dingto

Reviewed By: brecht

CC: dingto, venomgfx

Differential Revision: https://developer.blender.org/D317
2014-03-07 23:16:33 +01:00
1d01675833 Cuda use streams and async to avoid busywaiting
This switches api usage for cuda towards using more of the Async calls.

Updating only once every second is sufficiently cheap that I don't think it is worth doing it less often.

Reviewed By: brecht

Differential Revision: https://developer.blender.org/D262
2014-03-06 20:51:46 +01:00
6b1a4fc66e Cycle CUDA: revert the f1aeb2ccf4 and 84f958754 busywait fixes for now.
It's unclear what kind of impact they have on performance at the moment, so I
rather play it safe and postpone this for 2.71.

Ref T38679, Ref T38712
2014-02-19 16:08:08 +01:00
f1aeb2ccf4 this is an attempted Fix: T38679
Cycles GPU Performance Regression

From my testing this (what i should have done in the first place) reduces the regression a lot.
Lets hope it is enough or we have to go back to busy waiting.
2014-02-17 20:11:45 +01:00
84f9587540 Cuda use streams and async to avoid busywaiting
This is my first stab at this and is based on this IRC converstation:

<mib2berlin> brecht: this is meaning as reminder only, I know you have other things to do > http://openvidia.sourceforge.net/index.php/Optimization_Notes#avoiding_busy_waits
<brecht> mib2berlin: thanks, bookmarked

only tested on Ubuntu 14.04 / cuda 5.0 but ill do some more testing tomorrow.

Also unsure about the placement and the lifetime of the stream and the event. But creating / deleting these seems to incur a non trivial cost.

Reviewers: brecht

Reviewed By: brecht

CC: mib2berlin, dingto

Differential Revision: https://developer.blender.org/D262
2014-01-28 18:40:08 +01:00
241fccaf6a Fix T37817: cycles CUDA detection problem on Windows with non-ascii paths. 2014-01-11 00:47:58 +01:00
85a0c5d4e1 Cycles: network render code updated for latest changes and improved
This actually works somewhat now, although viewport rendering is broken and any
kind of network error or connection failure will kill Blender.

* Experimental WITH_CYCLES_NETWORK cmake option
* Networked Device is shown as an option next to CPU and GPU Compute
* Various updates to work with the latest Cycles code
* Locks and thread safety for RPC calls and tiles
* Refactored pointer mapping code
* Fix error in CPU brand string retrieval code

This includes work by Doug Gale, Martijn Berger and Brecht Van Lommel.

Reviewers: brecht

Differential Revision: http://developer.blender.org/D36
2013-12-07 12:26:58 +01:00
b5a5773fa9 Cycles / CUDA:
* Remove support for  CUDA Toolkit 4.x, only Toolkit 5.0 and above are supported now.
* Remove support for sm_1x cards (< Fermi) for good. We didn't officially support those cards for a few releases already, now remove some special code that was still there.
2013-10-08 15:29:28 +00:00
fa352bb749 Fix #35684: cycles unable to use full 6GB of memory on NVidia Titan GPU. We now
use arrays instead of textures for general storage on this card (image textures
are still stored as texture). Textures were found to be faster on older cards,
but the limits on 1D texture size have not increased along with the memory size,
which meant that the full 6 GB could not be used.

The performance actually seems to be slightly better with arrays in some tests
on Titan. For older cards there seems to be a bit of a mix, some are better and
others not. We may change those to use arrays too, but more testing is needed,
only Titan and Tesla K20 (sm_35) is changed for now.

The fact that arrays are faster is a bit surprising, as others found textures
to be faster on Kepler. However even if they were, the memory limitation is
more important to solve anyway.
https://research.nvidia.com/publication/understanding-efficiency-ray-traversal-gpus-kepler-and-fermi-addendum
2013-09-27 19:09:31 +00:00
29f6616d60 Cycles: viewport render now takes scene color management settings into account,
except for curves, that's still missing from the OpenColorIO GLSL shader.

The pixels are stored in a half float texture, converterd from full float with
native GPU instructions and SIMD on the CPU, so it should be pretty quick.
Using a GLSL shader is useful for GPU render because it avoids a copy through
CPU memory.
2013-08-30 23:49:38 +00:00
6785874e7a Fix #36137: cycles render not using all GPU's when the number of GPU's is larger
than the number of CPU threads
2013-08-30 23:09:22 +00:00
01e22d1b9f Cycles: more code refactoring to rename things internally as well. Also change
property name back so we keep compatibility.
2013-08-23 14:34:34 +00:00
b9ce231060 Cycles: relicense GNU GPL source code to Apache version 2.0.
More information in this post:
http://code.blender.org/

Thanks to all contributes for giving their permission!
2013-08-18 14:16:15 +00:00
743a7a4a4b Cycles:
* GPU kernel can now be compiled without __NON_PROGRESSIVE__ again, was broken after my last commit. Also add a check for have_error(), in case the GPU kernel comes without Non-Progressive, to avoid a crash.

* Don't compile progressive kernel twice on CPU, if __NON_PROGRESSIVE__ would be disabled there.
2013-08-09 20:03:49 +00:00
a18112249d Cycles / Non-Progressive integrator:
* Non-Progressive integrator is now available on the GPU (CUDA, sm_20 and above). 

Implementation details:
* kernel_path_trace() has been split up into two functions:
kernel_path_trace_non_progressive() and kernel_path_trace_progressive().

* We compile two CUDA kernel entry functions (in kernel.cu) for the two integrators, they are still inside one .cubin file but due to the kernel separation there should be no performance problem. I tested with the BMW file on my Geforce 540M and the render times were the same for 100 samples (1.57 min in my case).

This is part of my GSoC project, SVN merge of r59032 + manual merge of UI changes for this from my branch.
2013-08-09 18:47:25 +00:00
7902fa57b6 Code cleanup: cycles
* Reshuffle SSE #ifdefs to try to avoid compilation errors enabling SSE on 32 bit.
* Remove CUDA kernel launch size exception on Mac, is not needed.
* Make OSL file compilation quiet like c/cpp files.
2013-06-26 23:29:33 +00:00