Cycles: Render Passes
authorBrecht Van Lommel <brechtvanlommel@pandora.be>
Wed, 25 Jan 2012 17:23:52 +0000 (17:23 +0000)
committerBrecht Van Lommel <brechtvanlommel@pandora.be>
Wed, 25 Jan 2012 17:23:52 +0000 (17:23 +0000)
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

36 files changed:
intern/cycles/blender/addon/ui.py
intern/cycles/blender/blender_object.cpp
intern/cycles/blender/blender_session.cpp
intern/cycles/blender/blender_shader.cpp
intern/cycles/device/device_cpu.cpp
intern/cycles/kernel/CMakeLists.txt
intern/cycles/kernel/kernel.cl
intern/cycles/kernel/kernel.cpp
intern/cycles/kernel/kernel.cu
intern/cycles/kernel/kernel.h
intern/cycles/kernel/kernel_accumulate.h [new file with mode: 0644]
intern/cycles/kernel/kernel_emission.h
intern/cycles/kernel/kernel_film.h
intern/cycles/kernel/kernel_object.h
intern/cycles/kernel/kernel_optimized.cpp
intern/cycles/kernel/kernel_passes.h [new file with mode: 0644]
intern/cycles/kernel/kernel_path.h
intern/cycles/kernel/kernel_random.h
intern/cycles/kernel/kernel_shader.h
intern/cycles/kernel/kernel_types.h
intern/cycles/kernel/svm/svm_types.h
intern/cycles/render/buffers.cpp
intern/cycles/render/buffers.h
intern/cycles/render/film.cpp
intern/cycles/render/film.h
intern/cycles/render/object.cpp
intern/cycles/render/object.h
intern/cycles/render/shader.cpp
intern/cycles/render/shader.h
source/blender/blenkernel/BKE_node.h
source/blender/makesdna/DNA_scene_types.h
source/blender/makesrna/intern/rna_render.c
source/blender/makesrna/intern/rna_scene.c
source/blender/nodes/composite/node_composite_tree.c
source/blender/nodes/composite/nodes/node_composite_image.c
source/blender/render/intern/source/render_result.c

index 763cff0bd7d90bd2314873fb4d4d00dc7fd5994a..6353bf37a15b94473040f45e9844a20d2c67e86a 100644 (file)
@@ -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")
 
index c805bd0306349201c7330918e34ef71df509cd4b..afcfcd9506342b86f617ac76b482014ebfef6006 100644 (file)
@@ -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 */
index b052fb3e195a742b7907e310e53cfae4dad619a4..ff1c32831bb130279d51f38070286c38ad3245b3 100644 (file)
@@ -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()
index f7b750e6dab8b5cfdb1d86f7710c483c40d1a310..5310e35dc25b06201bb1ec755c8136f02e7d799f 100644 (file)
@@ -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()) {
index 25da978edd8f053c870f582435d7b622caf603ca..2ca599f6c674c07bb30fdd8a5e227e645a54dffc 100644 (file)
@@ -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);
                }
        }
index b6e749f3fcbbf9d2f790d790fb135f8d6b115d36..59fe9709e61abb5e82b1941cdd36d11fddb1fc78 100644 (file)
@@ -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
index 305d81339c2ada2151e0f49388983903fc306018..f98414d4f0221fe17205ebd9a74982f86701510b 100644 (file)
@@ -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,
index 9e0d252772b9bb32202d6de6d237baed306bd4ea..a93f6172d28522b4d3b64340a33356d78b0ddfa4 100644 (file)
@@ -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);
 }
