diff options
author | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2012-01-25 21:23:52 +0400 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@pandora.be> | 2012-01-25 21:23:52 +0400 |
commit | f99343d3b8676543e2bd6acd6ee2274c21b1b388 (patch) | |
tree | fd40cd33691a783f82cf877e774d1b3a64d91ed3 /intern/cycles/kernel | |
parent | 14f475fccad7158098ddecc285c617f990b2f8b5 (diff) |
Cycles: Render Passes
Currently supported passes:
* Combined, Z, Normal, Object Index, Material Index, Emission, Environment,
Diffuse/Glossy/Transmission x Direct/Indirect/Color
Not supported yet:
* UV, Vector, Mist
Only enabled for CPU devices at the moment, will do GPU tweaks tommorrow,
also for environment importance sampling.
Documentation:
http://wiki.blender.org/index.php/Doc:2.6/Manual/Render/Cycles/Passes
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/CMakeLists.txt | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.cl | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.cpp | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.cu | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.h | 8 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_accumulate.h | 283 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_emission.h | 23 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_film.h | 13 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_object.h | 10 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_optimized.cpp | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_passes.h | 146 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_path.h | 109 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_random.h | 44 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_shader.h | 108 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_types.h | 99 | ||||
-rw-r--r-- | intern/cycles/kernel/svm/svm_types.h | 25 |
16 files changed, 758 insertions, 128 deletions
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index b6e749f3fcb..59fe9709e61 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -15,6 +15,7 @@ set(SRC set(SRC_HEADERS kernel.h + kernel_accumulate.h kernel_bvh.h kernel_camera.h kernel_compat_cpu.h @@ -30,6 +31,7 @@ set(SRC_HEADERS kernel_mbvh.h kernel_montecarlo.h kernel_object.h + kernel_passes.h kernel_path.h kernel_qbvh.h kernel_random.h diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 305d81339c2..f98414d4f02 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -28,7 +28,7 @@ __kernel void kernel_ocl_path_trace( __constant KernelData *data, - __global float4 *buffer, + __global float *buffer, __global uint *rng_state, #define KERNEL_TEX(type, ttype, name) \ @@ -56,7 +56,7 @@ __kernel void kernel_ocl_path_trace( __kernel void kernel_ocl_tonemap( __constant KernelData *data, __global uchar4 *rgba, - __global float4 *buffer, + __global float *buffer, #define KERNEL_TEX(type, ttype, name) \ __global type *name, diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index 9e0d252772b..a93f6172d28 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -204,14 +204,14 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t /* Path Tracing */ -void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) +void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) { kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride); } /* Tonemapping */ -void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride) { kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride); } diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index 4e585fd563d..fc7992c87f0 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -26,7 +26,7 @@ #include "kernel_path.h" #include "kernel_displace.h" -extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride) +extern "C" __global__ void kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; int y = sy + blockDim.y*blockIdx.y + threadIdx.y; @@ -35,7 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_stat kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride); } -extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride) +extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; int y = sy + blockDim.y*blockIdx.y + threadIdx.y; diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h index df6b5ee92da..26c0bcd6d1a 100644 --- a/intern/cycles/kernel/kernel.h +++ b/intern/cycles/kernel/kernel.h @@ -36,17 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg); void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size); void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height); -void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, +void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride); -void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, +void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride); void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); #ifdef WITH_OPTIMIZED_KERNEL -void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, +void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride); -void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, +void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride); void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h new file mode 100644 index 00000000000..e71fad58c96 --- /dev/null +++ b/intern/cycles/kernel/kernel_accumulate.h @@ -0,0 +1,283 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + */ + +CCL_NAMESPACE_BEGIN + +/* BSDF Eval + * + * BSDF evaluation result, split per BSDF type. This is used to accumulate + * render passes separately. */ + +__device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass) +{ +#ifdef __PASSES__ + eval->use_light_pass = use_light_pass; + + if(eval->use_light_pass) { + eval->diffuse = make_float3(0.0f, 0.0f, 0.0f); + eval->glossy = make_float3(0.0f, 0.0f, 0.0f); + eval->transmission = make_float3(0.0f, 0.0f, 0.0f); + eval->transparent = make_float3(0.0f, 0.0f, 0.0f); + + if(type == CLOSURE_BSDF_TRANSPARENT_ID) + eval->transparent = value; + else if(CLOSURE_IS_BSDF_DIFFUSE(type)) + eval->diffuse = value; + else if(CLOSURE_IS_BSDF_GLOSSY(type)) + eval->glossy = value; + else + eval->transmission = value; + } + else + eval->diffuse = value; +#else + *eval = value; +#endif +} + +__device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + if(CLOSURE_IS_BSDF_DIFFUSE(type)) + eval->diffuse += value; + else if(CLOSURE_IS_BSDF_GLOSSY(type)) + eval->glossy += value; + else + eval->transmission += value; + + /* skipping transparent, this function is used by for eval(), will be zero then */ + } + else + eval->diffuse += value; +#else + *eval += value; +#endif +} + +__device_inline bool bsdf_eval_is_zero(BsdfEval *eval) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + return is_zero(eval->diffuse) + && is_zero(eval->glossy) + && is_zero(eval->transmission) + && is_zero(eval->transparent); + } + else + return is_zero(eval->diffuse); +#else + return is_zero(*eval); +#endif +} + +__device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value) +{ +#ifdef __PASSES__ + if(eval->use_light_pass) { + eval->diffuse *= value; + eval->glossy *= value; + eval->transmission *= value; + + /* skipping transparent, this function is used by for eval(), will be zero then */ + } + else + eval->diffuse *= value; +#else + *eval *= value; +#endif +} + +/* Path Radiance + * + * We accumulate different render passes separately. After summing at the end + * to get the combined result, it should be identical. We definte directly + * visible as the first non-transparent hit, while indirectly visible are the + * bounces after that. */ + +__device_inline void path_radiance_init(PathRadiance *L, int use_light_pass) +{ + /* clear all */ +#ifdef __PASSES__ + L->use_light_pass = use_light_pass; + + if(use_light_pass) { + L->indirect = make_float3(0.0f, 0.0f, 0.0f); + L->direct_throughput = make_float3(0.0f, 0.0f, 0.0f); + L->direct_emission = make_float3(0.0f, 0.0f, 0.0f); + + L->color_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->color_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->color_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->direct_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->direct_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->direct_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->indirect_diffuse = make_float3(0.0f, 0.0f, 0.0f); + L->indirect_glossy = make_float3(0.0f, 0.0f, 0.0f); + L->indirect_transmission = make_float3(0.0f, 0.0f, 0.0f); + + L->emission = make_float3(0.0f, 0.0f, 0.0f); + L->background = make_float3(0.0f, 0.0f, 0.0f); + } + else + L->emission = make_float3(0.0f, 0.0f, 0.0f); +#else + *L = make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput, + BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label) +{ + float inverse_pdf = 1.0f/bsdf_pdf; + +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) { + if(bsdf_label & LABEL_TRANSPARENT) { + /* transparent bounce before first hit */ + *throughput *= bsdf_eval->transparent*inverse_pdf; + } + else { + /* first on directly visible surface */ + float3 value = *throughput*inverse_pdf; + + L->indirect_diffuse = bsdf_eval->diffuse*value; + L->indirect_glossy = bsdf_eval->glossy*value; + L->indirect_transmission = bsdf_eval->transmission*value; + + *throughput = L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission; + + L->direct_throughput = *throughput; + } + } + else { + /* indirectly visible through BSDF */ + float3 sum = (bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission + bsdf_eval->transparent)*inverse_pdf; + *throughput *= sum; + } + } + else + *throughput *= bsdf_eval->diffuse*inverse_pdf; +#else + *throughput *= *bsdf_eval*inverse_pdf; +#endif +} + +__device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) + L->emission += throughput*value; + else if(bounce == 1) + L->direct_emission += throughput*value; + else + L->indirect += throughput*value; + } + else + L->emission += throughput*value; +#else + *L += throughput*value; +#endif +} + +__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) { + /* directly visible lighting */ + L->direct_diffuse += throughput*bsdf_eval->diffuse; + L->direct_glossy += throughput*bsdf_eval->glossy; + L->direct_transmission += throughput*bsdf_eval->transmission; + } + else { + /* indirectly visible lighting after BSDF bounce */ + float3 sum = bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission; + L->indirect += throughput*sum; + } + } + else + L->emission += throughput*bsdf_eval->diffuse; +#else + *L += throughput*(*bsdf_eval); +#endif +} + +__device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + if(bounce == 0) + L->background += throughput*value; + else if(bounce == 1) + L->direct_emission += throughput*value; + else + L->indirect += throughput*value; + } + else + L->emission += throughput*value; +#else + *L += throughput*value; +#endif +} + +__device_inline float3 safe_divide_color(float3 a, float3 b) +{ + float x, y, z; + + x = (b.x != 0.0f)? a.x/b.x: 0.0f; + y = (b.y != 0.0f)? a.y/b.y: 0.0f; + z = (b.z != 0.0f)? a.z/b.z: 0.0f; + + return make_float3(x, y, z); +} + +__device_inline float3 path_radiance_sum(PathRadiance *L) +{ +#ifdef __PASSES__ + if(L->use_light_pass) { + /* this division is a bit ugly, but means we only have to keep track of + only a single throughput further along the path, here we recover just + the indirect parth that is not influenced by any particular BSDF type */ + L->direct_emission = safe_divide_color(L->direct_emission, L->direct_throughput); + L->direct_diffuse += L->indirect_diffuse*L->direct_emission; + L->direct_glossy += L->indirect_glossy*L->direct_emission; + L->direct_transmission += L->indirect_transmission*L->direct_emission; + + L->indirect = safe_divide_color(L->indirect, L->direct_throughput); + L->indirect_diffuse *= L->indirect; + L->indirect_glossy *= L->indirect; + L->indirect_transmission *= L->indirect; + + return L->emission + L->background + + L->direct_diffuse + L->direct_glossy + L->direct_transmission + + L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission; + } + else + return L->emission; +#else + return *L; +#endif +} + +CCL_NAMESPACE_END + diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index 51698f3a9bd..b3a5b2bfcb4 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -57,7 +57,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando, } __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, - float randt, float rando, float randu, float randv, Ray *ray, float3 *eval) + float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval) { LightSample ls; @@ -83,32 +83,33 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, return false; /* evaluate closure */ - *eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D); + float3 light_eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D); - if(is_zero(*eval)) + if(is_zero(light_eval)) return false; /* todo: use visbility flag to skip lights */ /* evaluate BSDF at shading point */ float bsdf_pdf; - float3 bsdf_eval = shader_bsdf_eval(kg, sd, ls.D, &bsdf_pdf); - *eval *= bsdf_eval/pdf; - - if(is_zero(*eval)) - return false; + shader_bsdf_eval(kg, sd, ls.D, eval, &bsdf_pdf); if(ls.prim != ~0 || ls.type == LIGHT_BACKGROUND) { /* multiple importance sampling */ float mis_weight = power_heuristic(pdf, bsdf_pdf); - *eval *= mis_weight; + light_eval *= mis_weight; } /* todo: clean up these weights */ else if(ls.shader & SHADER_AREA_LIGHT) - *eval *= 0.25f; /* area lamp */ + light_eval *= 0.25f; /* area lamp */ else if(ls.t != FLT_MAX) - *eval *= 0.25f*M_1_PI_F; /* point lamp */ + light_eval *= 0.25f*M_1_PI_F; /* point lamp */ + + bsdf_eval_mul(eval, light_eval/pdf); + + if(bsdf_eval_is_zero(eval)) + return false; if(ls.shader & SHADER_CAST_SHADOW) { /* setup ray */ diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h index cd8acc9647a..d0fb5402291 100644 --- a/intern/cycles/kernel/kernel_film.h +++ b/intern/cycles/kernel/kernel_film.h @@ -48,15 +48,22 @@ __device uchar4 film_float_to_byte(float4 color) return result; } -__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +__device void kernel_film_tonemap(KernelGlobals *kg, + __global uchar4 *rgba, __global float *buffer, + int sample, int resolution, int x, int y, int offset, int stride) { + /* buffer offset */ int index = offset + x + y*stride; - float4 irradiance = buffer[index]; + rgba += index; + buffer += index*kernel_data.film.pass_stride; + + /* map colors */ + float4 irradiance = *(float4*)buffer; float4 float_result = film_map(kg, irradiance, sample); uchar4 byte_result = film_float_to_byte(float_result); - rgba[index] = byte_result; + *rgba = byte_result; } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_object.h b/intern/cycles/kernel/kernel_object.h index b4c42d3bacc..318a6fea489 100644 --- a/intern/cycles/kernel/kernel_object.h +++ b/intern/cycles/kernel/kernel_object.h @@ -64,5 +64,15 @@ __device_inline float object_surface_area(KernelGlobals *kg, int object) return f.x; } +__device_inline float object_pass_id(KernelGlobals *kg, int object) +{ + if(object == ~0) + return 0.0f; + + int offset = object*OBJECT_SIZE + OBJECT_PROPERTIES; + float4 f = kernel_tex_fetch(__objects, offset); + return f.y; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp index 50341021d9d..393686bb203 100644 --- a/intern/cycles/kernel/kernel_optimized.cpp +++ b/intern/cycles/kernel/kernel_optimized.cpp @@ -35,14 +35,14 @@ CCL_NAMESPACE_BEGIN /* Path Tracing */ -void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) +void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) { kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride); } /* Tonemapping */ -void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride) +void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride) { kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride); } diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h new file mode 100644 index 00000000000..0e775812eda --- /dev/null +++ b/intern/cycles/kernel/kernel_passes.h @@ -0,0 +1,146 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + */ + +CCL_NAMESPACE_BEGIN + +__device_inline void kernel_write_pass_float(__global float *buffer, int sample, float value) +{ + float *buf = buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_write_pass_float3(__global float *buffer, int sample, float3 value) +{ + float3 *buf = (float3*)buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_write_pass_float4(__global float *buffer, int sample, float4 value) +{ + float4 *buf = (float4*)buffer; + *buf = (sample == 0)? value: *buf + value; +} + +__device_inline void kernel_clear_passes(__global float *buffer, int sample, int pass_stride) +{ +#ifdef __PASSES__ + if(sample == 0 && pass_stride != 4) + for(int i = 4; i < pass_stride; i++) + buffer[i] = 0.0f; +#endif +} + +__device void kernel_write_data_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, + ShaderData *sd, int sample, int path_flag, float3 throughput) +{ +#ifdef __PASSES__ + if(!(path_flag & PATH_RAY_CAMERA)) + return; + + int flag = kernel_data.film.pass_flag; + + if(!(flag & PASS_ALL)) + return; + + /* todo: add alpha treshold */ + if(!(path_flag & PATH_RAY_TRANSPARENT)) { + if(sample == 0) { + if(flag & PASS_DEPTH) { + Transform tfm = kernel_data.cam.worldtocamera; + float depth = len(transform(&tfm, sd->P)); + + kernel_write_pass_float(buffer + kernel_data.film.pass_depth, sample, depth); + } + if(flag & PASS_OBJECT_ID) { + float id = object_pass_id(kg, sd->object); + kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, sample, id); + } + if(flag & PASS_MATERIAL_ID) { + float id = shader_pass_id(kg, sd); + kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, sample, id); + } + } + + if(flag & PASS_NORMAL) { + float3 normal = sd->N; + kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, sample, normal); + } + if(flag & PASS_UV) { + float3 uv = make_float3(0.0f, 0.0f, 0.0f); /* todo: request and lookup */ + kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, sample, uv); + } + } + + if(flag & (PASS_DIFFUSE_INDIRECT|PASS_DIFFUSE_COLOR|PASS_DIFFUSE_DIRECT)) + L->color_diffuse += shader_bsdf_diffuse(kg, sd)*throughput; + if(flag & (PASS_GLOSSY_INDIRECT|PASS_GLOSSY_COLOR|PASS_GLOSSY_DIRECT)) + L->color_glossy += shader_bsdf_glossy(kg, sd)*throughput; + if(flag & (PASS_TRANSMISSION_INDIRECT|PASS_TRANSMISSION_COLOR|PASS_TRANSMISSION_DIRECT)) + L->color_transmission += shader_bsdf_transmission(kg, sd)*throughput; +#endif +} + +__device void kernel_write_light_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, int sample) +{ +#ifdef __PASSES__ + int flag = kernel_data.film.pass_flag; + + if(!kernel_data.film.use_light_pass) + return; + + if(flag & PASS_DIFFUSE_INDIRECT) { + float3 color = safe_divide_color(L->indirect_diffuse, L->color_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, sample, color); + } + if(flag & PASS_GLOSSY_INDIRECT) { + float3 color = safe_divide_color(L->indirect_glossy, L->color_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, sample, color); + } + if(flag & PASS_TRANSMISSION_INDIRECT) { + float3 color = safe_divide_color(L->indirect_transmission, L->color_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, sample, color); + } + if(flag & PASS_DIFFUSE_DIRECT) { + float3 color = safe_divide_color(L->direct_diffuse, L->color_diffuse); + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, sample, color); + } + if(flag & PASS_GLOSSY_DIRECT) { + float3 color = safe_divide_color(L->direct_glossy, L->color_glossy); + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, sample, color); + } + if(flag & PASS_TRANSMISSION_DIRECT) { + float3 color = safe_divide_color(L->direct_transmission, L->color_transmission); + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, sample, color); + } + + if(flag & PASS_EMISSION) + kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, sample, L->emission); + if(flag & PASS_BACKGROUND) + kernel_write_pass_float3(buffer + kernel_data.film.pass_background, sample, L->background); + + if(flag & PASS_DIFFUSE_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, sample, L->color_diffuse); + if(flag & PASS_GLOSSY_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, sample, L->color_glossy); + if(flag & PASS_TRANSMISSION_COLOR) + kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, sample, L->color_transmission); +#endif +} + +CCL_NAMESPACE_END + diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index c80d2068506..c0bfa320405 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -25,39 +25,16 @@ #else #include "kernel_bvh.h" #endif +#include "kernel_accumulate.h" #include "kernel_camera.h" #include "kernel_shader.h" #include "kernel_light.h" #include "kernel_emission.h" #include "kernel_random.h" +#include "kernel_passes.h" CCL_NAMESPACE_BEGIN -#ifdef __MODIFY_TP__ -__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample) -{ - /* modify throughput to influence path termination probability, to avoid - darker regions receiving fewer samples than lighter regions. also RGB - are weighted differently. proper validation still remains to be done. */ - const float3 weights = make_float3(1.0f, 1.33f, 0.66f); - const float3 one = make_float3(1.0f, 1.0f, 1.0f); - const int minsample = 5; - const float minL = 0.1f; - - if(sample >= minsample) { - float3 L = buffer[offset + x + y*stride]; - float3 Lmin = make_float3(minL, minL, minL); - float correct = (float)(sample+1)/(float)sample; - - L = film_map(L*correct, sample); - - return weights/clamp(L, Lmin, one); - } - - return weights; -} -#endif - typedef struct PathState { uint flag; int bounce; @@ -168,7 +145,7 @@ __device_inline float path_state_terminate_probability(KernelGlobals *kg, PathSt return average(throughput); } -__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, float3 *light_L) +__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, BsdfEval *L_light) { if(ray->t == 0.0f) return false; @@ -208,7 +185,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra } if(!scene_intersect(kg, ray, PATH_RAY_SHADOW_TRANSPARENT, isect)) { - *light_L *= throughput; + bsdf_eval_mul(L_light, throughput); return false; } @@ -234,11 +211,14 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra return result; } -__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, float3 throughput) +__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer) { /* initialize */ - float3 L = make_float3(0.0f, 0.0f, 0.0f); - float Ltransparent = 0.0f; + PathRadiance L; + float3 throughput = make_float3(1.0f, 1.0f, 1.0f); + float L_transparent = 0.0f; + + path_radiance_init(&L, kernel_data.film.use_light_pass); #ifdef __EMISSION__ float ray_pdf = 0.0f; @@ -257,12 +237,12 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R if(!scene_intersect(kg, &ray, visibility, &isect)) { /* eval background shader if nothing hit */ if(kernel_data.background.transparent && (state.flag & PATH_RAY_CAMERA)) { - Ltransparent += average(throughput); + L_transparent += average(throughput); } else { /* sample background shader */ - float3 background_L = indirect_background(kg, &ray, state.flag, ray_pdf); - L += throughput*background_L; + float3 L_background = indirect_background(kg, &ray, state.flag, ray_pdf); + path_radiance_accum_background(&L, throughput, L_background, state.bounce); } break; @@ -274,19 +254,24 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R float rbsdf = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF); shader_eval_surface(kg, &sd, rbsdf, state.flag); + kernel_write_data_passes(kg, buffer, &L, &sd, sample, state.flag, throughput); + #ifdef __HOLDOUT__ if((sd.flag & SD_HOLDOUT) && (state.flag & PATH_RAY_CAMERA)) { float3 holdout_weight = shader_holdout_eval(kg, &sd); if(kernel_data.background.transparent) - Ltransparent += average(holdout_weight*throughput); + /* any throughput is ok, should all be identical here */ + L_transparent += average(holdout_weight*throughput); } #endif #ifdef __EMISSION__ /* emission */ - if(sd.flag & SD_EMISSION) - L += throughput*indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf); + if(sd.flag & SD_EMISSION) { + float3 emission = indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf); + path_radiance_accum_emission(&L, throughput, emission, state.bounce); + } #endif /* path termination. this is a strange place to put the termination, it's @@ -310,7 +295,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R float light_v = path_rng(kg, rng, sample, rng_offset + PRNG_LIGHT_V); Ray light_ray; - float3 light_L; + BsdfEval L_light; #ifdef __MULTI_LIGHT__ /* index -1 means randomly sample from distribution */ @@ -320,10 +305,10 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R #else const int i = -1; #endif - if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &light_L)) { + if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &L_light)) { /* trace shadow ray */ - if(!shadow_blocked(kg, &state, &light_ray, &isect, &light_L)) - L += throughput*light_L; + if(!shadow_blocked(kg, &state, &light_ray, &isect, &L_light)) + path_radiance_accum_light(&L, throughput, &L_light, state.bounce); } #ifdef __MULTI_LIGHT__ } @@ -338,7 +323,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R /* sample BSDF */ float bsdf_pdf; - float3 bsdf_eval; + BsdfEval bsdf_eval; float3 bsdf_omega_in; differential3 bsdf_domega_in; float bsdf_u = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF_U); @@ -350,11 +335,11 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R shader_release(kg, &sd); - if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) + if(bsdf_pdf == 0.0f || bsdf_eval_is_zero(&bsdf_eval)) break; /* modify throughput */ - throughput *= bsdf_eval/bsdf_pdf; + path_radiance_bsdf_bounce(&L, &throughput, &bsdf_eval, bsdf_pdf, state.bounce, label); /* set labels */ #if defined(__EMISSION__) || defined(__BACKGROUND__) @@ -374,18 +359,33 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R #endif } - return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent); + float3 L_sum = path_radiance_sum(&L); + + kernel_write_light_passes(kg, buffer, &L, sample); + + return make_float4(L_sum.x, L_sum.y, L_sum.z, 1.0f - L_transparent); } -__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride) +__device void kernel_path_trace(KernelGlobals *kg, + __global float *buffer, __global uint *rng_state, + int sample, int x, int y, int offset, int stride) { + /* buffer offset */ + int index = offset + x + y*stride; + int pass_stride = kernel_data.film.pass_stride; + + rng_state += index; + buffer += index*pass_stride; + + kernel_clear_passes(buffer, sample, pass_stride); + /* initialize random numbers */ RNG rng; float filter_u; float filter_v; - path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v); + path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v); /* sample camera ray */ Ray ray; @@ -396,23 +396,12 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, &ray); /* integrate */ -#ifdef __MODIFY_TP__ - float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample); - float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput; -#else - float3 throughput = make_float3(1.0f, 1.0f, 1.0f); - float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput); -#endif + float4 L = kernel_path_integrate(kg, &rng, sample, ray, buffer); /* accumulate result in output buffer */ - int index = offset + x + y*stride; - - if(sample == 0) - buffer[index] = L; - else - buffer[index] += L; + kernel_write_pass_float4(buffer, sample, L); - path_rng_end(kg, rng_state, rng, x, y, offset, stride); + path_rng_end(kg, rng_state, rng); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h index 41301ebd3dc..6d15100f8a3 100644 --- a/intern/cycles/kernel/kernel_random.h +++ b/intern/cycles/kernel/kernel_random.h @@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime #endif } -__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy) +__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy) { #ifdef __SOBOL_FULL_SCREEN__ uint px, py; @@ -135,19 +135,31 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, *rng ^= kernel_data.integrator.seed; - *fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x; - *fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y; + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x; + *fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y; + } #else - *rng = rng_state[offset + x + y*stride]; + *rng = *rng_state; *rng ^= kernel_data.integrator.seed; - *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); - *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); + *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + } #endif } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride) +__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) { /* nothing to do */ } @@ -163,21 +175,27 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension) return (float)*rng * (1.0f/(float)0xFFFFFFFF); } -__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy) +__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy) { /* load state */ - *rng = rng_state[offset + x + y*stride]; + *rng = *rng_state; *rng ^= kernel_data.integrator.seed; - *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); - *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + if(sample == 0) { + *fx = 0.5f; + *fy = 0.5f; + } + else { + *fx = path_rng(kg, rng, sample, PRNG_FILTER_U); + *fy = path_rng(kg, rng, sample, PRNG_FILTER_V); + } } -__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride) +__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng) { /* store state for next sample */ - rng_state[offset + x + y*stride] = rng; + *rng_state = rng; } #endif diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index 9c59e1566a9..1d2cf46aa56 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -75,7 +75,7 @@ __device_inline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, if(sd->shader & SHADER_SMOOTH_NORMAL) sd->N = triangle_smooth_normal(kg, sd->prim, sd->u, sd->v); - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __DPDU__ /* dPdu/dPdv */ @@ -166,7 +166,7 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, #endif } - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __DPDU__ /* dPdu/dPdv */ @@ -243,7 +243,7 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData sd->Ng = -sd->P; sd->I = -sd->P; sd->shader = kernel_data.background.shader; - sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK); + sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2); #ifdef __INSTANCING__ sd->object = ~0; @@ -275,8 +275,8 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData #ifdef __MULTI_CLOSURE__ -__device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf, - int skip_bsdf, float3 sum_eval, float sum_pdf, float sum_sample_weight) +__device_inline void _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf, + int skip_bsdf, BsdfEval *bsdf_eval, float sum_pdf, float sum_sample_weight) { for(int i = 0; i< sd->num_closure; i++) { if(i == skip_bsdf) @@ -293,7 +293,7 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float #endif if(bsdf_pdf != 0.0f) { - sum_eval += eval*sc->weight; + bsdf_eval_accum(bsdf_eval, sc->type, eval*sc->weight); sum_pdf += bsdf_pdf*sc->sample_weight; } @@ -302,25 +302,27 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float } *pdf = sum_pdf/sum_sample_weight; - return sum_eval; } #endif -__device float3 shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, - const float3 omega_in, float *pdf) +__device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd, + const float3 omega_in, BsdfEval *eval, float *pdf) { #ifdef __MULTI_CLOSURE__ - return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, make_float3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f); + bsdf_eval_init(eval, NBUILTIN_CLOSURES, make_float3(0.0f, 0.0f, 0.0f), kernel_data.film.use_light_pass); + + return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f); #else const ShaderClosure *sc = &sd->closure; + *pdf = 0.0f; - return svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight; + *eval = svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight; #endif } __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, - float randu, float randv, float3 *eval, + float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf) { #ifdef __MULTI_CLOSURE__ @@ -359,27 +361,28 @@ __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd, const ShaderClosure *sc = &sd->closure[sampled]; int label; + float3 eval; *pdf = 0.0f; #ifdef __OSL__ - label = OSLShader::bsdf_sample(sd, sc, randu, randv, *eval, *omega_in, *domega_in, *pdf); + label = OSLShader::bsdf_sample(sd, sc, randu, randv, eval, *omega_in, *domega_in, *pdf); #else - label = svm_bsdf_sample(sd, sc, randu, randv, eval, omega_in, domega_in, pdf); + label = svm_bsdf_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf); #endif - *eval *= sc->weight; + bsdf_eval_init(bsdf_eval, sc->type, eval*sc->weight, kernel_data.film.use_light_pass); if(sd->num_closure > 1 && *pdf != 0.0f) { float sweight = sc->sample_weight; - *eval = _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, *eval, *pdf*sweight, sweight); + _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, bsdf_eval, *pdf*sweight, sweight); } return label; #else /* sample the single closure that we picked */ *pdf = 0.0f; - int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, eval, omega_in, domega_in, pdf); - *eval *= sd->closure.weight; + int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, bsdf_eval, omega_in, domega_in, pdf); + *bsdf_eval *= sd->closure.weight; return label; #endif } @@ -421,6 +424,68 @@ __device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd) #endif } +__device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_DIFFUSE(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_DIFFUSE(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_GLOSSY(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_GLOSSY(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} + +__device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd) +{ +#ifdef __MULTI_CLOSURE__ + float3 eval = make_float3(0.0f, 0.0f, 0.0f); + + for(int i = 0; i< sd->num_closure; i++) { + ShaderClosure *sc = &sd->closure[i]; + + if(CLOSURE_IS_BSDF_TRANSMISSION(sc->type)) + eval += sc->weight; + } + + return eval; +#else + if(CLOSURE_IS_BSDF_TRANSMISSION(sd->closure.type)) + return sd->closure.weight; + else + return make_float3(0.0f, 0.0f, 0.0f); +#endif +} /* Emission */ @@ -588,12 +653,17 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect) int prim = kernel_tex_fetch(__prim_index, isect->prim); float4 Ns = kernel_tex_fetch(__tri_normal, prim); int shader = __float_as_int(Ns.w); - int flag = kernel_tex_fetch(__shader_flag, shader & SHADER_MASK); + int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)*2); return (flag & SD_HAS_SURFACE_TRANSPARENT) != 0; } #endif +__device int shader_pass_id(KernelGlobals *kg, ShaderData *sd) +{ + return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2 + 1); +} + /* Free ShaderData */ __device void shader_release(KernelGlobals *kg, ShaderData *sd) diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index 008ec0bdf28..b4b1da83162 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -70,6 +70,9 @@ CCL_NAMESPACE_BEGIN #ifdef __KERNEL_ADV_SHADING__ #define __MULTI_CLOSURE__ #define __TRANSPARENT_SHADOWS__ +#ifdef __KERNEL_CPU__ +#define __PASSES__ +#endif #endif //#define __MULTI_LIGHT__ @@ -150,6 +153,75 @@ typedef enum ClosureLabel { LABEL_STOP = 2048 } ClosureLabel; +/* Render Passes */ + +typedef enum PassType { + PASS_NONE = 0, + PASS_COMBINED = 1, + PASS_DEPTH = 2, + PASS_NORMAL = 8, + PASS_UV = 16, + PASS_OBJECT_ID = 32, + PASS_MATERIAL_ID = 64, + PASS_DIFFUSE_COLOR = 128, + PASS_GLOSSY_COLOR = 256, + PASS_TRANSMISSION_COLOR = 512, + PASS_DIFFUSE_INDIRECT = 1024, + PASS_GLOSSY_INDIRECT = 2048, + PASS_TRANSMISSION_INDIRECT = 4096, + PASS_DIFFUSE_DIRECT = 8192, + PASS_GLOSSY_DIRECT = 16384, + PASS_TRANSMISSION_DIRECT = 32768, + PASS_EMISSION = 65536, + PASS_BACKGROUND = 131072 +} PassType; + +#define PASS_ALL (~0) + +#ifdef __PASSES__ + +typedef float3 PathThroughput; + +struct PathRadiance { + int use_light_pass; + + float3 indirect; + float3 direct_throughput; + float3 direct_emission; + + float3 color_diffuse; + float3 color_glossy; + float3 color_transmission; + + float3 direct_diffuse; + float3 direct_glossy; + float3 direct_transmission; + + float3 indirect_diffuse; + float3 indirect_glossy; + float3 indirect_transmission; + + float3 emission; + float3 background; +}; + +struct BsdfEval { + int use_light_pass; + + float3 diffuse; + float3 glossy; + float3 transmission; + float3 transparent; +}; + +#else + +typedef float3 PathThroughput; +typedef float3 PathRadiance; +typedef float3 BsdfEval; + +#endif + /* Shader Flag */ typedef enum ShaderFlag { @@ -353,7 +425,32 @@ typedef struct KernelCamera { typedef struct KernelFilm { float exposure; - int pad1, pad2, pad3; + int pass_flag; + int pass_stride; + int use_light_pass; + + int pass_combined; + int pass_depth; + int pass_normal; + int pass_pad; + + int pass_uv; + int pass_object_id; + int pass_material_id; + int pass_diffuse_color; + + int pass_glossy_color; + int pass_transmission_color; + int pass_diffuse_indirect; + int pass_glossy_indirect; + + int pass_transmission_indirect; + int pass_diffuse_direct; + int pass_glossy_direct; + int pass_transmission_direct; + + int pass_emission; + int pass_background; } KernelFilm; typedef struct KernelBackground { diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h index ff6e9b94a0f..533a2944557 100644 --- a/intern/cycles/kernel/svm/svm_types.h +++ b/intern/cycles/kernel/svm/svm_types.h @@ -266,22 +266,26 @@ typedef enum ShaderType { typedef enum ClosureType { CLOSURE_BSDF_ID, + CLOSURE_BSDF_DIFFUSE_ID, CLOSURE_BSDF_OREN_NAYAR_ID, - CLOSURE_BSDF_TRANSLUCENT_ID, + CLOSURE_BSDF_REFLECTION_ID, - CLOSURE_BSDF_REFRACTION_ID, - CLOSURE_BSDF_GLASS_ID, - CLOSURE_BSDF_TRANSPARENT_ID, CLOSURE_BSDF_MICROFACET_GGX_ID, - CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID, CLOSURE_BSDF_MICROFACET_BECKMANN_ID, - CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID, CLOSURE_BSDF_WARD_ID, - CLOSURE_BSDF_ASHIKHMIN_VELVET_ID, - CLOSURE_BSDF_WESTIN_BACKSCATTER_ID, CLOSURE_BSDF_WESTIN_SHEEN_ID, + CLOSURE_BSDF_TRANSLUCENT_ID, + CLOSURE_BSDF_REFRACTION_ID, + CLOSURE_BSDF_WESTIN_BACKSCATTER_ID, + CLOSURE_BSDF_ASHIKHMIN_VELVET_ID, + CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID, + CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID, + CLOSURE_BSDF_GLASS_ID, + + CLOSURE_BSDF_TRANSPARENT_ID, + CLOSURE_BSSRDF_CUBIC_ID, CLOSURE_EMISSION_ID, CLOSURE_DEBUG_ID, @@ -297,7 +301,10 @@ typedef enum ClosureType { } ClosureType; /* watch this, being lazy with memory usage */ -#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_WESTIN_SHEEN_ID) +#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_TRANSPARENT_ID) +#define CLOSURE_IS_BSDF_DIFFUSE(type) (type >= CLOSURE_BSDF_DIFFUSE_ID && type <= CLOSURE_BSDF_OREN_NAYAR_ID) +#define CLOSURE_IS_BSDF_GLOSSY(type) (type >= CLOSURE_BSDF_REFLECTION_ID && type <= CLOSURE_BSDF_WESTIN_SHEEN_ID) +#define CLOSURE_IS_BSDF_TRANSMISSION(type) (type >= CLOSURE_BSDF_TRANSLUCENT_ID && type <= CLOSURE_BSDF_GLASS_ID) #define CLOSURE_IS_VOLUME(type) (type >= CLOSURE_VOLUME_ID && type <= CLOSURE_VOLUME_ISOTROPIC_ID) #define CLOSURE_IS_EMISSION(type) (type == CLOSURE_EMISSION_ID) #define CLOSURE_IS_HOLDOUT(type) (type == CLOSURE_HOLDOUT_ID) |