diff options
26 files changed, 505 insertions, 66 deletions
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index 0a3cffd5071..46a49f52de0 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -155,6 +155,10 @@ class CyclesWorldSettings(bpy.types.PropertyGroup): @classmethod def register(cls): bpy.types.World.cycles = PointerProperty(type=cls, name="Cycles World Settings", description="Cycles world settings") + cls.sample_as_light = BoolProperty(name="Sample as Lamp", description="Use direct light sampling for the environment, enabling for non-solid colors is recommended", + default=False) + cls.sample_map_resolution = IntProperty(name="Map Resolution", description="Importance map size is resolution x resolution; higher values potentially produce less noise, at the cost of memory and speed", + default=256, min=4, max=8096) @classmethod def unregister(cls): diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py index 70f38fa7e8c..67ff79a2037 100644 --- a/intern/cycles/blender/addon/ui.py +++ b/intern/cycles/blender/addon/ui.py @@ -453,10 +453,38 @@ class CyclesWorld_PT_surface(CyclesButtonsPanel, Panel): layout = self.layout world = context.world + if not panel_node_draw(layout, world, 'OUTPUT_WORLD', 'Surface'): layout.prop(world, "horizon_color", text="Color") +class CyclesWorld_PT_settings(CyclesButtonsPanel, Panel): + bl_label = "Settings" + bl_context = "world" + bl_options = {'DEFAULT_CLOSED'} + + @classmethod + def poll(cls, context): + return context.world and CyclesButtonsPanel.poll(context) + + def draw(self, context): + layout = self.layout + + world = context.world + cworld = world.cycles + + split = layout.split() + col = split.column() + + col.prop(cworld, "sample_as_light") + row = col.row() + row.active = cworld.sample_as_light + row.prop(cworld, "sample_map_resolution") + + col = split.column() + col.label() + + class CyclesWorld_PT_volume(CyclesButtonsPanel, Panel): bl_label = "Volume" bl_context = "world" diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp index 608cb33eadd..bc7868e0192 100644 --- a/intern/cycles/blender/blender_object.cpp +++ b/intern/cycles/blender/blender_object.cpp @@ -16,10 +16,13 @@ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. */ +#include "graph.h" #include "light.h" #include "mesh.h" #include "object.h" #include "scene.h" +#include "nodes.h" +#include "shader.h" #include "blender_sync.h" #include "blender_util.h" @@ -152,6 +155,35 @@ void BlenderSync::sync_light(BL::Object b_parent, int b_index, BL::Object b_ob, light->tag_update(scene); } +void BlenderSync::sync_background_light() +{ + BL::World b_world = b_scene.world(); + + PointerRNA cworld = RNA_pointer_get(&b_world.ptr, "cycles"); + bool sample_as_light = get_boolean(cworld, "sample_as_light"); + + if(sample_as_light) { + /* test if we need to sync */ + Light *light; + ObjectKey key(b_world, 0, b_world); + + if(light_map.sync(&light, b_world, b_world, key) || + world_recalc || + b_world.ptr.data != world_map) + { + light->type = LIGHT_BACKGROUND; + light->map_resolution = get_int(cworld, "sample_map_resolution"); + light->shader = scene->default_background; + + light->tag_update(scene); + light_map.set_recalc(b_world); + } + } + + world_map = b_world.ptr.data; + world_recalc = false; +} + /* Object */ void BlenderSync::sync_object(BL::Object b_parent, int b_index, BL::Object b_ob, Transform& tfm, uint layer_flag) @@ -263,6 +295,8 @@ void BlenderSync::sync_objects(BL::SpaceView3D b_v3d) } } + sync_background_light(); + /* handle removed data and modified pointers */ if(light_map.post_sync()) scene->light_manager->tag_update(scene); diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp index f5ffc321476..1ce134f3094 100644 --- a/intern/cycles/blender/blender_shader.cpp +++ b/intern/cycles/blender/blender_shader.cpp @@ -700,9 +700,6 @@ void BlenderSync::sync_world() if(background->modified(prevbackground)) background->tag_update(scene); - - world_map = b_world.ptr.data; - world_recalc = false; } /* Sync Lamps */ diff --git a/intern/cycles/blender/blender_sync.h b/intern/cycles/blender/blender_sync.h index 7b65376bd84..5e76a0a0b1e 100644 --- a/intern/cycles/blender/blender_sync.h +++ b/intern/cycles/blender/blender_sync.h @@ -80,6 +80,7 @@ private: Mesh *sync_mesh(BL::Object b_ob, bool object_updated); void sync_object(BL::Object b_parent, int b_index, BL::Object b_object, Transform& tfm, uint layer_flag); void sync_light(BL::Object b_parent, int b_index, BL::Object b_ob, Transform& tfm); + void sync_background_light(); /* util */ void find_shader(BL::ID id, vector<uint>& used_shaders, int default_shader); diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index f4b2b7a8269..25da978edd8 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -217,7 +217,7 @@ public: #ifdef WITH_OPTIMIZED_KERNEL if(system_cpu_support_optimized()) { for(int x = task.shader_x; x < task.shader_x + task.shader_w; x++) { - kernel_cpu_optimized_shader(kg, (uint4*)task.shader_input, (float3*)task.shader_output, task.shader_eval_type, x); + kernel_cpu_optimized_shader(kg, (uint4*)task.shader_input, (float4*)task.shader_output, task.shader_eval_type, x); if(tasks.worker_cancel()) break; @@ -227,7 +227,7 @@ public: #endif { for(int x = task.shader_x; x < task.shader_x + task.shader_w; x++) { - kernel_cpu_shader(kg, (uint4*)task.shader_input, (float3*)task.shader_output, task.shader_eval_type, x); + kernel_cpu_shader(kg, (uint4*)task.shader_input, (float4*)task.shader_output, task.shader_eval_type, x); if(tasks.worker_cancel()) break; diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 479cf9b2e64..305d81339c2 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -80,7 +80,7 @@ __kernel void kernel_ocl_tonemap( kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride); } -/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float3 *output, int type, int sx) +/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float *output, int type, int sx) { int x = sx + get_global_id(0); diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index e66ddd86cd6..9e0d252772b 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -218,7 +218,7 @@ void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sam /* Shader Evaluation */ -void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float3 *output, int type, int i) +void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index c97aeb67548..4e585fd563d 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -44,7 +44,7 @@ extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x, y, offset, stride); } -extern "C" __global__ void kernel_cuda_shader(uint4 *input, float3 *output, int type, int sx) +extern "C" __global__ void kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h index 20d43c91169..df6b5ee92da 100644 --- a/intern/cycles/kernel/kernel.h +++ b/intern/cycles/kernel/kernel.h @@ -40,7 +40,7 @@ void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_ int sample, int x, int y, int offset, int stride); 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_shader(KernelGlobals *kg, uint4 *input, float3 *output, +void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); #ifdef WITH_OPTIMIZED_KERNEL @@ -48,7 +48,7 @@ void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int sample, int x, int y, int offset, int stride); 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_shader(KernelGlobals *kg, uint4 *input, float3 *output, +void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i); #endif diff --git a/intern/cycles/kernel/kernel_compat_cpu.h b/intern/cycles/kernel/kernel_compat_cpu.h index 783ae519845..79f894bfdac 100644 --- a/intern/cycles/kernel/kernel_compat_cpu.h +++ b/intern/cycles/kernel/kernel_compat_cpu.h @@ -141,6 +141,7 @@ template<typename T> struct texture_image { }; typedef texture<float4> texture_float4; +typedef texture<float2> texture_float2; typedef texture<float> texture_float; typedef texture<uint> texture_uint; typedef texture<int> texture_int; diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index 40129a2f68f..cc719bfadbc 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -45,6 +45,7 @@ /* Textures */ typedef texture<float4, 1> texture_float4; +typedef texture<float2, 1> texture_float2; typedef texture<float, 1> texture_float; typedef texture<uint, 1> texture_uint; typedef texture<int, 1> texture_int; diff --git a/intern/cycles/kernel/kernel_differential.h b/intern/cycles/kernel/kernel_differential.h index 4e2b1ea7d13..5b4290a7722 100644 --- a/intern/cycles/kernel/kernel_differential.h +++ b/intern/cycles/kernel/kernel_differential.h @@ -71,8 +71,8 @@ __device void differential_dudv(differential *du, differential *dv, float3 dPdu, * and the same for dudy and dvdy. the denominator is the same for both * solutions, so we compute it only once. * - * dP.dx = dPdu * dudx + dPdv * dvdx; - * dP.dy = dPdu * dudy + dPdv * dvdy; */ + * dP.dx = dPdu * dudx + dPdv * dvdx; + * dP.dy = dPdu * dudy + dPdv * dvdy; */ float det = (dPdu.x*dPdv.y - dPdv.x*dPdu.y); diff --git a/intern/cycles/kernel/kernel_displace.h b/intern/cycles/kernel/kernel_displace.h index c39e5e43dbb..73666892cf3 100644 --- a/intern/cycles/kernel/kernel_displace.h +++ b/intern/cycles/kernel/kernel_displace.h @@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN -__device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float3 *output, ShaderEvalType type, int i) +__device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float4 *output, ShaderEvalType type, int i) { ShaderData sd; uint4 in = input[i]; @@ -62,7 +62,7 @@ __device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float3 *ou } /* write output */ - output[i] = out; + output[i] = make_float4(out.x, out.y, out.z, 0.0f); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index b81db721eb3..51698f3a9bd 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -25,21 +25,31 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando, { /* setup shading at emitter */ ShaderData sd; - - shader_setup_from_sample(kg, &sd, ls->P, ls->Ng, I, ls->shader, ls->object, ls->prim, u, v); - ls->Ng = sd.Ng; - - /* no path flag, we're evaluating this for all closures. that's weak but - we'd have to do multiple evaluations otherwise */ - shader_eval_surface(kg, &sd, rando, 0); - float3 eval; - /* evaluate emissive closure */ - if(sd.flag & SD_EMISSION) - eval = shader_emissive_eval(kg, &sd); - else - eval = make_float3(0.0f, 0.0f, 0.0f); + if(ls->type == LIGHT_BACKGROUND) { + Ray ray; + ray.D = ls->D; + ray.P = ls->P; + ray.dP.dx = make_float3(0.0f, 0.0f, 0.0f); + ray.dP.dy = make_float3(0.0f, 0.0f, 0.0f); + shader_setup_from_background(kg, &sd, &ray); + eval = shader_eval_background(kg, &sd, 0); + } + else { + shader_setup_from_sample(kg, &sd, ls->P, ls->Ng, I, ls->shader, ls->object, ls->prim, u, v); + ls->Ng = sd.Ng; + + /* no path flag, we're evaluating this for all closures. that's weak but + we'd have to do multiple evaluations otherwise */ + shader_eval_surface(kg, &sd, rando, 0); + + /* evaluate emissive closure */ + if(sd.flag & SD_EMISSION) + eval = shader_emissive_eval(kg, &sd); + else + eval = make_float3(0.0f, 0.0f, 0.0f); + } shader_release(kg, &sd); @@ -51,25 +61,31 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, { LightSample ls; + float pdf = -1.0f; + #ifdef __MULTI_LIGHT__ if(lindex != -1) { /* sample position on a specified light */ - light_select(kg, lindex, randu, randv, sd->P, &ls); + light_select(kg, lindex, randu, randv, sd->P, &ls, &pdf); } else #endif { /* sample a light and position on int */ - light_sample(kg, randt, randu, randv, sd->P, &ls); + light_sample(kg, randt, randu, randv, sd->P, &ls, &pdf); } /* compute pdf */ - float pdf = light_sample_pdf(kg, &ls, -ls.D, ls.t); + if(pdf < 0.0f) + pdf = light_sample_pdf(kg, &ls, -ls.D, ls.t); + + if(pdf == 0.0f) + return false; /* evaluate closure */ *eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D); - if(is_zero(*eval) || pdf == 0.0f) + if(is_zero(*eval)) return false; /* todo: use visbility flag to skip lights */ @@ -83,7 +99,7 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex, if(is_zero(*eval)) return false; - if(ls.prim != ~0) { + if(ls.prim != ~0 || ls.type == LIGHT_BACKGROUND) { /* multiple importance sampling */ float mis_weight = power_heuristic(pdf, bsdf_pdf); *eval *= mis_weight; @@ -125,7 +141,8 @@ __device float3 indirect_emission(KernelGlobals *kg, ShaderData *sd, float t, in float3 L = shader_emissive_eval(kg, sd); if(!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_SAMPLE_AS_LIGHT)) { - /* multiple importance sampling */ + /* multiple importance sampling, get triangle light pdf, + and compute weight with respect to BSDF pdf */ float pdf = triangle_light_pdf(kg, sd->Ng, sd->I, t); float mis_weight = power_heuristic(bsdf_pdf, pdf); @@ -135,5 +152,34 @@ __device float3 indirect_emission(KernelGlobals *kg, ShaderData *sd, float t, in return L; } +/* Indirect Background */ + +__device float3 indirect_background(KernelGlobals *kg, Ray *ray, int path_flag, float bsdf_pdf) +{ +#ifdef __BACKGROUND__ + /* evaluate background closure */ + ShaderData sd; + shader_setup_from_background(kg, &sd, ray); + float3 L = shader_eval_background(kg, &sd, path_flag); + shader_release(kg, &sd); + + /* check if background light exists or if we should skip pdf */ + int res = kernel_data.integrator.pdf_background_res; + + if(!(path_flag & PATH_RAY_MIS_SKIP) && res) { + /* multiple importance sampling, get background light pdf for ray + direction, and compute weight with respect to BSDF pdf */ + float pdf = background_light_pdf(kg, ray->D); + float mis_weight = power_heuristic(bsdf_pdf, pdf); + + return L*mis_weight; + } + + return L; +#else + return make_float3(0.8f, 0.8f, 0.8f); +#endif +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h index d5d47b28d59..4c2b69c2716 100644 --- a/intern/cycles/kernel/kernel_light.h +++ b/intern/cycles/kernel/kernel_light.h @@ -26,6 +26,7 @@ typedef struct LightSample { int object; int prim; int shader; + LightType type; } LightSample; /* Regular Light */ @@ -58,13 +59,125 @@ __device float3 area_light_sample(float3 axisu, float3 axisv, float randu, float return axisu*randu + axisv*randv; } +__device float3 background_light_sample(KernelGlobals *kg, float randu, float randv, float *pdf) +{ + /* for the following, the CDF values are actually a pair of floats, with the + function value as X and the actual CDF as Y. The last entry's function + value is the CDF total. */ + int res = kernel_data.integrator.pdf_background_res; + int cdf_count = res + 1; + + /* this is basically std::lower_bound as used by pbrt */ + int first = 0; + int count = res; + + while(count > 0) { + int step = count >> 1; + int middle = first + step; + + if(kernel_tex_fetch(__light_background_marginal_cdf, middle).y < randv) { + first = middle + 1; + count -= step + 1; + } + else + count = step; + } + + int index_v = max(0, first - 1); + kernel_assert(index_v >= 0 && index_v < res); + + float2 cdf_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v); + float2 cdf_next_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v + 1); + float2 cdf_last_v = kernel_tex_fetch(__light_background_marginal_cdf, res); + + /* importance-sampled V direction */ + float dv = (randv - cdf_v.y) / (cdf_next_v.y - cdf_v.y); + float v = (index_v + dv) / res; + + /* this is basically std::lower_bound as used by pbrt */ + first = 0; + count = res; + while(count > 0) { + int step = count >> 1; + int middle = first + step; + + if(kernel_tex_fetch(__light_background_conditional_cdf, index_v * cdf_count + middle).y < randu) { + first = middle + 1; + count -= step + 1; + } + else + count = step; + } + + int index_u = max(0, first - 1); + kernel_assert(index_u >= 0 && index_u < res); + + float2 cdf_u = kernel_tex_fetch(__light_background_conditional_cdf, index_v * cdf_count + index_u); + float2 cdf_next_u = kernel_tex_fetch(__light_background_conditional_cdf, index_v * cdf_count + index_u + 1); + float2 cdf_last_u = kernel_tex_fetch(__light_background_conditional_cdf, index_v * cdf_count + res); + + /* importance-sampled U direction */ + float du = (randu - cdf_u.y) / (cdf_next_u.y - cdf_u.y); + float u = (index_u + du) / res; + + /* spherical coordinates */ + float theta = v * M_PI_F; + float phi = u * M_PI_F * 2.0f; + + /* compute pdf */ + float denom = cdf_last_u.x * cdf_last_v.x; + float sin_theta = sinf(theta); + + if(sin_theta == 0.0f || denom == 0.0f) + *pdf = 0.0f; + else + *pdf = (cdf_u.x * cdf_v.x)/(2.0f * M_PI_F * M_PI_F * sin_theta * denom); + + *pdf *= kernel_data.integrator.pdf_lights; + + /* compute direction */ + return spherical_to_direction(theta, phi); +} + +__device float background_light_pdf(KernelGlobals *kg, float3 direction) +{ + float2 uv = direction_to_equirectangular(direction); + int res = kernel_data.integrator.pdf_background_res; + + float sin_theta = sinf(uv.y * M_PI_F); + + if(sin_theta == 0.0f) + return 0.0f; + + int index_u = clamp((int)(uv.x * res), 0, res - 1); + int index_v = clamp((int)(uv.y * res), 0, res - 1); + + /* pdfs in V direction */ + float2 cdf_last_u = kernel_tex_fetch(__light_background_conditional_cdf, index_v * (res + 1) + res); + float2 cdf_last_v = kernel_tex_fetch(__light_background_marginal_cdf, res); + + float denom = cdf_last_u.x * cdf_last_v.x; + + if(denom == 0.0f) + return 0.0f; + + /* pdfs in U direction */ + float2 cdf_u = kernel_tex_fetch(__light_background_conditional_cdf, index_v * (res + 1) + index_u); + float2 cdf_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v); + + float pdf = (cdf_u.x * cdf_v.x)/(2.0f * M_PI_F * M_PI_F * sin_theta * denom); + + return pdf * kernel_data.integrator.pdf_lights; +} + __device void regular_light_sample(KernelGlobals *kg, int point, - float randu, float randv, float3 P, LightSample *ls) + float randu, float randv, float3 P, LightSample *ls, float *pdf) { float4 data0 = kernel_tex_fetch(__light_data, point*LIGHT_SIZE + 0); float4 data1 = kernel_tex_fetch(__light_data, point*LIGHT_SIZE + 1); LightType type = (LightType)__float_as_int(data0.x); + ls->type = type; if(type == LIGHT_DISTANT) { /* distant light */ @@ -79,6 +192,15 @@ __device void regular_light_sample(KernelGlobals *kg, int point, ls->D = -D; ls->t = FLT_MAX; } + else if(type == LIGHT_BACKGROUND) { + /* infinite area light (e.g. light dome or env light) */ + float3 D = background_light_sample(kg, randu, randv, pdf); + + ls->P = D; + ls->Ng = D; + ls->D = -D; + ls->t = FLT_MAX; + } else { ls->P = make_float3(data0.y, data0.z, data0.w); @@ -139,6 +261,7 @@ __device void triangle_light_sample(KernelGlobals *kg, int prim, int object, ls->object = object; ls->prim = prim; ls->t = 0.0f; + ls->type = LIGHT_AREA; #ifdef __INSTANCING__ /* instance transform */ @@ -192,7 +315,7 @@ __device int light_distribution_sample(KernelGlobals *kg, float randt) /* Generic Light */ -__device void light_sample(KernelGlobals *kg, float randt, float randu, float randv, float3 P, LightSample *ls) +__device void light_sample(KernelGlobals *kg, float randt, float randu, float randv, float3 P, LightSample *ls, float *pdf) { /* sample index */ int index = light_distribution_sample(kg, randt); @@ -207,7 +330,7 @@ __device void light_sample(KernelGlobals *kg, float randt, float randu, float ra } else { int point = -prim-1; - regular_light_sample(kg, point, randu, randv, P, ls); + regular_light_sample(kg, point, randu, randv, P, ls, pdf); } /* compute incoming direction and distance */ @@ -227,9 +350,9 @@ __device float light_sample_pdf(KernelGlobals *kg, LightSample *ls, float3 I, fl return pdf; } -__device void light_select(KernelGlobals *kg, int index, float randu, float randv, float3 P, LightSample *ls) +__device void light_select(KernelGlobals *kg, int index, float randu, float randv, float3 P, LightSample *ls, float *pdf) { - regular_light_sample(kg, index, randu, randv, P, ls); + regular_light_sample(kg, index, randu, randv, P, ls, pdf); } __device float light_select_pdf(KernelGlobals *kg, LightSample *ls, float3 I, float t) diff --git a/intern/cycles/kernel/kernel_montecarlo.h b/intern/cycles/kernel/kernel_montecarlo.h index df291b66b23..9776baf65e4 100644 --- a/intern/cycles/kernel/kernel_montecarlo.h +++ b/intern/cycles/kernel/kernel_montecarlo.h @@ -104,13 +104,13 @@ __device_inline void sample_uniform_hemisphere(const float3 N, __device float3 sample_uniform_sphere(float u1, float u2) { - float z = 1.0f - 2.0f*u1; - float r = sqrtf(fmaxf(0.0f, 1.0f - z*z)); - float phi = 2.0f*M_PI_F*u2; - float x = r*cosf(phi); - float y = r*sinf(phi); + float z = 1.0f - 2.0f*u1; + float r = sqrtf(fmaxf(0.0f, 1.0f - z*z)); + float phi = 2.0f*M_PI_F*u2; + float x = r*cosf(phi); + float y = r*sinf(phi); - return make_float3(x, y, z); + return make_float3(x, y, z); } __device float power_heuristic(float a, float b) @@ -203,6 +203,28 @@ __device float3 spherical_to_direction(float theta, float phi) cosf(theta)); } +/* Equirectangular */ + +__device float2 direction_to_equirectangular(float3 dir) +{ + float u = (atan2f(dir.y, dir.x) + M_PI_F)/(2.0f*M_PI_F); + float v = atan2f(dir.z, hypotf(dir.x, dir.y))/M_PI_F + 0.5f; + + return make_float2(u, v); +} + +__device float3 equirectangular_to_direction(float u, float v) +{ + /* XXX check correctness? */ + float theta = M_PI_F*v; + float phi = 2.0f*M_PI_F*u; + + return make_float3( + sin(theta)*cos(phi), + sin(theta)*sin(phi), + cos(theta)); +} + CCL_NAMESPACE_END #endif /* __KERNEL_MONTECARLO_CL__ */ diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp index c437e06adfa..50341021d9d 100644 --- a/intern/cycles/kernel/kernel_optimized.cpp +++ b/intern/cycles/kernel/kernel_optimized.cpp @@ -49,7 +49,7 @@ void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffe /* Shader Evaluate */ -void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float3 *output, int type, int i) +void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index d27ad861c6a..c80d2068506 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -260,14 +260,9 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R Ltransparent += average(throughput); } else { -#ifdef __BACKGROUND__ - ShaderData sd; - shader_setup_from_background(kg, &sd, &ray); - L += throughput*shader_eval_background(kg, &sd, state.flag); - shader_release(kg, &sd); -#else - L += throughput*make_float3(0.8f, 0.8f, 0.8f); -#endif + /* sample background shader */ + float3 background_L = indirect_background(kg, &ray, state.flag, ray_pdf); + L += throughput*background_L; } break; @@ -362,7 +357,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R throughput *= bsdf_eval/bsdf_pdf; /* set labels */ -#ifdef __EMISSION__ +#if defined(__EMISSION__) || defined(__BACKGROUND__) ray_pdf = bsdf_pdf; #endif diff --git a/intern/cycles/kernel/kernel_textures.h b/intern/cycles/kernel/kernel_textures.h index 2bfb1b3b88e..ca7ae432efa 100644 --- a/intern/cycles/kernel/kernel_textures.h +++ b/intern/cycles/kernel/kernel_textures.h @@ -33,6 +33,8 @@ KERNEL_TEX(float4, texture_float4, __attributes_float3) /* lights */ KERNEL_TEX(float4, texture_float4, __light_distribution) KERNEL_TEX(float4, texture_float4, __light_data) +KERNEL_TEX(float2, texture_float2, __light_background_marginal_cdf) +KERNEL_TEX(float2, texture_float2, __light_background_conditional_cdf) /* shaders */ KERNEL_TEX(uint4, texture_uint4, __svm_nodes) diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index 2c03a34df1f..008ec0bdf28 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -165,6 +165,7 @@ typedef enum ShaderFlag { typedef enum LightType { LIGHT_POINT, LIGHT_DISTANT, + LIGHT_BACKGROUND, LIGHT_AREA } LightType; @@ -379,18 +380,19 @@ typedef struct KernelIntegrator { int num_all_lights; float pdf_triangles; float pdf_lights; + int pdf_background_res; /* bounces */ int min_bounce; int max_bounce; - int max_diffuse_bounce; - int max_glossy_bounce; - int max_transmission_bounce; + int max_diffuse_bounce; + int max_glossy_bounce; + int max_transmission_bounce; /* transparent */ - int transparent_min_bounce; - int transparent_max_bounce; + int transparent_min_bounce; + int transparent_max_bounce; int transparent_shadows; /* caustics */ diff --git a/intern/cycles/kernel/svm/svm_image.h b/intern/cycles/kernel/svm/svm_image.h index 62e24166970..073021bdd54 100644 --- a/intern/cycles/kernel/svm/svm_image.h +++ b/intern/cycles/kernel/svm/svm_image.h @@ -175,9 +175,8 @@ __device void svm_node_tex_environment(KernelGlobals *kg, ShaderData *sd, float decode_node_uchar4(node.z, &co_offset, &out_offset, &alpha_offset, &srgb); float3 co = stack_load_float3(stack, co_offset); - float u = (atan2f(co.y, co.x) + M_PI_F)/(2*M_PI_F); - float v = atan2f(co.z, hypotf(co.x, co.y))/M_PI_F + 0.5f; - float4 f = svm_image_texture(kg, id, u, v); + float2 uv = direction_to_equirectangular(co); + float4 f = svm_image_texture(kg, id, uv.x, uv.y); float3 r = make_float3(f.x, f.y, f.z); if(srgb) { diff --git a/intern/cycles/render/light.cpp b/intern/cycles/render/light.cpp index c8206365373..eea5cfb0137 100644 --- a/intern/cycles/render/light.cpp +++ b/intern/cycles/render/light.cpp @@ -26,8 +26,74 @@ #include "util_foreach.h" #include "util_progress.h" +#include "kernel_montecarlo.h" + CCL_NAMESPACE_BEGIN +static void dump_background_pixels(Device *device, DeviceScene *dscene, int res, vector<float3>& pixels) +{ + /* create input */ + int width = res; + int height = res; + + device_vector<uint4> d_input; + device_vector<float4> d_output; + + uint4 *d_input_data = d_input.resize(width*height); + + for(int y = 0; y < height; y++) { + for(int x = 0; x < width; x++) { + float u = x/(float)width; + float v = y/(float)height; + float3 D = -equirectangular_to_direction(u, v); + + uint4 in = make_uint4(__float_as_int(D.x), __float_as_int(D.y), __float_as_int(D.z), 0); + d_input_data[x + y*width] = in; + } + } + + /* compute on device */ + float4 *d_output_data = d_output.resize(width*height); + memset((void*)d_output.data_pointer, 0, d_output.memory_size()); + + device->const_copy_to("__data", &dscene->data, sizeof(dscene->data)); + + device->mem_alloc(d_input, MEM_READ_ONLY); + device->mem_copy_to(d_input); + device->mem_alloc(d_output, MEM_WRITE_ONLY); + + DeviceTask main_task(DeviceTask::SHADER); + main_task.shader_input = d_input.device_pointer; + main_task.shader_output = d_output.device_pointer; + main_task.shader_eval_type = SHADER_EVAL_BACKGROUND; + main_task.shader_x = 0; + main_task.shader_w = width*height; + + list<DeviceTask> split_tasks; + main_task.split_max_size(split_tasks, 128*128); + + foreach(DeviceTask& task, split_tasks) { + device->task_add(task); + device->task_wait(); + } + + device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4)); + device->mem_free(d_input); + device->mem_free(d_output); + + d_output_data = reinterpret_cast<float4*>(d_output.data_pointer); + + pixels.resize(width*height); + + for(int y = 0; y < height; y++) { + for(int x = 0; x < width; x++) { + pixels[y*width + x].x = d_output_data[y*width + x].x; + pixels[y*width + x].y = d_output_data[y*width + x].y; + pixels[y*width + x].z = d_output_data[y*width + x].z; + } + } +} + /* Light */ Light::Light() @@ -44,6 +110,8 @@ Light::Light() axisv = make_float3(0.0f, 0.0f, 0.0f); sizev = 1.0f; + map_resolution = 512; + cast_shadow = true; shader = 0; } @@ -66,6 +134,8 @@ LightManager::~LightManager() void LightManager::device_update_distribution(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress) { + progress.set_status("Updating Lights", "Computing distribution"); + /* option to always sample all point lights */ bool multi_light = false; @@ -232,6 +302,99 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen dscene->light_distribution.clear(); } +void LightManager::device_update_background(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress) +{ + KernelIntegrator *kintegrator = &dscene->data.integrator; + Light *background_light = NULL; + + /* find background light */ + foreach(Light *light, scene->lights) { + if(light->type == LIGHT_BACKGROUND) { + background_light = light; + break; + } + } + + /* no background light found, signal renderer to skip sampling */ + if(!background_light) { + kintegrator->pdf_background_res = 0; + return; + } + + progress.set_status("Updating Lights", "Importance map"); + + assert(kintegrator->use_direct_light); + + /* get the resolution from the light's size (we stuff it in there) */ + int res = background_light->map_resolution; + kintegrator->pdf_background_res = res; + + assert(res > 0); + + vector<float3> pixels; + dump_background_pixels(device, dscene, res, pixels); + + if(progress.get_cancel()) + return; + + /* build row distributions and column distribution for the infinite area environment light */ + int cdf_count = res + 1; + float2 *marg_cdf = dscene->light_background_marginal_cdf.resize(cdf_count); + float2 *cond_cdf = dscene->light_background_conditional_cdf.resize(cdf_count * cdf_count); + + /* conditional CDFs (rows, U direction) */ + for(int i = 0; i < res; i++) { + float sin_theta = sinf(M_PI_F * (i + 0.5f) / res); + float3 env_color = pixels[i * res]; + float ave_luminamce = average(env_color); + + cond_cdf[i * cdf_count].x = ave_luminamce * sin_theta; + cond_cdf[i * cdf_count].y = 0.0f; + + for(int j = 1; j < res; j++) { + env_color = pixels[i * res + j]; + ave_luminamce = average(env_color); + + cond_cdf[i * cdf_count + j].x = ave_luminamce * sin_theta; + cond_cdf[i * cdf_count + j].y = cond_cdf[i * cdf_count + j - 1].y + cond_cdf[i * cdf_count + j - 1].x / res; + } + + float cdf_total = cond_cdf[i * cdf_count + res - 1].y + cond_cdf[i * cdf_count + res - 1].x / res; + + /* stuff the total into the brightness value for the last entry, because + we are going to normalize the CDFs to 0.0 to 1.0 afterwards */ + cond_cdf[i * cdf_count + res].x = cdf_total; + + if(cdf_total > 0.0f) + for(int j = 1; j < res; j++) + cond_cdf[i * cdf_count + j].y /= cdf_total; + + cond_cdf[i * cdf_count + res].y = 1.0f; + } + + /* marginal CDFs (column, V direction, sum of rows) */ + marg_cdf[0].x = cond_cdf[res].x; + marg_cdf[0].y = 0.0f; + + for(int i = 1; i < res; i++) { + marg_cdf[i].x = cond_cdf[i * cdf_count + res].x; + marg_cdf[i].y = marg_cdf[i - 1].y + marg_cdf[i - 1].x / res; + } + + float cdf_total = marg_cdf[res - 1].y + marg_cdf[res - 1].x / res; + marg_cdf[res].x = cdf_total; + + if(cdf_total > 0.0f) + for(int i = 1; i < res; i++) + marg_cdf[i].y /= cdf_total; + + marg_cdf[res].y = 1.0f; + + /* update device */ + device->tex_alloc("__light_background_marginal_cdf", dscene->light_background_marginal_cdf); + device->tex_alloc("__light_background_conditional_cdf", dscene->light_background_conditional_cdf); +} + void LightManager::device_update_points(Device *device, DeviceScene *dscene, Scene *scene) { if(scene->lights.size() == 0) @@ -264,6 +427,14 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce light_data[i*LIGHT_SIZE + 2] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); light_data[i*LIGHT_SIZE + 3] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); } + else if(light->type == LIGHT_BACKGROUND) { + shader_id &= ~SHADER_AREA_LIGHT; + + light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), 0.0f, 0.0f, 0.0f); + light_data[i*LIGHT_SIZE + 1] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); + light_data[i*LIGHT_SIZE + 2] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); + light_data[i*LIGHT_SIZE + 3] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); + } else if(light->type == LIGHT_AREA) { float3 axisu = light->axisu*(light->sizeu*light->size); float3 axisv = light->axisv*(light->sizev*light->size); @@ -291,6 +462,9 @@ void LightManager::device_update(Device *device, DeviceScene *dscene, Scene *sce device_update_distribution(device, dscene, scene, progress); if(progress.get_cancel()) return; + device_update_background(device, dscene, scene, progress); + if(progress.get_cancel()) return; + need_update = false; } @@ -298,9 +472,13 @@ void LightManager::device_free(Device *device, DeviceScene *dscene) { device->tex_free(dscene->light_distribution); device->tex_free(dscene->light_data); + device->tex_free(dscene->light_background_marginal_cdf); + device->tex_free(dscene->light_background_conditional_cdf); dscene->light_distribution.clear(); dscene->light_data.clear(); + dscene->light_background_marginal_cdf.clear(); + dscene->light_background_conditional_cdf.clear(); } void LightManager::tag_update(Scene *scene) diff --git a/intern/cycles/render/light.h b/intern/cycles/render/light.h index 19cbcb55386..0ed143f5ad1 100644 --- a/intern/cycles/render/light.h +++ b/intern/cycles/render/light.h @@ -46,6 +46,8 @@ public: float3 axisv; float sizev; + int map_resolution; + bool cast_shadow; int shader; @@ -68,6 +70,7 @@ public: protected: void device_update_points(Device *device, DeviceScene *dscene, Scene *scene); void device_update_distribution(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress); + void device_update_background(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress); }; CCL_NAMESPACE_END diff --git a/intern/cycles/render/mesh_displace.cpp b/intern/cycles/render/mesh_displace.cpp index c4f3e43bfba..a6f8e3f6be8 100644 --- a/intern/cycles/render/mesh_displace.cpp +++ b/intern/cycles/render/mesh_displace.cpp @@ -89,7 +89,7 @@ bool MeshManager::displace(Device *device, Scene *scene, Mesh *mesh, Progress& p return false; /* run device task */ - device_vector<float3> d_output; + device_vector<float4> d_output; d_output.resize(d_input.size()); device->mem_alloc(d_input, MEM_READ_ONLY); @@ -106,7 +106,7 @@ bool MeshManager::displace(Device *device, Scene *scene, Mesh *mesh, Progress& p device->task_add(task); device->task_wait(); - device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float3)); + device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4)); device->mem_free(d_input); device->mem_free(d_output); @@ -118,7 +118,7 @@ bool MeshManager::displace(Device *device, Scene *scene, Mesh *mesh, Progress& p done.resize(mesh->verts.size(), false); int k = 0; - float3 *offset = (float3*)d_output.data_pointer; + float4 *offset = (float4*)d_output.data_pointer; for(size_t i = 0; i < mesh->triangles.size(); i++) { Mesh::Triangle t = mesh->triangles[i]; @@ -130,7 +130,8 @@ bool MeshManager::displace(Device *device, Scene *scene, Mesh *mesh, Progress& p for(int j = 0; j < 3; j++) { if(!done[t.v[j]]) { done[t.v[j]] = true; - mesh->verts[t.v[j]] += offset[k++]; + float3 off = float4_to_float3(offset[k++]); + mesh->verts[t.v[j]] += off; } } } diff --git a/intern/cycles/render/scene.h b/intern/cycles/render/scene.h index 17bd7e20129..4a5224f8a94 100644 --- a/intern/cycles/render/scene.h +++ b/intern/cycles/render/scene.h @@ -78,6 +78,8 @@ public: /* lights */ device_vector<float4> light_distribution; device_vector<float4> light_data; + device_vector<float2> light_background_marginal_cdf; + device_vector<float2> light_background_conditional_cdf; /* shaders */ device_vector<uint4> svm_nodes; |