index 4e585fd563db8671c28e385f36277bf0a6f6f74e..fc7992c87f0ab58a73ab6c0cd53e22a1dcea27ad 100644 (file)
@@ -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;
index df6b5ee92dab7596824e971be8c8a4edb500719c..26c0bcd6d1a531d8b32af904d581db1e4284fe55 100644 (file)
@@ -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 (file)
index 0000000..e71fad5
--- /dev/null
@@ -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
+
index 51698f3a9bd428bd4611d50737bd5096b2b87333..b3a5b2bfcb41da8734afd591330423485551e685 100644 (file)
@@ -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 */
index cd8acc9647abf76b5ef38b0586940b82b66f876a..d0fb5402291c1c8608317db8a3fda1d0c6f9d7f8 100644 (file)
@@ -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
index b4c42d3bacc489ba81d070b0492d996cd8d773bf..318a6fea489cdb7ff88f6b77b4b74dd1d6000829 100644 (file)
@@ -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
 
index 50341021d9d8416c4fff7e4da87fbd1396473457..393686bb2031209c4e7cd3a10018c81b7263ae6e 100644 (file)
@@ -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 (file)
index 0000000..0e77581
--- /dev/null
@@ -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
+
index c80d2068506d41f3fe37d4f8cb0e94baceedec1b..c0bfa32040523388b7a1a4df90dee2a53df8d9d3 100644 (file)
 #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
index 41301ebd3dcee79ad4238005cad39a5616c48537..6d15100f8a32870c3d19bd101167d6055e2f2818 100644 (file)
@@ -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
index 9c59e1566a94157c555181b1223d7551eaebaf32..1d2cf46aa566da67b863769d8b6c3d7b1a41a3e9 100644 (file)
@@ -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)
index 008ec0bdf2842a6aab155fb4fc3a0642affbcd7f..b4b1da831627bcd93e040cbc2c25cd2dd011881b 100644 (file)
@@ -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 {
index ff6e9b94a0f106ed9385ec1586155b1f60a9b869..533a2944557defa0ffde284646d498658a776f77 100644 (file)
@@ -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)
index a6bbbc91901e536db3bc759fa582e9dfe7788e85..08dda9441119dba565a792c44ab9e570b86155c2 100644 (file)
@@ -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"
 
 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 */
index f4a9b37c09ba4f1133bb1604fce79b657470c4ba..3062e5ae3e4701bcea5175bd17583f00ad5cb4a3 100644 (file)
 
 #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();
index 0ae2866f182262d495cb6b28d4dbafd2150b0cea..bc51384b873f3ff758b8f65dd8696cdb26aa20bf 100644 (file)
 #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)
index df24fad3725f4d5026eab3bd3ef89ecae811b849..511ad316460483909e67e101f41ca9ddc538b062 100644 (file)
@@ -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();
index 3a9f0add735572e2b4e38f58c4a4d7a9c8a6ea8c..f83c85c632d03d3b43f1f0c674354084a99e3ce9 100644 (file)
@@ -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++;
 
index 7fe83cf7d914f1afcc67f1adc108bbd83f7a8c3d..14da2cfb35d4f4fa420be4cd0e3114f72a75acd3 100644 (file)
@@ -41,6 +41,7 @@ public:
        Transform tfm;
        BoundBox bounds;
        ustring name;
+       int pass_id;
        vector<ParamValue> attributes;
        uint visibility;
 
