[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [59034] trunk/blender: Cycles / Non-Progressive integrator:
Thomas Dinges
blender at dingto.org
Fri Aug 9 20:47:25 CEST 2013
Revision: 59034
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=59034
Author: dingto
Date: 2013-08-09 18:47:25 +0000 (Fri, 09 Aug 2013)
Log Message:
-----------
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.
Revision Links:
--------------
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=59032
Modified Paths:
--------------
trunk/blender/intern/cycles/blender/addon/ui.py
trunk/blender/intern/cycles/blender/blender_sync.cpp
trunk/blender/intern/cycles/device/device_cuda.cpp
trunk/blender/intern/cycles/device/device_task.h
trunk/blender/intern/cycles/kernel/kernel.cl
trunk/blender/intern/cycles/kernel/kernel.cpp
trunk/blender/intern/cycles/kernel/kernel.cu
trunk/blender/intern/cycles/kernel/kernel_path.h
trunk/blender/intern/cycles/kernel/kernel_shader.h
trunk/blender/intern/cycles/kernel/kernel_sse2.cpp
trunk/blender/intern/cycles/kernel/kernel_sse3.cpp
trunk/blender/intern/cycles/kernel/kernel_types.h
trunk/blender/intern/cycles/render/session.cpp
Property Changed:
----------------
trunk/blender/
Property changes on: trunk/blender
___________________________________________________________________
Modified: svn:mergeinfo
- /branches/ge_dev:58091-58422
/branches/ge_harmony:42255,42279-42282,42286,42302,42338,42349,42616,42620,42698-42699,42739,42753,42773-42774,42832,44568,44597-44598,44793-44794
/branches/soc-2011-cucumber:37517,38166-38167,38177,38179-38180,38187,38242,38384,38387,38403-38404,38407,38968,38970,38973,39045,40845,42997-42998,43439
/branches/soc-2011-tomato:42376,42378-42379,42383,42385,42395,42397-42400,42407,42411,42418,42443-42444,42446,42467,42472,42486,42650-42652,42654-42655,42709-42710,42733-42734,42801,43872,44130,44141,44147-44149,44151-44152,44229-44230,45623-45625,46037,48089,48092,48551-48552,48679,48790,48792-48793,49076,49087,49292,49294,49466,49894,50052,50126,52854-52856,54573
/branches/soc-2013-depsgraph_mt:57516
/branches/soc-2013-dingto:57424,57487,57507,57525,57599,57670,57918-57919,57981,58091,58245,58253,58587,58772,58774-58775,58828,58835
/tags/blender-2.67b-release/blender:57122
+ /branches/ge_dev:58091-58422
/branches/ge_harmony:42255,42279-42282,42286,42302,42338,42349,42616,42620,42698-42699,42739,42753,42773-42774,42832,44568,44597-44598,44793-44794
/branches/soc-2011-cucumber:37517,38166-38167,38177,38179-38180,38187,38242,38384,38387,38403-38404,38407,38968,38970,38973,39045,40845,42997-42998,43439
/branches/soc-2011-tomato:42376,42378-42379,42383,42385,42395,42397-42400,42407,42411,42418,42443-42444,42446,42467,42472,42486,42650-42652,42654-42655,42709-42710,42733-42734,42801,43872,44130,44141,44147-44149,44151-44152,44229-44230,45623-45625,46037,48089,48092,48551-48552,48679,48790,48792-48793,49076,49087,49292,49294,49466,49894,50052,50126,52854-52856,54573
/branches/soc-2013-depsgraph_mt:57516
/branches/soc-2013-dingto:57424,57487,57507,57525,57599,57670,57918-57919,57981,58091,58245,58253,58587,58772,58774-58775,58828,58835,59032
/tags/blender-2.67b-release/blender:57122
Modified: trunk/blender/intern/cycles/blender/addon/ui.py
===================================================================
--- trunk/blender/intern/cycles/blender/addon/ui.py 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/blender/addon/ui.py 2013-08-09 18:47:25 UTC (rev 59034)
@@ -67,9 +67,7 @@
row.operator("render.cycles_sampling_preset_add", text="", icon="ZOOMOUT").remove_active = True
row = layout.row()
- sub = row.row()
- sub.active = (device_type == 'NONE' or cscene.device == 'CPU')
- sub.prop(cscene, "progressive")
+ row.prop(cscene, "progressive")
if not cscene.progressive:
row.prop(cscene, "squared_samples")
@@ -82,7 +80,7 @@
sub.prop(cscene, "seed")
sub.prop(cscene, "sample_clamp")
- if cscene.progressive or (device_type != 'NONE' and cscene.device == 'GPU'):
+ if cscene.progressive:
col = split.column()
sub = col.column(align=True)
sub.label(text="Samples:")
@@ -656,7 +654,6 @@
lamp = context.lamp
clamp = lamp.cycles
cscene = context.scene.cycles
- device_type = context.user_preferences.system.compute_device_type
layout.prop(lamp, "type", expand=True)
@@ -675,7 +672,7 @@
sub.prop(lamp, "size", text="Size X")
sub.prop(lamp, "size_y", text="Size Y")
- if not cscene.progressive and (device_type == 'NONE' or cscene.device == 'CPU'):
+ if not cscene.progressive:
col.prop(clamp, "samples")
col = split.column()
@@ -864,7 +861,6 @@
world = context.world
cworld = world.cycles
cscene = context.scene.cycles
- device_type = context.user_preferences.system.compute_device_type
col = layout.column()
@@ -872,7 +868,7 @@
sub = col.row(align=True)
sub.active = cworld.sample_as_light
sub.prop(cworld, "sample_map_resolution")
- if not cscene.progressive and (device_type == 'NONE' or cscene.device == 'CPU'):
+ if not cscene.progressive:
sub.prop(cworld, "samples")
Modified: trunk/blender/intern/cycles/blender/blender_sync.cpp
===================================================================
--- trunk/blender/intern/cycles/blender/blender_sync.cpp 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/blender/blender_sync.cpp 2013-08-09 18:47:25 UTC (rev 59034)
@@ -420,7 +420,7 @@
preview_aa_samples = preview_aa_samples * preview_aa_samples;
}
- if(get_boolean(cscene, "progressive") == 0 && params.device.type == DEVICE_CPU) {
+ if(get_boolean(cscene, "progressive") == 0) {
if(background) {
params.samples = aa_samples;
}
Modified: trunk/blender/intern/cycles/device/device_cuda.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_cuda.cpp 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/device/device_cuda.cpp 2013-08-09 18:47:25 UTC (rev 59034)
@@ -558,7 +558,7 @@
}
}
- void path_trace(RenderTile& rtile, int sample)
+ void path_trace(RenderTile& rtile, int sample, bool progressive)
{
if(have_error())
return;
@@ -570,7 +570,10 @@
CUdeviceptr d_rng_state = cuda_device_ptr(rtile.rng_state);
/* get kernel function */
- cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"))
+ if(progressive)
+ cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_progressive"))
+ else
+ cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_non_progressive"))
/* pass in parameters */
int offset = 0;
@@ -914,6 +917,8 @@
if(task->type == DeviceTask::PATH_TRACE) {
RenderTile tile;
+ bool progressive = task->integrator_progressive;
+
/* keep rendering tiles until done */
while(task->acquire_tile(this, tile)) {
int start_sample = tile.start_sample;
@@ -925,7 +930,7 @@
break;
}
- path_trace(tile, sample);
+ path_trace(tile, sample, progressive);
tile.sample = sample + 1;
Modified: trunk/blender/intern/cycles/device/device_task.h
===================================================================
--- trunk/blender/intern/cycles/device/device_task.h 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/device/device_task.h 2013-08-09 18:47:25 UTC (rev 59034)
@@ -65,6 +65,7 @@
boost::function<bool(void)> get_cancel;
bool need_finish_queue;
+ bool integrator_progressive;
protected:
double last_update_time;
};
Modified: trunk/blender/intern/cycles/kernel/kernel.cl
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cl 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/kernel/kernel.cl 2013-08-09 18:47:25 UTC (rev 59034)
@@ -51,7 +51,7 @@
int y = sy + get_global_id(1);
if(x < sx + sw && y < sy + sh)
- kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
+ kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
}
__kernel void kernel_ocl_tonemap(
Modified: trunk/blender/intern/cycles/kernel/kernel.cpp
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cpp 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/kernel/kernel.cpp 2013-08-09 18:47:25 UTC (rev 59034)
@@ -90,7 +90,10 @@
void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
- kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
+ if(kernel_data.integrator.progressive)
+ kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
+ else
+ kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */
Modified: trunk/blender/intern/cycles/kernel/kernel.cu
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cu 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/kernel/kernel.cu 2013-08-09 18:47:25 UTC (rev 59034)
@@ -26,15 +26,24 @@
#include "kernel_path.h"
#include "kernel_displace.h"
-extern "C" __global__ void kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
+extern "C" __global__ void kernel_cuda_path_trace_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
if(x < sx + sw && y < sy + sh)
- kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
+ kernel_path_trace_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
}
+extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
+{
+ int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
+ int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
+
+ if(x < sx + sw && y < sy + sh)
+ kernel_path_trace_non_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
+}
+
extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
Modified: trunk/blender/intern/cycles/kernel/kernel_path.h
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_path.h 2013-08-09 17:30:21 UTC (rev 59033)
+++ trunk/blender/intern/cycles/kernel/kernel_path.h 2013-08-09 18:47:25 UTC (rev 59034)
@@ -1134,20 +1134,8 @@
#endif
-__device void kernel_path_trace(KernelGlobals *kg,
- __global float *buffer, __global uint *rng_state,
- int sample, int x, int y, int offset, int stride)
+__device_inline void kernel_path_trace_setup(KernelGlobals *kg, __global uint *rng_state, int sample, int x, int y, RNG *rng, Ray *ray)
{
- /* buffer offset */
- int index = offset + x + y*stride;
- int pass_stride = kernel_data.film.pass_stride;
-
- rng_state += index;
- buffer += index*pass_stride;
-
- /* initialize random numbers */
- RNG rng;
-
float filter_u;
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list