diff options
author | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2012-01-25 21:23:52 +0400 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2012-01-25 21:23:52 +0400 |
commit | f99343d3b8676543e2bd6acd6ee2274c21b1b388 (patch) | |
tree | fd40cd33691a783f82cf877e774d1b3a64d91ed3 /intern | |
parent | 14f475fccad7158098ddecc285c617f990b2f8b5 (diff) |
Cycles: Render Passes
Currently supported passes:
* Combined, Z, Normal, Object Index, Material Index, Emission, Environment,
Diffuse/Glossy/Transmission x Direct/Indirect/Color
Not supported yet:
* UV, Vector, Mist
Only enabled for CPU devices at the moment, will do GPU tweaks tommorrow,
also for environment importance sampling.
Documentation:
http://wiki.blender.org/index.php/Doc:2.6/Manual/Render/Cycles/Passes
Diffstat (limited to 'intern')
29 files changed, 1225 insertions, 187 deletions
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py index 763cff0bd7d..6353bf37a15 100644 --- a/intern/cycles/blender/addon/ui.py +++ b/intern/cycles/blender/addon/ui.py @@ -183,6 +183,38 @@ class CyclesRender_PT_layers(CyclesButtonsPanel, Panel): layout.separator() + split = layout.split() + + col = split.column() + col.label(text="Passes:") + col.prop(rl, "use_pass_combined") + col.prop(rl, "use_pass_z") + col.prop(rl, "use_pass_normal") + col.prop(rl, "use_pass_object_index") + col.prop(rl, "use_pass_material_index") + col.prop(rl, "use_pass_emit") + col.prop(rl, "use_pass_environment") + + col = split.column() + col.label() + col.label(text="Diffuse:") + row = col.row(align=True) + row.prop(rl, "use_pass_diffuse_direct", text="Direct", toggle=True) + row.prop(rl, "use_pass_diffuse_indirect", text="Indirect", toggle=True) + row.prop(rl, "use_pass_diffuse_color", text="Color", toggle=True) + col.label(text="Glossy:") + row = col.row(align=True) + row.prop(rl, "use_pass_glossy_direct", text="Direct", toggle=True) + row.prop(rl, "use_pass_glossy_indirect", text="Indirect", toggle=True) + row.prop(rl, "use_pass_glossy_color", text="Color", toggle=True) + col.label(text="Transmission:") + row = col.row(align=True) + row.prop(rl, "use_pass_transmission_direct", text="Direct", toggle=True) + row.prop(rl, "use_pass_transmission_indirect", text="Indirect", toggle=True) + row.prop(rl, "use_pass_transmission_color", text="Color", toggle=True) + + layout.separator() + rl = rd.layers[0] layout.prop(rl, "material_override", text="Material") diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp index c805bd03063..afcfcd95063 100644 --- a/intern/cycles/blender/blender_object.cpp +++ b/intern/cycles/blender/blender_object.cpp @@ -214,6 +214,7 @@ void BlenderSync::sync_object(BL::Object b_parent, int b_index, BL::Object b_ob, /* object sync */ if(object_updated || (object->mesh && object->mesh->need_update)) { object->name = b_ob.name().c_str(); + object->pass_id = b_ob.pass_index(); object->tfm = tfm; /* visibility flags for both parent */ diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp index b052fb3e195..ff1c32831bb 100644 --- a/intern/cycles/blender/blender_session.cpp +++ b/intern/cycles/blender/blender_session.cpp @@ -116,9 +116,68 @@ void BlenderSession::free_session() delete session; } +static PassType get_pass_type(BL::RenderPass b_pass) +{ + switch(b_pass.type()) { + case BL::RenderPass::type_COMBINED: + return PASS_COMBINED; + + case BL::RenderPass::type_Z: + return PASS_DEPTH; + case BL::RenderPass::type_NORMAL: + return PASS_NORMAL; + case BL::RenderPass::type_OBJECT_INDEX: + return PASS_OBJECT_ID; + case BL::RenderPass::type_UV: + return PASS_UV; + case BL::RenderPass::type_MATERIAL_INDEX: + return PASS_MATERIAL_ID; + + case BL::RenderPass::type_DIFFUSE_DIRECT: + return PASS_DIFFUSE_DIRECT; + case BL::RenderPass::type_GLOSSY_DIRECT: + return PASS_GLOSSY_DIRECT; + case BL::RenderPass::type_TRANSMISSION_DIRECT: + return PASS_TRANSMISSION_DIRECT; + + case BL::RenderPass::type_DIFFUSE_INDIRECT: + return PASS_DIFFUSE_INDIRECT; + case BL::RenderPass::type_GLOSSY_INDIRECT: + return PASS_GLOSSY_INDIRECT; + case BL::RenderPass::type_TRANSMISSION_INDIRECT: + return PASS_TRANSMISSION_INDIRECT; + + case BL::RenderPass::type_DIFFUSE_COLOR: + return PASS_DIFFUSE_COLOR; + case BL::RenderPass::type_GLOSSY_COLOR: + return PASS_GLOSSY_COLOR; + case BL::RenderPass::type_TRANSMISSION_COLOR: + return PASS_TRANSMISSION_COLOR; + + case BL::RenderPass::type_EMIT: + return PASS_EMISSION; + case BL::RenderPass::type_ENVIRONMENT: + return PASS_BACKGROUND; + + case BL::RenderPass::type_DIFFUSE: + case BL::RenderPass::type_SHADOW: + case BL::RenderPass::type_AO: + case BL::RenderPass::type_COLOR: + case BL::RenderPass::type_REFRACTION: + case BL::RenderPass::type_SPECULAR: + case BL::RenderPass::type_REFLECTION: + case BL::RenderPass::type_VECTOR: + case BL::RenderPass::type_MIST: + return PASS_NONE; + } + + return PASS_NONE; +} + void BlenderSession::render() { /* get buffer parameters */ + SessionParams session_params = BlenderSync::get_session_params(b_userpref, b_scene, background); BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height); int w = buffer_params.width, h = buffer_params.height; @@ -143,6 +202,25 @@ void BlenderSession::render() /* set layer */ b_rlay = *b_iter; + /* add passes */ + if(session_params.device.type == DEVICE_CPU) { /* todo */ + BL::RenderLayer::passes_iterator b_pass_iter; + + for(b_rlay.passes.begin(b_pass_iter); b_pass_iter != b_rlay.passes.end(); ++b_pass_iter) { + BL::RenderPass b_pass(*b_pass_iter); + PassType pass_type = get_pass_type(b_pass); + + if(pass_type != PASS_NONE) + Pass::add(pass_type, buffer_params.passes); + } + } + + scene->film->passes = buffer_params.passes; + scene->film->need_update = true; + + /* update session */ + session->reset(buffer_params, session_params.samples); + /* update scene */ sync->sync_data(b_v3d, active); @@ -165,22 +243,41 @@ void BlenderSession::write_render_result() { /* get state */ RenderBuffers *buffers = session->buffers; + + /* copy data from device */ + if(!buffers->copy_from_device()) + return; + + BufferParams& params = buffers->params; float exposure = scene->film->exposure; double total_time, sample_time; int sample; + session->progress.get_sample(sample, total_time, sample_time); - /* get pixels */ - float4 *pixels = buffers->copy_from_device(exposure, sample); + vector<float> pixels(params.width*params.height*4); - if(!pixels) - return; + /* copy each pass */ + BL::RenderLayer::passes_iterator b_iter; + + for(b_rlay.passes.begin(b_iter); b_iter != b_rlay.passes.end(); ++b_iter) { + BL::RenderPass b_pass(*b_iter); - /* write pixels */ - rna_RenderLayer_rect_set(&b_rlay.ptr, (float*)pixels); - RE_engine_update_result((RenderEngine*)b_engine.ptr.data, (RenderResult*)b_rr.ptr.data); + /* find matching pass type */ + PassType pass_type = get_pass_type(b_pass); + int components = b_pass.channels(); + + /* copy pixels */ + if(buffers->get_pass(pass_type, exposure, sample, components, &pixels[0])) + rna_RenderPass_rect_set(&b_pass.ptr, &pixels[0]); + } - delete [] pixels; + /* copy combined pass */ + if(buffers->get_pass(PASS_COMBINED, exposure, sample, 4, &pixels[0])) + rna_RenderLayer_rect_set(&b_rlay.ptr, &pixels[0]); + + /* tag result as updated */ + RE_engine_update_result((RenderEngine*)b_engine.ptr.data, (RenderResult*)b_rr.ptr.data); } void BlenderSession::synchronize() diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp index f7b750e6dab..5310e35dc25 100644 --- a/intern/cycles/blender/blender_shader.cpp +++ b/intern/cycles/blender/blender_shader.cpp @@ -636,6 +636,7 @@ void BlenderSync::sync_materials() ShaderGraph *graph = new ShaderGraph(); shader->name = b_mat->name().c_str(); + shader->pass_id = b_mat->pass_index(); /* create nodes */ if(b_mat->use_nodes() && b_mat->node_tree()) { diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index 25da978edd8..2ca599f6c67 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -162,7 +162,7 @@ public: if(system_cpu_support_optimized()) { for(int y = task.y; y < task.y + task.h; y++) { for(int x = task.x; x < task.x + task.w; x++) - kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, + kernel_cpu_optimized_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y, task.offset, task.stride); if(tasks.worker_cancel()) @@ -174,7 +174,7 @@ public: { for(int y = task.y; y < task.y + task.h; y++) { for(int x = task.x; x < task.x + task.w; x++) - kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, + kernel_cpu_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y, task.offset, task.stride); if(tasks.worker_cancel()) @@ -194,7 +194,7 @@ public: if(system_cpu_support_optimized()) { for(int y = task.y; y < task.y + task.h; y++) for(int x = task.x; x < task.x + task.w; x++) - kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, + kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer, task.sample, task.resolution, x, y, task.offset, task.stride); } else @@ -202,7 +202,7 @@ public: { for(int y = task.y; y < task.y + task.h; y++) for(int x = task.x; x < task.x + task.w; x++) - kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, + kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer, task.sample, task.resolution, x, y, task.offset, task.stride); } } diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index b6e749f3fcb..59fe9709e61 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -15,6 +15,7 @@ set(SRC set(SRC_HEADERS kernel.h + kernel_accumulate.h kernel_bvh.h kernel_camera.h kernel_compat_cpu.h @@ -30,6 +31,7 @@ set(SRC_HEADERS kernel_mbvh.h kernel_montecarlo.h kernel_object.h + kernel_passes.h kernel_path.h kernel_qbvh.h kernel_random.h diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 305d81339c2..f98414d4f02 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -28,7 +28,7 @@ __kernel void kernel_ocl_path_trace( __constant KernelData *data, - __global float4 *buffer, + __global float *buffer, __global uint *rng_state, #define KERNEL_TEX(type, ttype, name) \ @@ -56,7 +56,7 @@ __kernel void kernel_ocl_path_trace( __kernel void kernel_ocl_tonemap( __constant KernelData *data, __global uchar4 *rgba, - __global float4 *buffer, + __global float *buffer, #define KERNEL_TEX(type, ttype, name) \ __global type *name, diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index 9e0d252772b..a93f6172d28 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -204,14 +204,14 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t /* Path Tracing */ -void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) +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); } /* Tonemapping */ -void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride) { kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride); } diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index 4e585fd563d..fc7992c87f0 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -26,7 +26,7 @@ #include "kernel_path.h" #include "kernel_displace.h" -extern "C" __global__ void kernel_cuda_path_trace(float4 *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(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; @@ -35,7 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_stat kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride); } -extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride) +extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int resolution, 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; diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h index df6b5ee92da..26c0bcd6d1a 100644 --- a/intern/cycles/kernel/kernel.h +++ b/intern/cycles/kernel/kernel.h @@ -36,17 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg); void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size); void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height); -void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, +void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride); -void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, +void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride); void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); #ifdef WITH_OPTIMIZED_KERNEL -void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, +void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride); -void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, +void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride); void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h new file mode 100644 index 00000000000..e71fad58c96 --- /dev/null +++ b/intern/cycles/kernel/kernel_accumulate.h @@ -0,0 +1,283 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + */ + +CCL_NAMESPACE_BEGIN + +/* BSDF Eval + * + * BSDF evaluation result, split per BSDF type. This is used to accumulate + * render passes separately. */ + +__device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass) +{ +#ifdef __PASSES__ + eval->use_light_pass = use_light_pass; + + if(eval->use_light_pass) { + eval->diffuse = make_float3(0.0f, 0.0f, 0.0f); + eval->glossy = make_float3(0.0f, 0.0f, 0.0f); + eval->transmission = make_float3(0.0f, 0.0f, 0.0f); + eval->transparent = make_float3(0.0f, 0.0f, 0.0f); + + if(type == CLOSURE_BSDF_TRANSPARENT_ID) + eval->transparent = value; + else if(CLOSURE_IS_BSDF_DIFFUSE(type)) + eval->diffuse = value; + else if(CLOSURE_IS_BSDF_GLOSSY(type)) + eval->glossy = value; + else + eval->transmission = value; + } + else + eval->diffuse = value; +#else + *eval = value; +#endif +} + +__device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + if(CLOSURE_IS_BSDF_DIFFUSE(type)) + eval->diffuse += value; + else if(CLOSURE_IS_BSDF_GLOSSY(type)) + eval->glossy += value; + else + eval->transmission += value; + + /* skipping transparent, this function is used by for eval(), will be zero then */ + } + else + eval->diffuse += value; +#else + *eval += value; +#endif +} + +__device_inline bool bsdf_eval_is_zero(BsdfEval *eval) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + return is_zero(eval->diffuse) + && is_zero(eval->glossy) + && is_zero(eval->transmission) + && is_zero(eval->transparent); + } + else + return is_zero(eval->diffuse); +#else + return is_zero(*eval); +#endif +} + +__device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + eval->diffuse *= value; + eval->glossy *= value; + eval->transmission *= value; + + /* skipping transparent, this function is used by for eval(), will be zero then */ + } + else + eval->diffuse *= value; +#else + *eval *= value; +#endif +} + +/* Path Radiance + * + * We accumulate different render passes separately. After summing at the end + * to get the combined result, it should be identical. We definte directly + * visible as the first non-transparent hit, while indirectly visible are the + * bounces after that. */ + +__device_inline void path_radiance_init(PathRadiance *L, int use_light_pass) +{ + /* clear all */ +#ifdef __PASSES__ + L->use_light_pass = use_light_pass; + + if(use_light_pass) { + L->indirect = make_float3(0.0f, 0.0f, 0.0f); + L->direct_throughput = make_float3(0.0f, 0.0f, 0.0f); + L->direct_emission = make_float3(0.0f, 0.0f, 0.0f); + + L->color_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->color_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->color_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->direct_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->direct_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->direct_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->indirect_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->indirect_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->indirect_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->emission = make_float3(0.0f, 0.0f, 0.0f); + L->background = make_float3(0.0f, 0.0f, 0.0f); + } + else + L->emission = make_float3(0.0f, 0.0f, 0.0f); +#else + *L = make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput, + BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label) +{ + float inverse_pdf = 1.0f/bsdf_pdf; + +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) { + if(bsdf_label & LABEL_TRANSPARENT) { + /* transparent bounce before first hit */ + *throughput *= bsdf_eval->transparent*inverse_pdf; + } + else { + /* first on directly visible surface */ + float3 value = *throughput*inverse_pdf; + + L->indirect_diffuse = bsdf_eval->diffuse*value; + L->indirect_glossy = bsdf_eval->glossy*value; + L->indirect_transmission = bsdf_eval->transmission*value; + + *throughput = L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission; + + L->direct_throughput = *throughput; + } + } + else { + /* indirectly visible through BSDF */ + float3 sum = (bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission + bsdf_eval->transparent)*inverse_pdf; + *throughput *= sum; + } + } + else + *throughput *= bsdf_eval->diffuse*inverse_pdf; +#else + *throughput *= *bsdf_eval*inverse_pdf; +#endif +} + +__device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) + L->emission += throughput*value; + else if(bounce == 1) + L->direct_emission += throughput*value; + else + L->indirect += throughput*value; + } + else + L->emission += throughput*value; +#else + *L += throughput*value; +#endif +} + +__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) { + /* directly visible lighting */ + L->direct_diffuse += throughput*bsdf_eval->diffuse; + L->direct_glossy += throughput*bsdf_eval->glossy; + L->direct_transmission += throughput*bsdf_eval->transmission; + } + else { + /* indirectly visible lighting after BSDF bounce */ + float3 sum = bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission; + L->indirect += throughput*sum; + } + } + else + L->emission += throughput*bsdf_eval->diffuse; +#else + *L += throughput*(*bsdf_eval); +#endif +} + +__device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) + L->background += throughput*value; + else if(bounce == 1) + L->direct_emission += throughput*value; + else + L->indirect += throughput*value; + } + else + L->emission += throughput*value; +#else + *L += throughput*value; +#endif +} + +__device_inline float3 safe_divide_color(float3 a, float3 b) +{ + float x, y, z; + + x = (b.x != 0.0f)? a.x/b.x: 0.0f; + y = (b.y != 0.0f)? a.y/b.y: 0.0f; + z = (b.z != 0.0f)? a.z/b.z: 0.0f; + + return make_float3(x, y, z); +} + +__device_inline float3 path_radiance_sum(PathRadiance *L) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + /* this division is a bit ugly, but means we only have to keep track of + only a single throughput further along the path, here we recover just + the indirect parth that is not influenced by any particular BSDF type */ + L->direct_emission = safe_divide_color(L->direct_emission, L->direct_throughput); + L->direct_diffuse += L->indirect_diffuse*L->direct_emission; + L->direct_glossy += L->indirect_glossy*L->direct_emission; + L->direct_transmission += L->indirect_transmission*L->direct_emission; + + L->indirect = safe_divide_color(L->indirect, L->direct_throughput); + L->indirect_diffuse *= L->indirect; + L->indirect_glossy *= L->indirect; + L->indirect_transmission *= L->indirect; + + return L->emission + L->background + + L->direct_diffuse + L->direct_glossy + L->direct_transmission + + L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission; + } + else + return L->emission; +#else + return *L; +#endif +} + +CCL_NAMESPACE_END + diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index 51698f3a9bd..b3a5b2bfcb4 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -57,7 +57,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando, } __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, - float randt, float rando, float randu, float randv, Ray *ray, float3 *eval) + float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval) { LightSample ls; @@ -83,32 +83,33 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, return false; /* evaluate closure */ - *eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D); + float3 light_eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D); - if(is_zero(*eval)) + if(is_zero(light_eval)) return false; /* todo: use visbility flag to skip lights */ /* evaluate BSDF at shading point */ float bsdf_pdf; - float3 bsdf_eval = shader_bsdf_eval(kg, sd, ls.D, &bsdf_pdf); - *eval *= bsdf_eval/pdf; - - if(is_zero(*eval)) - return false; + shader_bsdf_eval(kg, sd, ls.D, eval, &bsdf_pdf); if(ls.prim != ~0 || ls.type == LIGHT_BACKGROUND) { /* multiple importance sampling */ float mis_weight = power_heuristic(pdf, bsdf_pdf); - *eval *= mis_weight; + light_eval *= mis_weight; } /* todo: clean up these weights */ else if(ls.shader & SHADER_AREA_LIGHT) - *eval *= 0.25f; /* area lamp */ + light_eval *= 0.25f; /* area lamp */ else if(ls.t != FLT_MAX) - *eval *= 0.25f*M_1_PI_F; /* point lamp */ + light_eval *= 0.25f*M_1_PI_F; /* point lamp */ + + bsdf_eval_mul(eval, light_eval/pdf); + + if(bsdf_eval_is_zero(eval)) + return false; if(ls.shader & SHADER_CAST_SHADOW) { /* setup ray */ diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h index cd8acc9647a..d0fb5402291 100644 --- a/intern/cycles/kernel/kernel_film.h +++ b/intern/cycles/kernel/kernel_film.h @@ -48,15 +48,22 @@ __device uchar4 film_float_to_byte(float4 color) return result; } -__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +__device void kernel_film_tonemap(KernelGlobals *kg, + __global uchar4 *rgba, __global float *buffer, + int sample, int resolution, int x, int y, int offset, int stride) { + /* buffer offset */ int index = offset + x + y*stride; - float4 irradiance = buffer[index]; + rgba += index; + buffer += index*kernel_data.film.pass_stride; + + /* map colors */ + float4 irradiance = *(float4*)buffer; float4 float_result = film_map(kg, irradiance, sample); uchar4 byte_result = film_float_to_byte(float_result); - rgba[index] = byte_result; + *rgba = byte_result; } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_object.h b/intern/cycles/kernel/kernel_object.h index b4c42d3bacc..318a6fea489 100644 --- a/intern/cycles/kernel/kernel_object.h +++ b/intern/cycles/kernel/kernel_object.h @@ -64,5 +64,15 @@ __device_inline float object_surface_area(KernelGlobals *kg, int object) return f.x; } +__device_inline float object_pass_id(KernelGlobals *kg, int object) +{ + if(object == ~0) + return 0.0f; + + int offset = object*OBJECT_SIZE + OBJECT_PROPERTIES; + float4 f = kernel_tex_fetch(__objects, offset); + return f.y; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp index 50341021d9d..393686bb203 100644 --- a/intern/cycles/kernel/kernel_optimized.cpp +++ b/intern/cycles/kernel/kernel_optimized.cpp @@ -35,14 +35,14 @@ CCL_NAMESPACE_BEGIN /* Path Tracing */ -void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) +void kernel_cpu_optimized_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); } /* Tonemapping */ -void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride) { kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride); } diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h new file mode 100644 index 00000000000..0e775812eda --- /dev/null +++ b/intern/cycles/kernel/kernel_passes.h @@ -0,0 +1,146 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + */ + +CCL_NAMESPACE_BEGIN + +__device_inline void kernel_write_pass_float(__global float *buffer, int sample, float value) +{ + float *buf = buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_write_pass_float3(__global float *buffer, int sample, float3 value) +{ + float3 *buf = (float3*)buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_write_pass_float4(__global float *buffer, int sample, float4 value) +{ + float4 *buf = (float4*)buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_clear_passes(__global float *buffer, int sample, int pass_stride) +{ +#ifdef __PASSES__ + if(sample == 0 && pass_stride != 4) + for(int i = 4; i < pass_stride; i++) + buffer[i] = 0.0f; +#endif +} + +__device void kernel_write_data_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, + ShaderData *sd, int sample, int path_flag, float3 throughput) +{ +#ifdef __PASSES__ + if(!(path_flag & PATH_RAY_CAMERA)) + return; + + int flag = kernel_data.film.pass_flag; + + if(!(flag & PASS_ALL)) + return; + + /* todo: add alpha treshold */ + if(!(path_flag & PATH_RAY_TRANSPARENT)) { + if(sample == 0) { + if(flag & PASS_DEPTH) { + Transform tfm = kernel_data.cam.worldtocamera; + float depth = len(transform(&tfm, sd->P)); + + kernel_write_pass_float(buffer + kernel_data.film.pass_depth, sample, depth); + } + if(flag & PASS_OBJECT_ID) { + float id = object_pass_id(kg, sd->object); + kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, sample, id); + } + if(flag & PASS_MATERIAL_ID) { + float id = shader_pass_id(kg, sd); + kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, sample, id); + } + } + + if(flag & PASS_NORMAL) { + float3 normal = sd->N; + kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, sample, normal); + } + if(flag & PASS_UV) { + float3 uv = make_float3(0.0f, 0.0f, 0.0f); /* todo: request and lookup */ + kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, sample, uv); + } + } + + if(flag & (PASS_DIFFUSE_INDIRECT|PASS_DIFFUSE_COLOR|PASS_DIFFUSE_DIRECT)) + L->color_diffuse += shader_bsdf_diffuse(kg, sd)*throughput; + if(flag & (PASS_GLOSSY_INDIRECT|PASS_GLOSSY_COLOR|PASS_GLOSSY_DIRECT)) + L->color_glossy += shader_bsdf_glossy(kg, sd)*throughput; + if(flag & (PASS_TRANSMISSION_INDIRECT|PASS_TRANSMISSION_COLOR|PASS_TRANSMISSION_DIRECT)) + L->color_transmission += shader_bsdf_transmission(kg, sd)*throughput; +#endif +} + +__device void kernel_write_light_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, int sample) +{ +#ifdef __PASSES__ + int flag = kernel_data.film.pass_flag; + + if(!kernel_data.film.use_light_pass) + return; + + if(flag & PASS_DIFFUSE_INDIRECT) { + float3 color = safe_divide_color(L->indirect_diffuse, L->color_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, sample, color); + } + if(flag & PASS_GLOSSY_INDIRECT) { + float3 color = safe_divide_color(L->indirect_glossy, L->color_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, sample, color); + } + if(flag & PASS_TRANSMISSION_INDIRECT) { + float3 color = safe_divide_color(L->indirect_transmission, L->color_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, sample, color); + } + if(flag & PASS_DIFFUSE_DIRECT) { + float3 color = safe_divide_color(L->direct_diffuse, L->color_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, sample, color); + } + if(flag & PASS_GLOSSY_DIRECT) { + float3 color = safe_divide_color(L->direct_glossy, L->color_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, sample, color); + } + if(flag & PASS_TRANSMISSION_DIRECT) { + float3 color = safe_divide_color(L->direct_transmission, L->color_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, sample, color); + } + + if(flag & PASS_EMISSION) + kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, sample, L->emission); + if(flag & PASS_BACKGROUND) + kernel_write_pass_float3(buffer + kernel_data.film.pass_background, sample, L->background); + + if(flag & PASS_DIFFUSE_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, sample, L->color_diffuse); + if(flag & PASS_GLOSSY_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, sample, L->color_glossy); + if(flag & PASS_TRANSMISSION_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, sample, L->color_transmission); +#endif +} + +CCL_NAMESPACE_END + diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index c80d2068506..c0bfa320405 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -25,39 +25,16 @@ #else #include "kernel_bvh.h" #endif +#include "kernel_accumulate.h" #include "kernel_camera.h" #include "kernel_shader.h" #include "kernel_light.h" #include "kernel_emission.h" #include "kernel_random.h" +#include "kernel_passes.h" CCL_NAMESPACE_BEGIN -#ifdef __MODIFY_TP__ -__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample) -{ - /* modify throughput to influence path termination probability, to avoid - darker regions receiving fewer samples than lighter regions. also RGB - are weighted differently. proper validation still remains to be done. */ - const float3 weights = make_float3(1.0f, 1.33f, 0.66f); - const float3 one = make_float3(1.0f, 1.0f, 1.0f); - const int minsample = 5; - const float minL = 0.1f; - - if(sample >= minsample) { - float3 L = buffer[offset + x + y*stride]; - float3 Lmin = make_float3(minL, minL, minL); - float correct = (float)(sample+1)/(float)sample; - - L = film_map(L*correct, sample); - - return weights/clamp(L, Lmin, one); - } - - return weights; -} -#endif - typedef struct PathState { uint flag; int bounce; @@ -168,7 +145,7 @@ __device_inline float path_state_terminate_probability(KernelGlobals *kg, PathSt return average(throughput); } -__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, float3 *light_L) +__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, BsdfEval *L_light) { if(ray->t == 0.0f) return false; @@ -208,7 +185,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra } if(!scene_intersect(kg, ray, PATH_RAY_SHADOW_TRANSPARENT, isect)) { - *light_L *= throughput; + bsdf_eval_mul(L_light, throughput); return false; } @@ -234,11 +211,14 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra return result; } -__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, float3 throughput) +__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer) { /* initialize */ - float3 L = make_float3(0.0f, 0.0f, 0.0f); - float Ltransparent = 0.0f; + PathRadiance L; + float3 throughput = make_float3(1.0f, 1.0f, 1.0f); + float L_transparent = 0.0f; + + path_radiance_init(&L, kernel_data.film.use_light_pass); #ifdef __EMISSION__ float ray_pdf = 0.0f; @@ -257,12 +237,12 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R if(!scene_intersect(kg, &ray, visibility, &isect)) { /* eval background shader if nothing hit */ if(kernel_data.background.transparent && (state.flag & PATH_RAY_CAMERA)) { - Ltransparent += average(throughput); + L_transparent += average(throughput); } else { /* sample background shader */ - float3 background_L = indirect_background(kg, &ray, state.flag, ray_pdf); - L += throughput*background_L; + float3 L_background = indirect_background(kg, &ray, state.flag, ray_pdf); + path_radiance_accum_background(&L, throughput, L_background, state.bounce); } break; @@ -274,19 +254,24 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R float rbsdf = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF); shader_eval_surface(kg, &sd, rbsdf, state.flag); + kernel_write_data_passes(kg, buffer, &L, &sd, sample, state.flag, throughput); + #ifdef __HOLDOUT__ if((sd.flag & SD_HOLDOUT) && (state.flag & PATH_RAY_CAMERA)) { float3 holdout_weight = shader_holdout_eval(kg, &sd); if(kernel_data.background.transparent) - Ltransparent += average(holdout_weight*throughput); + /* any throughput is ok, should all be identical here */ + L_transparent += average(holdout_weight*throughput); } #endif #ifdef __EMISSION__ /* emission */ - if(sd.flag & SD_EMISSION) - L += throughput*indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf); + if(sd.flag & SD_EMISSION) { + float3 emission = indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf); + path_radiance_accum_emission(&L, throughput, emission, state.bounce); + } #endif /* path termination. this is a strange place to put the termination, it's @@ -310,7 +295,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R float light_v = path_rng(kg, rng, sample, rng_offset + PRNG_LIGHT_V); Ray light_ray; - float3 light_L; + BsdfEval L_light; #ifdef __MULTI_LIGHT__ /* index -1 means randomly sample from distribution */ @@ -320,10 +305,10 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R #else const int i = -1; #endif - if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &light_L)) { + if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &L_light)) { /* trace shadow ray */ - if(!shadow_blocked(kg, &state, &light_ray, &isect, &light_L)) - L += throughput*light_L; + if(!shadow_blocked(kg, &state, &light_ray, &isect, &L_light)) + path_radiance_accum_light(&L, throughput, &L_light, state.bounce); } #ifdef __MULTI_LIGHT__ } @@ -338,7 +323,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R /* sample BSDF */ float bsdf_pdf; - float3 bsdf_eval; + BsdfEval bsdf_eval; float3 bsdf_omega_in; differential3 bsdf_domega_in; float bsdf_u = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF_U); @@ -350,11 +335,11 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R shader_release(kg, &sd); - if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) + if(bsdf_pdf == 0.0f || bsdf_eval_is_zero(&bsdf_eval)) break; /* modify throughput */ - throughput *= bsdf_eval/bsdf_pdf; + path_radiance_bsdf_bounce(&L, &throughput, &bsdf_eval, bsdf_pdf, state.bounce, label); /* set labels */ #if defined(__EMISSION__) || defined(__BACKGROUND__) @@ -374,18 +359,33 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R #endif } - return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent); + float3 L_sum = path_radiance_sum(&L); + + kernel_write_light_passes(kg, buffer, &L, sample); + + return make_float4(L_sum.x, L_sum.y, L_sum.z, 1.0f - L_transparent); } -__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride) +__device void kernel_path_trace(KernelGlobals *kg, + __global float *buffer, __global uint *rng_state, + int sample, int x, int y, int offset, int stride) { + /* buffer offset */ + int index = offset + x + y*stride; + int pass_stride = kernel_data.film.pass_stride; + + rng_state += index; + buffer += index*pass_stride; + + kernel_clear_passes(buffer, sample, pass_stride); + /* initialize random numbers */ RNG rng; float filter_u; float filter_v; - path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v); + path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v); /* sample camera ray */ Ray ray; @@ -396,23 +396,12 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, &ray); /* integrate */ -#ifdef __MODIFY_TP__ - float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample); - float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput; -#else - float3 throughput = make_float3(1.0f, 1.0f, 1.0f); - float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput); -#endif + float4 L = kernel_path_integrate(kg, &rng, sample, ray, buffer); /* accumulate result in output buffer */ - int index = offset + x + y*stride; - - if(sample == 0) - buffer[index] = L; - else - buffer[index] += L; + kernel_write_pass_float4(buffer, sample, L); - path_rng_end(kg, rng_state, rng, x, y, offset, stride); + path_rng_end(kg, rng_state, rng); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h index 41301ebd3dc..6d15100f8a3 100644 --- a/intern/cycles/kernel/kernel_random.h +++ b/intern/cycles/kernel/kernel_random.h @@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime #endif } -__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy) +__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy) { #ifdef __SOBOL_FULL_SCREEN__ uint px, py; @@ -135,19 +135,31 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, *rng ^= kernel_data.integrator.seed; - *fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x; - *fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y; + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x; + *fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y; + } #else - *rng = rng_state[offset + x + y*stride]; + *rng = *rng_state; *rng ^= kernel_data.integrator.seed; - *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); - *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); + *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + } #endif } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride) +__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) { /* nothing to do */ } @@ -163,21 +175,27 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension) return (float)*rng * (1.0f/(float)0xFFFFFFFF); } -__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy) +__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy) { /* load state */ - *rng = rng_state[offset + x + y*stride]; + *rng = *rng_state; *rng ^= kernel_data.integrator.seed; - *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); - *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); + *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + } } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride) +__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) { /* store state for next sample */ - rng_state[offset + x + y*stride] = rng; + *rng_state = rng; } #endif diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index 9c59e1566a9..1d2cf46aa56 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -75,7 +75,7 @@ __device_inline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, if(sd->shader & SHADER_SMOOTH_NORMAL) sd->N = triangle_smooth_normal(kg, sd->prim, sd->u, sd->v); - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __DPDU__ /* dPdu/dPdv */ @@ -166,7 +166,7 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, #endif } - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __DPDU__ /* dPdu/dPdv */ @@ -243,7 +243,7 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData sd->Ng = -sd->P; sd->I = -sd->P; sd->shader = kernel_data.background.shader; - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __INSTANCING__ sd->object = ~0; @@ -275,8 +275,8 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData #ifdef __MULTI_CLOSURE__ -__device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf, - int skip_bsdf, float3 sum_eval, float sum_pdf, float sum_sample_weight) +__device_inline void _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf, + int skip_bsdf, BsdfEval *bsdf_eval, float sum_pdf, float sum_sample_weight) { for(int i = 0; i< sd->num_closure; i++) { if(i == skip_bsdf) @@ -293,7 +293,7 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float #endif if(bsdf_pdf != 0.0f) { - sum_eval += eval*sc->weight; + bsdf_eval_accum(bsdf_eval, sc->type, eval*sc->weight); sum_pdf += bsdf_pdf*sc->sample_weight; } @@ -302,25 +302,27 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float } *pdf = sum_pdf/sum_sample_weight; - return sum_eval; } #endif -__device float3 shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, - const float3 omega_in, float *pdf) +__device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, + const float3 omega_in, BsdfEval *eval, float *pdf) { #ifdef __MULTI_CLOSURE__ - return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, make_float3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f); + bsdf_eval_init(eval, NBUILTIN_CLOSURES, make_float3(0.0f, 0.0f, 0.0f), kernel_data.film.use_light_pass); + + return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f); #else const ShaderClosure *sc = &sd->closure; + *pdf = 0.0f; - return svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight; + *eval = svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight; #endif } __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, - float randu, float randv, float3 *eval, + float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf) { #ifdef __MULTI_CLOSURE__ @@ -359,27 +361,28 @@ __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc = &sd->closure[sampled]; int label; + float3 eval; *pdf = 0.0f; #ifdef __OSL__ - label = OSLShader::bsdf_sample(sd, sc, randu, randv, *eval, *omega_in, *domega_in, *pdf); + label = OSLShader::bsdf_sample(sd, sc, randu, randv, eval, *omega_in, *domega_in, *pdf); #else - label = svm_bsdf_sample(sd, sc, randu, randv, eval, omega_in, domega_in, pdf); + label = svm_bsdf_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf); #endif - *eval *= sc->weight; + bsdf_eval_init(bsdf_eval, sc->type, eval*sc->weight, kernel_data.film.use_light_pass); if(sd->num_closure > 1 && *pdf != 0.0f) { float sweight = sc->sample_weight; - *eval = _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, *eval, *pdf*sweight, sweight); + _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, bsdf_eval, *pdf*sweight, sweight); } return label; #else /* sample the single closure that we picked */ *pdf = 0.0f; - int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, eval, omega_in, domega_in, pdf); - *eval *= sd->closure.weight; + int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, bsdf_eval, omega_in, domega_in, pdf); + *bsdf_eval *= sd->closure.weight; return label; #endif } @@ -421,6 +424,68 @@ __device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd) #endif } +__device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_DIFFUSE(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_DIFFUSE(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_GLOSSY(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_GLOSSY(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_TRANSMISSION(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_TRANSMISSION(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} /* Emission */ @@ -588,12 +653,17 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect) int prim = kernel_tex_fetch(__prim_index, isect->prim); float4 Ns = kernel_tex_fetch(__tri_normal, prim); int shader = __float_as_int(Ns.w); - int flag = kernel_tex_fetch(__shader_flag, shader & SHADER_MASK); + int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)*2); return (flag & SD_HAS_SURFACE_TRANSPARENT) != 0; } #endif +__device int shader_pass_id(KernelGlobals *kg, ShaderData *sd) +{ + return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2 + 1); +} + /* Free ShaderData */ __device void shader_release(KernelGlobals *kg, ShaderData *sd) diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index 008ec0bdf28..b4b1da83162 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -70,6 +70,9 @@ CCL_NAMESPACE_BEGIN #ifdef __KERNEL_ADV_SHADING__ #define __MULTI_CLOSURE__ #define __TRANSPARENT_SHADOWS__ +#ifdef __KERNEL_CPU__ +#define __PASSES__ +#endif #endif //#define __MULTI_LIGHT__ @@ -150,6 +153,75 @@ typedef enum ClosureLabel { LABEL_STOP = 2048 } ClosureLabel; +/* Render Passes */ + +typedef enum PassType { + PASS_NONE = 0, + PASS_COMBINED = 1, + PASS_DEPTH = 2, + PASS_NORMAL = 8, + PASS_UV = 16, + PASS_OBJECT_ID = 32, + PASS_MATERIAL_ID = 64, + PASS_DIFFUSE_COLOR = 128, + PASS_GLOSSY_COLOR = 256, + PASS_TRANSMISSION_COLOR = 512, + PASS_DIFFUSE_INDIRECT = 1024, + PASS_GLOSSY_INDIRECT = 2048, + PASS_TRANSMISSION_INDIRECT = 4096, + PASS_DIFFUSE_DIRECT = 8192, + PASS_GLOSSY_DIRECT = 16384, + PASS_TRANSMISSION_DIRECT = 32768, + PASS_EMISSION = 65536, + PASS_BACKGROUND = 131072 +} PassType; + +#define PASS_ALL (~0) + +#ifdef __PASSES__ + +typedef float3 PathThroughput; + +struct PathRadiance { + int use_light_pass; + + float3 indirect; + float3 direct_throughput; + float3 direct_emission; + + float3 color_diffuse; + float3 color_glossy; + float3 color_transmission; + + float3 direct_diffuse; + float3 direct_glossy; + float3 direct_transmission; + + float3 indirect_diffuse; + float3 indirect_glossy; + float3 indirect_transmission; + + float3 emission; + float3 background; +}; + +struct BsdfEval { + int use_light_pass; + + float3 diffuse; + float3 glossy; + float3 transmission; + float3 transparent; +}; + +#else + +typedef float3 PathThroughput; +typedef float3 PathRadiance; +typedef float3 BsdfEval; + +#endif + /* Shader Flag */ typedef enum ShaderFlag { @@ -353,7 +425,32 @@ typedef struct KernelCamera { typedef struct KernelFilm { float exposure; - int pad1, pad2, pad3; + int pass_flag; + int pass_stride; + int use_light_pass; + + int pass_combined; + int pass_depth; + int pass_normal; + int pass_pad; + + int pass_uv; + int pass_object_id; + int pass_material_id; + int pass_diffuse_color; + + int pass_glossy_color; + int pass_transmission_color; + int pass_diffuse_indirect; + int pass_glossy_indirect; + + int pass_transmission_indirect; + int pass_diffuse_direct; + int pass_glossy_direct; + int pass_transmission_direct; + + int pass_emission; + int pass_background; } KernelFilm; typedef struct KernelBackground { diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h index ff6e9b94a0f..533a2944557 100644 --- a/intern/cycles/kernel/svm/svm_types.h +++ b/intern/cycles/kernel/svm/svm_types.h @@ -266,22 +266,26 @@ typedef enum ShaderType { typedef enum ClosureType { CLOSURE_BSDF_ID, + CLOSURE_BSDF_DIFFUSE_ID, CLOSURE_BSDF_OREN_NAYAR_ID, - CLOSURE_BSDF_TRANSLUCENT_ID, + CLOSURE_BSDF_REFLECTION_ID, - CLOSURE_BSDF_REFRACTION_ID, - CLOSURE_BSDF_GLASS_ID, - CLOSURE_BSDF_TRANSPARENT_ID, CLOSURE_BSDF_MICROFACET_GGX_ID, - CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID, CLOSURE_BSDF_MICROFACET_BECKMANN_ID, - CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID, CLOSURE_BSDF_WARD_ID, - CLOSURE_BSDF_ASHIKHMIN_VELVET_ID, - CLOSURE_BSDF_WESTIN_BACKSCATTER_ID, CLOSURE_BSDF_WESTIN_SHEEN_ID, + CLOSURE_BSDF_TRANSLUCENT_ID, + CLOSURE_BSDF_REFRACTION_ID, + CLOSURE_BSDF_WESTIN_BACKSCATTER_ID, + CLOSURE_BSDF_ASHIKHMIN_VELVET_ID, + CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID, + CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID, + CLOSURE_BSDF_GLASS_ID, + + CLOSURE_BSDF_TRANSPARENT_ID, + CLOSURE_BSSRDF_CUBIC_ID, CLOSURE_EMISSION_ID, CLOSURE_DEBUG_ID, @@ -297,7 +301,10 @@ typedef enum ClosureType { } ClosureType; /* watch this, being lazy with memory usage */ -#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_WESTIN_SHEEN_ID) +#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_TRANSPARENT_ID) +#define CLOSURE_IS_BSDF_DIFFUSE(type) (type >= CLOSURE_BSDF_DIFFUSE_ID && type <= CLOSURE_BSDF_OREN_NAYAR_ID) +#define CLOSURE_IS_BSDF_GLOSSY(type) (type >= CLOSURE_BSDF_REFLECTION_ID && type <= CLOSURE_BSDF_WESTIN_SHEEN_ID) +#define CLOSURE_IS_BSDF_TRANSMISSION(type) (type >= CLOSURE_BSDF_TRANSLUCENT_ID && type <= CLOSURE_BSDF_GLASS_ID) #define CLOSURE_IS_VOLUME(type) (type >= CLOSURE_VOLUME_ID && type <= CLOSURE_VOLUME_ISOTROPIC_ID) #define CLOSURE_IS_EMISSION(type) (type == CLOSURE_EMISSION_ID) #define CLOSURE_IS_HOLDOUT(type) (type == CLOSURE_HOLDOUT_ID) diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp index a6bbbc91901..08dda944111 100644 --- a/intern/cycles/render/buffers.cpp +++ b/intern/cycles/render/buffers.cpp @@ -22,6 +22,7 @@ #include "device.h" #include "util_debug.h" +#include "util_foreach.h" #include "util_hash.h" #include "util_image.h" #include "util_math.h" @@ -31,6 +32,48 @@ CCL_NAMESPACE_BEGIN +/* Buffer Params */ + +BufferParams::BufferParams() +{ + width = 0; + height = 0; + + full_x = 0; + full_y = 0; + full_width = 0; + full_height = 0; + + Pass::add(PASS_COMBINED, passes); +} + +void BufferParams::get_offset_stride(int& offset, int& stride) +{ + offset = -(full_x + full_y*width); + stride = width; +} + +bool BufferParams::modified(const BufferParams& params) +{ + return !(full_x == params.full_x + && full_y == params.full_y + && width == params.width + && height == params.height + && full_width == params.full_width + && full_height == params.full_height + && Pass::equals(passes, params.passes)); +} + +int BufferParams::get_passes_size() +{ + int size = 0; + + foreach(Pass& pass, passes) + size += pass.components; + + return size; +} + /* Render Buffers */ RenderBuffers::RenderBuffers(Device *device_) @@ -64,7 +107,7 @@ void RenderBuffers::reset(Device *device, BufferParams& params_) device_free(); /* allocate buffer */ - buffer.resize(params.width, params.height); + buffer.resize(params.width*params.height*params.get_passes_size()); device->mem_alloc(buffer, MEM_READ_WRITE); device->mem_zero(buffer); @@ -82,31 +125,76 @@ void RenderBuffers::reset(Device *device, BufferParams& params_) device->mem_copy_to(rng_state); } -float4 *RenderBuffers::copy_from_device(float exposure, int sample) +bool RenderBuffers::copy_from_device() { if(!buffer.device_pointer) - return NULL; + return false; device->mem_copy_from(buffer, 0, params.width, params.height, sizeof(float4)); - float4 *out = new float4[params.width*params.height]; - float4 *in = (float4*)buffer.data_pointer; - float scale = 1.0f/(float)sample; - - for(int i = params.width*params.height - 1; i >= 0; i--) { - float4 rgba = in[i]*scale; + return true; +} + +bool RenderBuffers::get_pass(PassType type, float exposure, int sample, int components, float *pixels) +{ + int pass_offset = 0; + + foreach(Pass& pass, params.passes) { + if(pass.type != type) { + pass_offset += pass.components; + continue; + } + + float *in = (float*)buffer.data_pointer + pass_offset; + int pass_stride = params.get_passes_size(); + + float scale = (pass.filter)? 1.0f/(float)sample: 1.0f; + float scale_exposure = (pass.exposure)? scale*exposure: scale; + + int size = params.width*params.height; + + if(components == 1) { + assert(pass.components == components); + + /* scalar */ + for(int i = 0; i < size; i++, in += pass_stride, pixels++) { + float f = *in; + + pixels[0] = f*scale_exposure; + } + } + else if(components == 3) { + assert(pass.components == 4); + + /* RGB/vector */ + for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) { + float3 f = make_float3(in[0], in[1], in[2]); + + pixels[0] = f.x*scale_exposure; + pixels[1] = f.y*scale_exposure; + pixels[2] = f.z*scale_exposure; + } + } + else if(components == 4) { + assert(pass.components == components); + + /* RGBA */ + for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) { + float4 f = make_float4(in[0], in[1], in[2], in[3]); - rgba.x = rgba.x*exposure; - rgba.y = rgba.y*exposure; - rgba.z = rgba.z*exposure; + pixels[0] = f.x*scale_exposure; + pixels[1] = f.y*scale_exposure; + pixels[2] = f.z*scale_exposure; - /* clamp since alpha might be > 1.0 due to russian roulette */ - rgba.w = clamp(rgba.w, 0.0f, 1.0f); + /* clamp since alpha might be > 1.0 due to russian roulette */ + pixels[3] = clamp(f.w*scale, 0.0f, 1.0f); + } + } - out[i] = rgba; + return true; } - return out; + return false; } /* Display Buffer */ diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h index f4a9b37c09b..3062e5ae3e4 100644 --- a/intern/cycles/render/buffers.h +++ b/intern/cycles/render/buffers.h @@ -21,6 +21,10 @@ #include "device_memory.h" +#include "film.h" + +#include "kernel_types.h" + #include "util_string.h" #include "util_thread.h" #include "util_types.h" @@ -45,32 +49,16 @@ public: int full_width; int full_height; - BufferParams() - { - width = 0; - height = 0; - - full_x = 0; - full_y = 0; - full_width = 0; - full_height = 0; - } - - void get_offset_stride(int& offset, int& stride) - { - offset = -(full_x + full_y*width); - stride = width; - } - - bool modified(const BufferParams& params) - { - return !(full_x == params.full_x - && full_y == params.full_y - && width == params.width - && height == params.height - && full_width == params.full_width - && full_height == params.full_height); - } + /* passes */ + vector<Pass> passes; + + /* functions */ + BufferParams(); + + void get_offset_stride(int& offset, int& stride); + bool modified(const BufferParams& params); + void add_pass(PassType type); + int get_passes_size(); }; /* Render Buffers */ @@ -80,7 +68,7 @@ public: /* buffer parameters */ BufferParams params; /* float buffer */ - device_vector<float4> buffer; + device_vector<float> buffer; /* random number generator state */ device_vector<uint> rng_state; /* mutex, must be locked manually by callers */ @@ -90,7 +78,9 @@ public: ~RenderBuffers(); void reset(Device *device, BufferParams& params); - float4 *copy_from_device(float exposure, int sample); + + bool copy_from_device(); + bool get_pass(PassType type, float exposure, int sample, int components, float *pixels); protected: void device_free(); diff --git a/intern/cycles/render/film.cpp b/intern/cycles/render/film.cpp index 0ae2866f182..bc51384b873 100644 --- a/intern/cycles/render/film.cpp +++ b/intern/cycles/render/film.cpp @@ -21,11 +21,111 @@ #include "film.h" #include "scene.h" +#include "util_foreach.h" + CCL_NAMESPACE_BEGIN +/* Pass */ + +void Pass::add(PassType type, vector<Pass>& passes) +{ + Pass pass; + + pass.type = type; + pass.filter = true; + pass.exposure = false; + + switch(type) { + case PASS_NONE: + pass.components = 0; + break; + case PASS_COMBINED: + pass.components = 4; + pass.exposure = true; + break; + case PASS_DEPTH: + pass.components = 1; + pass.filter = false; + break; + case PASS_NORMAL: + pass.components = 4; + break; + case PASS_UV: + pass.components = 4; + break; + case PASS_OBJECT_ID: + pass.components = 1; + pass.filter = false; + break; + case PASS_MATERIAL_ID: + pass.components = 1; + pass.filter = false; + break; + case PASS_DIFFUSE_COLOR: + pass.components = 4; + break; + case PASS_GLOSSY_COLOR: + pass.components = 4; + break; + case PASS_TRANSMISSION_COLOR: + pass.components = 4; + break; + case PASS_DIFFUSE_INDIRECT: + pass.components = 4; + pass.exposure = true; + break; + case PASS_GLOSSY_INDIRECT: + pass.components = 4; + pass.exposure = true; + break; + case PASS_TRANSMISSION_INDIRECT: + pass.components = 4; + pass.exposure = true; + break; + case PASS_DIFFUSE_DIRECT: + pass.components = 4; + pass.exposure = true; + break; + case PASS_GLOSSY_DIRECT: + pass.components = 4; + pass.exposure = true; + break; + case PASS_TRANSMISSION_DIRECT: + pass.components = 4; + pass.exposure = true; + break; + + case PASS_EMISSION: + pass.components = 4; + pass.exposure = true; + break; + case PASS_BACKGROUND: + pass.components = 4; + pass.exposure = true; + break; + } + + passes.push_back(pass); +} + +bool Pass::equals(const vector<Pass>& A, const vector<Pass>& B) +{ + if(A.size() != B.size()) + return false; + + for(int i = 0; i < A.size(); i++) + if(A[i].type != B[i].type) + return false; + + return true; +} + +/* Film */ + Film::Film() { exposure = 0.8f; + Pass::add(PASS_COMBINED, passes); need_update = true; } @@ -42,6 +142,82 @@ void Film::device_update(Device *device, DeviceScene *dscene) /* update __data */ kfilm->exposure = exposure; + kfilm->pass_flag = 0; + kfilm->pass_stride = 0; + kfilm->use_light_pass = 0; + + foreach(Pass& pass, passes) { + kfilm->pass_flag |= pass.type; + + switch(pass.type) { + case PASS_COMBINED: + kfilm->pass_combined = kfilm->pass_stride; + break; + case PASS_DEPTH: + kfilm->pass_depth = kfilm->pass_stride; + break; + case PASS_NORMAL: + kfilm->pass_normal = kfilm->pass_stride; + break; + case PASS_UV: + kfilm->pass_uv = kfilm->pass_stride; + break; + case PASS_OBJECT_ID: + kfilm->pass_object_id = kfilm->pass_stride; + break; + case PASS_MATERIAL_ID: + kfilm->pass_material_id = kfilm->pass_stride; + break; + case PASS_DIFFUSE_COLOR: + kfilm->pass_diffuse_color = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_GLOSSY_COLOR: + kfilm->pass_glossy_color = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_TRANSMISSION_COLOR: + kfilm->pass_transmission_color = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_DIFFUSE_INDIRECT: + kfilm->pass_diffuse_indirect = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_GLOSSY_INDIRECT: + kfilm->pass_glossy_indirect = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_TRANSMISSION_INDIRECT: + kfilm->pass_transmission_indirect = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_DIFFUSE_DIRECT: + kfilm->pass_diffuse_direct = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_GLOSSY_DIRECT: + kfilm->pass_glossy_direct = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_TRANSMISSION_DIRECT: + kfilm->pass_transmission_direct = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + + case PASS_EMISSION: + kfilm->pass_emission = kfilm->pass_stride; + kfilm->use_light_pass = 1; + break; + case PASS_BACKGROUND: + kfilm->pass_background = kfilm->pass_stride; + kfilm->use_light_pass = 1; + case PASS_NONE: + break; + } + + kfilm->pass_stride += pass.components; + } need_update = false; } @@ -52,7 +228,8 @@ void Film::device_free(Device *device, DeviceScene *dscene) bool Film::modified(const Film& film) { - return !(exposure == film.exposure); + return !(exposure == film.exposure + && Pass::equals(passes, film.passes)); } void Film::tag_update(Scene *scene) diff --git a/intern/cycles/render/film.h b/intern/cycles/render/film.h index df24fad3725..511ad316460 100644 --- a/intern/cycles/render/film.h +++ b/intern/cycles/render/film.h @@ -20,6 +20,9 @@ #define __FILM_H__ #include "util_string.h" +#include "util_vector.h" + +#include "kernel_types.h" CCL_NAMESPACE_BEGIN @@ -27,9 +30,21 @@ class Device; class DeviceScene; class Scene; +class Pass { +public: + PassType type; + int components; + bool filter; + bool exposure; + + static void add(PassType type, vector<Pass>& passes); + static bool equals(const vector<Pass>& A, const vector<Pass>& B); +}; + class Film { public: float exposure; + vector<Pass> passes; bool need_update; Film(); diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp index 3a9f0add735..f83c85c632d 100644 --- a/intern/cycles/render/object.cpp +++ b/intern/cycles/render/object.cpp @@ -36,6 +36,7 @@ Object::Object() mesh = NULL; tfm = transform_identity(); visibility = ~0; + pass_id = 0; } Object::~Object() @@ -135,6 +136,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene /* todo: correct for displacement, and move to a better place */ float uniform_scale; float surface_area = 0.0f; + float pass_id = ob->pass_id; if(transform_uniform_scale(tfm, uniform_scale)) { map<Mesh*, float>::iterator it = surface_area_map.find(mesh); @@ -171,7 +173,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene memcpy(&objects[offset], &tfm, sizeof(float4)*4); memcpy(&objects[offset+4], &itfm, sizeof(float4)*4); memcpy(&objects[offset+8], &ntfm, sizeof(float4)*4); - objects[offset+12] = make_float4(surface_area, 0.0f, 0.0f, 0.0f); + objects[offset+12] = make_float4(surface_area, pass_id, 0.0f, 0.0f); i++; diff --git a/intern/cycles/render/object.h b/intern/cycles/render/object.h index 7fe83cf7d91..14da2cfb35d 100644 --- a/intern/cycles/render/object.h +++ b/intern/cycles/render/object.h @@ -41,6 +41,7 @@ public: Transform tfm; BoundBox bounds; ustring name; + int pass_id; vector<ParamValue> attributes; uint visibility; diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp index 6e827ec94bb..12968a79ab2 100644 --- a/intern/cycles/render/shader.cpp +++ b/intern/cycles/render/shader.cpp @@ -35,6 +35,7 @@ CCL_NAMESPACE_BEGIN Shader::Shader() { name = ""; + pass_id = 0; graph = NULL; graph_bump = NULL; @@ -167,7 +168,7 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc if(scene->shaders.size() == 0) return; - uint shader_flag_size = scene->shaders.size()*2; + uint shader_flag_size = scene->shaders.size()*4; uint *shader_flag = dscene->shader_flag.resize(shader_flag_size); uint i = 0; @@ -184,7 +185,9 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc flag |= SD_HOMOGENEOUS_VOLUME; shader_flag[i++] = flag; + shader_flag[i++] = shader->pass_id; shader_flag[i++] = flag; + shader_flag[i++] = shader->pass_id; } device->tex_alloc("__shader_flag", dscene->shader_flag); diff --git a/intern/cycles/render/shader.h b/intern/cycles/render/shader.h index 45efa123ef6..35f3cfe27f5 100644 --- a/intern/cycles/render/shader.h +++ b/intern/cycles/render/shader.h @@ -47,6 +47,7 @@ class Shader { public: /* name */ string name; + int pass_id; /* shader graph */ ShaderGraph *graph; |