index 6e827ec94bb623af7d499a30c6a57b51c0686987..12968a79ab2b26d343ea2056b29d0595adfc8dbd 100644 (file)
@@ -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);
index 45efa123ef678e1ab5d57c9d0e4eeb3f59bc3155..35f3cfe27f5e0d973ef0e965ea67d20060d97ded 100644 (file)
@@ -47,6 +47,7 @@ class Shader {
 public:
        /* name */
        string name;
+       int pass_id;
 
        /* shader graph */
        ShaderGraph *graph;
index 2fa587e168a58959e6c2d73529675972c4134b2c..0cccd8a366b57b2ddfb81da170e22ee695972650 100644 (file)
@@ -557,25 +557,34 @@ void                      ntreeGPUMaterialNodes(struct bNodeTree *ntree, struct GPUMaterial *mat);
 /* ************** COMPOSITE NODES *************** */
 
 /* output socket defines */
-#define RRES_OUT_IMAGE         0
-#define RRES_OUT_ALPHA         1
-#define RRES_OUT_Z                     2
-#define RRES_OUT_NORMAL                3
-#define RRES_OUT_UV                    4
-#define RRES_OUT_VEC           5
-#define RRES_OUT_RGBA          6
-#define RRES_OUT_DIFF          7
-#define RRES_OUT_SPEC          8
-#define RRES_OUT_SHADOW                9
-#define RRES_OUT_AO                    10
-#define RRES_OUT_REFLECT       11
-#define RRES_OUT_REFRACT       12
-#define RRES_OUT_INDIRECT      13
-#define RRES_OUT_INDEXOB       14
-#define RRES_OUT_INDEXMA       15
-#define RRES_OUT_MIST          16
-#define RRES_OUT_EMIT          17
-#define RRES_OUT_ENV           18
+#define RRES_OUT_IMAGE                         0
+#define RRES_OUT_ALPHA                         1
+#define RRES_OUT_Z                                     2
+#define RRES_OUT_NORMAL                                3
+#define RRES_OUT_UV                                    4
+#define RRES_OUT_VEC                           5
+#define RRES_OUT_RGBA                          6
+#define RRES_OUT_DIFF                          7
+#define RRES_OUT_SPEC                          8
+#define RRES_OUT_SHADOW                                9
+#define RRES_OUT_AO                                    10
+#define RRES_OUT_REFLECT                       11
+#define RRES_OUT_REFRACT                       12
+#define RRES_OUT_INDIRECT                      13
+#define RRES_OUT_INDEXOB                       14
+#define RRES_OUT_INDEXMA                       15
+#define RRES_OUT_MIST                          16
+#define RRES_OUT_EMIT                          17
+#define RRES_OUT_ENV                           18
+#define RRES_OUT_DIFF_DIRECT           19
+#define RRES_OUT_DIFF_INDIRECT         20
+#define RRES_OUT_DIFF_COLOR                    21
+#define RRES_OUT_GLOSSY_DIRECT         22
+#define RRES_OUT_GLOSSY_INDIRECT       23
+#define RRES_OUT_GLOSSY_COLOR          24
+#define RRES_OUT_TRANSM_DIRECT         25
+#define RRES_OUT_TRANSM_INDIRECT       26
+#define RRES_OUT_TRANSM_COLOR          27
 
 /* note: types are needed to restore callbacks, don't change values */
 #define CMP_NODE_VIEWER                201
index 6b13951348af761f42845718dfd4a372583e9ffb..1f42118dbc07eb8ef3217996b7866548a2642609 100644 (file)
@@ -202,25 +202,34 @@ typedef struct SceneRenderLayer {
 #define SCE_LAY_NEG_ZMASK      0x80000
 
 /* srl->passflag */
-#define SCE_PASS_COMBINED              (1<<0)
-#define SCE_PASS_Z                             (1<<1)
-#define SCE_PASS_RGBA                  (1<<2)
-#define SCE_PASS_DIFFUSE               (1<<3)
-#define SCE_PASS_SPEC                  (1<<4)
-#define SCE_PASS_SHADOW                        (1<<5)
-#define SCE_PASS_AO                            (1<<6)
-#define SCE_PASS_REFLECT               (1<<7)
-#define SCE_PASS_NORMAL                        (1<<8)
-#define SCE_PASS_VECTOR                        (1<<9)
-#define SCE_PASS_REFRACT               (1<<10)
-#define SCE_PASS_INDEXOB               (1<<11)
-#define SCE_PASS_UV                            (1<<12)
-#define SCE_PASS_INDIRECT              (1<<13)
-#define SCE_PASS_MIST                  (1<<14)
-#define SCE_PASS_RAYHITS               (1<<15)
-#define SCE_PASS_EMIT                  (1<<16)
-#define SCE_PASS_ENVIRONMENT   (1<<17)
-#define SCE_PASS_INDEXMA       (1<<18)
+#define SCE_PASS_COMBINED                      (1<<0)
+#define SCE_PASS_Z                                     (1<<1)
+#define SCE_PASS_RGBA                          (1<<2)
+#define SCE_PASS_DIFFUSE                       (1<<3)
+#define SCE_PASS_SPEC                          (1<<4)
+#define SCE_PASS_SHADOW                                (1<<5)
+#define SCE_PASS_AO                                    (1<<6)
+#define SCE_PASS_REFLECT                       (1<<7)
+#define SCE_PASS_NORMAL                                (1<<8)
+#define SCE_PASS_VECTOR                                (1<<9)
+#define SCE_PASS_REFRACT                       (1<<10)
+#define SCE_PASS_INDEXOB                       (1<<11)
+#define SCE_PASS_UV                                    (1<<12)
+#define SCE_PASS_INDIRECT                      (1<<13)
+#define SCE_PASS_MIST                          (1<<14)
+#define SCE_PASS_RAYHITS                       (1<<15)
+#define SCE_PASS_EMIT                          (1<<16)
+#define SCE_PASS_ENVIRONMENT           (1<<17)
+#define SCE_PASS_INDEXMA                       (1<<18)
+#define SCE_PASS_DIFFUSE_DIRECT                (1<<19)
+#define SCE_PASS_DIFFUSE_INDIRECT      (1<<20)
+#define SCE_PASS_DIFFUSE_COLOR         (1<<21)
+#define SCE_PASS_GLOSSY_DIRECT         (1<<22)
+#define SCE_PASS_GLOSSY_INDIRECT       (1<<23)
+#define SCE_PASS_GLOSSY_COLOR          (1<<24)
+#define SCE_PASS_TRANSM_DIRECT         (1<<25)
+#define SCE_PASS_TRANSM_INDIRECT       (1<<26)
+#define SCE_PASS_TRANSM_COLOR          (1<<27)
 
 /* note, srl->passflag is treestore element 'nr' in outliner, short still... */
 
index f00b97994d5bfd49e55ecaaffd47b7916f9e1c52..b88e452588ea903b381d4c7487ecf4cfbe96191d 100644 (file)
@@ -472,6 +472,15 @@ static void rna_def_render_pass(BlenderRNA *brna)
                {SCE_PASS_EMIT, "EMIT", 0, "Emit", ""},
                {SCE_PASS_ENVIRONMENT, "ENVIRONMENT", 0, "Environment", ""},
                {SCE_PASS_INDEXMA, "MATERIAL_INDEX", 0, "Material Index", ""},
+               {SCE_PASS_DIFFUSE_DIRECT, "DIFFUSE_DIRECT", 0, "Diffuse Direct", ""},
+               {SCE_PASS_DIFFUSE_INDIRECT, "DIFFUSE_INDIRECT", 0, "Diffuse Indirect", ""},
+               {SCE_PASS_DIFFUSE_COLOR, "DIFFUSE_COLOR", 0, "Diffuse Color", ""},
+               {SCE_PASS_GLOSSY_DIRECT, "GLOSSY_DIRECT", 0, "Glossy Direct", ""},
+               {SCE_PASS_GLOSSY_INDIRECT, "GLOSSY_INDIRECT", 0, "Glossy Indirect", ""},
+               {SCE_PASS_GLOSSY_COLOR, "GLOSSY_COLOR", 0, "Glossy Color", ""},
+               {SCE_PASS_TRANSM_DIRECT, "TRANSMISSION_DIRECT", 0, "Transmission Direct", ""},
+               {SCE_PASS_TRANSM_INDIRECT, "TRANSMISSION_INDIRECT", 0, "Transmission Indirect", ""},
+               {SCE_PASS_TRANSM_COLOR, "TRANSMISSION_COLOR", 0, "Transmission Color", ""},
                {0, NULL, 0, NULL, NULL}};
        
        srna= RNA_def_struct(brna, "RenderPass", NULL);
index cd033b1f329fd9da840513ed6fed3c82e9573126..c1df493eea5e4c7f6de403acb91e0463bf6fe71d 100644 (file)
@@ -2065,6 +2065,60 @@ void rna_def_render_layer_common(StructRNA *srna, int scene)
        RNA_def_property_ui_icon(prop, ICON_RESTRICT_RENDER_OFF, 1);
        if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
        else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_diffuse_direct", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_DIFFUSE_DIRECT);
+       RNA_def_property_ui_text(prop, "Diffuse Direct", "Deliver diffuse direct pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_diffuse_indirect", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_DIFFUSE_INDIRECT);
+       RNA_def_property_ui_text(prop, "Diffuse Indirect", "Deliver diffuse indirect pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_diffuse_color", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_DIFFUSE_COLOR);
+       RNA_def_property_ui_text(prop, "Diffuse Color", "Deliver diffuse color pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_glossy_direct", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_GLOSSY_DIRECT);
+       RNA_def_property_ui_text(prop, "Glossy Direct", "Deliver glossy direct pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_glossy_indirect", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_GLOSSY_INDIRECT);
+       RNA_def_property_ui_text(prop, "Glossy Indirect", "Deliver glossy indirect pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_glossy_color", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_GLOSSY_COLOR);
+       RNA_def_property_ui_text(prop, "Glossy Color", "Deliver glossy color pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_transmission_direct", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_TRANSM_DIRECT);
+       RNA_def_property_ui_text(prop, "Transmission Direct", "Deliver transmission direct pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_transmission_indirect", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_TRANSM_INDIRECT);
+       RNA_def_property_ui_text(prop, "Transmission Indirect", "Deliver transmission indirect pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
+
+       prop= RNA_def_property(srna, "use_pass_transmission_color", PROP_BOOLEAN, PROP_NONE);
+       RNA_def_property_boolean_sdna(prop, NULL, "passflag", SCE_PASS_TRANSM_COLOR);
+       RNA_def_property_ui_text(prop, "Transmission Color", "Deliver transmission color pass");
+       if(scene) RNA_def_property_update(prop, NC_SCENE|ND_RENDER_OPTIONS, "rna_SceneRenderLayer_pass_update");
+       else RNA_def_property_clear_flag(prop, PROP_EDITABLE);
 }
 
 static void rna_def_scene_game_recast_data(BlenderRNA *brna)
index 93456d39c9cf374f7d5a836a80e375e3161124b2..dfc0bcb3f755c0ff77e21438cfc802ca7be04ede 100644 (file)
@@ -694,7 +694,28 @@ static void force_hidden_passes(bNode *node, int passflag)
        if(!(passflag & SCE_PASS_EMIT)) sock->flag |= SOCK_UNAVAIL;
        sock= BLI_findlink(&node->outputs, RRES_OUT_ENV);
        if(!(passflag & SCE_PASS_ENVIRONMENT)) sock->flag |= SOCK_UNAVAIL;
-       
+
+       sock= BLI_findlink(&node->outputs, RRES_OUT_DIFF_DIRECT);
+       if(!(passflag & SCE_PASS_DIFFUSE_DIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_DIFF_INDIRECT);
+       if(!(passflag & SCE_PASS_DIFFUSE_INDIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_DIFF_COLOR);
+       if(!(passflag & SCE_PASS_DIFFUSE_COLOR)) sock->flag |= SOCK_UNAVAIL;
+
+       sock= BLI_findlink(&node->outputs, RRES_OUT_GLOSSY_DIRECT);
+       if(!(passflag & SCE_PASS_GLOSSY_DIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_GLOSSY_INDIRECT);
+       if(!(passflag & SCE_PASS_GLOSSY_INDIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_GLOSSY_COLOR);
+       if(!(passflag & SCE_PASS_GLOSSY_COLOR)) sock->flag |= SOCK_UNAVAIL;
+
+       sock= BLI_findlink(&node->outputs, RRES_OUT_TRANSM_DIRECT);
+       if(!(passflag & SCE_PASS_TRANSM_DIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_TRANSM_INDIRECT);
+       if(!(passflag & SCE_PASS_TRANSM_INDIRECT)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_TRANSM_COLOR);
+       if(!(passflag & SCE_PASS_TRANSM_COLOR)) sock->flag |= SOCK_UNAVAIL;
+       sock= BLI_findlink(&node->outputs, RRES_OUT_TRANSM_COLOR);
 }
 
 /* based on rules, force sockets hidden always */
index 57b5cec4256d1b76f645ccbc9285f04518a670be..997bd72d5c496c7ad4a0342cd0ad2816d20393f2 100644 (file)
 /* **************** IMAGE (and RenderResult, multilayer image) ******************** */
 
 static bNodeSocketTemplate cmp_node_rlayers_out[]= {
-       {       SOCK_RGBA, 0, "Image",          0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 0, "Alpha",         1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 0, "Z",                     1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_VECTOR, 0, "Normal",       0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_VECTOR, 0, "UV",           1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_VECTOR, 0, "Speed",        1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Color",          0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Diffuse",        0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Specular",       0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Shadow",         0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "AO",                     0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Reflect",        0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Refract",        0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Indirect",       0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 0, "IndexOB",       0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 0, "IndexMA",       0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_FLOAT, 0, "Mist",          0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Emit",           0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
-       {       SOCK_RGBA, 0, "Environment",0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Image",                                  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_FLOAT, 0, "Alpha",                                 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_FLOAT, 0, "Z",                                             1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_VECTOR, 0, "Normal",                               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_VECTOR, 0, "UV",                                   1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_VECTOR, 0, "Speed",                                1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Color",                                  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Diffuse",                                0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Specular",                               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Shadow",                                 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "AO",                                             0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Reflect",                                0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Refract",                                0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Indirect",                               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_FLOAT, 0, "IndexOB",                               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_FLOAT, 0, "IndexMA",                               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_FLOAT, 0, "Mist",                                  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Emit",                                   0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Environment",                    0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Diffuse Direct",                 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Diffuse Indirect",               0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Diffuse Color",                  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Glossy Direct",                  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Glossy Indirect",                0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Glossy Color",                   0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Transmission Direct",    0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Transmission Indirect",  0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+       {       SOCK_RGBA, 0, "Transmission Color",             0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
        {       -1, 0, ""       }
 };
 
@@ -238,6 +247,24 @@ static void outputs_multilayer_get(RenderData *rd, RenderLayer *rl, bNodeStack *
                out[RRES_OUT_EMIT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_EMIT);
        if(out[RRES_OUT_ENV]->hasoutput)
                out[RRES_OUT_ENV]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_ENVIRONMENT);
+       if(out[RRES_OUT_DIFF_DIRECT]->hasoutput)
+               out[RRES_OUT_DIFF_DIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_DIFFUSE_DIRECT);
+       if(out[RRES_OUT_DIFF_INDIRECT]->hasoutput)
+               out[RRES_OUT_DIFF_INDIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_DIFFUSE_INDIRECT);
+       if(out[RRES_OUT_DIFF_COLOR]->hasoutput)
+               out[RRES_OUT_DIFF_COLOR]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_DIFFUSE_COLOR);
+       if(out[RRES_OUT_GLOSSY_DIRECT]->hasoutput)
+               out[RRES_OUT_GLOSSY_DIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_GLOSSY_DIRECT);
+       if(out[RRES_OUT_GLOSSY_INDIRECT]->hasoutput)
+               out[RRES_OUT_GLOSSY_INDIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_GLOSSY_INDIRECT);
+       if(out[RRES_OUT_GLOSSY_COLOR]->hasoutput)
+               out[RRES_OUT_GLOSSY_COLOR]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_GLOSSY_COLOR);
+       if(out[RRES_OUT_TRANSM_DIRECT]->hasoutput)
+               out[RRES_OUT_TRANSM_DIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_TRANSM_DIRECT);
+       if(out[RRES_OUT_TRANSM_INDIRECT]->hasoutput)
+               out[RRES_OUT_TRANSM_INDIRECT]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_TRANSM_INDIRECT);
+       if(out[RRES_OUT_TRANSM_COLOR]->hasoutput)
+               out[RRES_OUT_TRANSM_COLOR]->data= compbuf_multilayer_get(rd, rl, ima, iuser, SCE_PASS_TRANSM_COLOR);
 }
 
 
@@ -402,6 +429,24 @@ static void node_composit_rlayers_out(RenderData *rd, RenderLayer *rl, bNodeStac
                out[RRES_OUT_EMIT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_EMIT);
        if(out[RRES_OUT_ENV]->hasoutput)
                out[RRES_OUT_ENV]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_ENVIRONMENT);
+       if(out[RRES_OUT_DIFF_DIRECT]->hasoutput)
+               out[RRES_OUT_DIFF_DIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_DIFFUSE_DIRECT);
+       if(out[RRES_OUT_DIFF_INDIRECT]->hasoutput)
+               out[RRES_OUT_DIFF_INDIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_DIFFUSE_INDIRECT);
+       if(out[RRES_OUT_DIFF_COLOR]->hasoutput)
+               out[RRES_OUT_DIFF_COLOR]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_DIFFUSE_COLOR);
+       if(out[RRES_OUT_GLOSSY_DIRECT]->hasoutput)
+               out[RRES_OUT_GLOSSY_DIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_GLOSSY_DIRECT);
+       if(out[RRES_OUT_GLOSSY_INDIRECT]->hasoutput)
+               out[RRES_OUT_GLOSSY_INDIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_GLOSSY_INDIRECT);
+       if(out[RRES_OUT_GLOSSY_COLOR]->hasoutput)
+               out[RRES_OUT_GLOSSY_COLOR]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_GLOSSY_COLOR);
+       if(out[RRES_OUT_TRANSM_DIRECT]->hasoutput)
+               out[RRES_OUT_TRANSM_DIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_TRANSM_DIRECT);
+       if(out[RRES_OUT_TRANSM_INDIRECT]->hasoutput)
+               out[RRES_OUT_TRANSM_INDIRECT]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_TRANSM_INDIRECT);
+       if(out[RRES_OUT_TRANSM_COLOR]->hasoutput)
+               out[RRES_OUT_TRANSM_COLOR]->data= compbuf_from_pass(rd, rl, rectx, recty, SCE_PASS_TRANSM_COLOR);
 }
 
 static void node_composit_exec_rlayers(void *data, bNode *node, bNodeStack **UNUSED(in), bNodeStack **out)
index e825c9d842e93fb261fb6c755d1f556883af7e05..30853e8e1f6991c5d2b4053f6bfb55c90da565d2 100644 (file)
@@ -224,6 +224,60 @@ static const char *get_pass_name(int passtype, int channel)
                if(channel==1) return "Rayhits.G";
                return "Rayhits.B";
        }
+       if(passtype == SCE_PASS_DIFFUSE_DIRECT) {
+               if(channel==-1) return "DiffDir";
+               if(channel==0) return "DiffDir.R";
+               if(channel==1) return "DiffDir.G";
+               return "DiffDir.B";
+       }
+       if(passtype == SCE_PASS_DIFFUSE_INDIRECT) {
+               if(channel==-1) return "DiffInd";
+               if(channel==0) return "DiffInd.R";
+               if(channel==1) return "DiffInd.G";
+               return "DiffInd.B";
+       }
+       if(passtype == SCE_PASS_DIFFUSE_COLOR) {
+               if(channel==-1) return "DiffCol";
+               if(channel==0) return "DiffCol.R";
+               if(channel==1) return "DiffCol.G";
+               return "DiffCol.B";
+       }
+       if(passtype == SCE_PASS_GLOSSY_DIRECT) {
+               if(channel==-1) return "GlossDir";
+               if(channel==0) return "GlossDir.R";
+               if(channel==1) return "GlossDir.G";
+               return "GlossDir.B";
+       }
+       if(passtype == SCE_PASS_GLOSSY_INDIRECT) {
+               if(channel==-1) return "GlossInd";
+               if(channel==0) return "GlossInd.R";
+               if(channel==1) return "GlossInd.G";
+               return "GlossInd.B";
+       }
+       if(passtype == SCE_PASS_GLOSSY_COLOR) {
+               if(channel==-1) return "GlossCol";
+               if(channel==0) return "GlossCol.R";
+               if(channel==1) return "GlossCol.G";
+               return "GlossCol.B";
+       }
+       if(passtype == SCE_PASS_TRANSM_DIRECT) {
+               if(channel==-1) return "TransDir";
+               if(channel==0) return "TransDir.R";
+               if(channel==1) return "TransDir.G";
+               return "TransDir.B";
+       }
+       if(passtype == SCE_PASS_TRANSM_INDIRECT) {
+               if(channel==-1) return "TransInd";
+               if(channel==0) return "TransInd.R";
+               if(channel==1) return "TransInd.G";
+               return "TransInd.B";
+       }
+       if(passtype == SCE_PASS_TRANSM_COLOR) {
+               if(channel==-1) return "TransCol";
+               if(channel==0) return "TransCol.R";
+               if(channel==1) return "TransCol.G";
+               return "TransCol.B";
+       }
        return "Unknown";
 }
 
@@ -286,6 +340,34 @@ static int passtype_from_name(const char *str)
        
        if(strcmp(str, "RayHits")==0)
                return SCE_PASS_RAYHITS;
+
+       if(strcmp(str, "DiffDir")==0)
+               return SCE_PASS_DIFFUSE_DIRECT;
+
+       if(strcmp(str, "DiffInd")==0)
+               return SCE_PASS_DIFFUSE_INDIRECT;
+
+       if(strcmp(str, "DiffCol")==0)
+               return SCE_PASS_DIFFUSE_COLOR;
+
+       if(strcmp(str, "GlossDir")==0)
+               return SCE_PASS_GLOSSY_DIRECT;
+
+       if(strcmp(str, "GlossInd")==0)
+               return SCE_PASS_GLOSSY_INDIRECT;
+
+       if(strcmp(str, "GlossCol")==0)
+               return SCE_PASS_GLOSSY_COLOR;
+
+       if(strcmp(str, "TransDir")==0)
+               return SCE_PASS_TRANSM_DIRECT;
+
+       if(strcmp(str, "TransInd")==0)
+               return SCE_PASS_TRANSM_INDIRECT;
+
+       if(strcmp(str, "TransCol")==0)
+               return SCE_PASS_TRANSM_COLOR;
+
        return 0;
 }
 
@@ -430,6 +512,24 @@ RenderResult *render_result_new(Render *re, rcti *partrct, int crop, int savebuf
                        render_layer_add_pass(rr, rl, 1, SCE_PASS_MIST);
                if(rl->passflag & SCE_PASS_RAYHITS)
                        render_layer_add_pass(rr, rl, 4, SCE_PASS_RAYHITS);
+               if(srl->passflag  & SCE_PASS_DIFFUSE_DIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_DIFFUSE_DIRECT);
+               if(srl->passflag  & SCE_PASS_DIFFUSE_INDIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_DIFFUSE_INDIRECT);
+               if(srl->passflag  & SCE_PASS_DIFFUSE_COLOR)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_DIFFUSE_COLOR);
+               if(srl->passflag  & SCE_PASS_GLOSSY_DIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_GLOSSY_DIRECT);
+               if(srl->passflag  & SCE_PASS_GLOSSY_INDIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_GLOSSY_INDIRECT);
+               if(srl->passflag  & SCE_PASS_GLOSSY_COLOR)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_GLOSSY_COLOR);
+               if(srl->passflag  & SCE_PASS_TRANSM_DIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_TRANSM_DIRECT);
+               if(srl->passflag  & SCE_PASS_TRANSM_INDIRECT)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_TRANSM_INDIRECT);
+               if(srl->passflag  & SCE_PASS_TRANSM_COLOR)
+                       render_layer_add_pass(rr, rl, 3, SCE_PASS_TRANSM_COLOR);
                
        }
        /* sss, previewrender and envmap don't do layers, so we make a default one */