Sample as Lamp option for world shaders, to enable multiple importance sampling.
authorBrecht Van Lommel <brechtvanlommel@pandora.be>
Fri, 20 Jan 2012 17:49:17 +0000 (17:49 +0000)
committerBrecht Van Lommel <brechtvanlommel@pandora.be>
Fri, 20 Jan 2012 17:49:17 +0000 (17:49 +0000)
By default lighting from the world is computed solely with indirect light
sampling. However for more complex environment maps this can be too noisy, as
sampling the BSDF may not easily find the highlights in the environment map
image. By enabling this option, the world background will be sampled as a lamp,
with lighter parts automatically given more samples.

Map Resolution specifies the size of the importance map (res x res). Before
rendering starts, an importance map is generated by "baking" a grayscale image
from the world shader. This will then be used to determine which parts of the
background are light and so should receive more samples than darker parts.
Higher resolutions will result in more accurate sampling but take more setup
time and memory.

Patch by Mike Farnsworth, thanks!

26 files changed:
intern/cycles/blender/addon/properties.py
intern/cycles/blender/addon/ui.py
intern/cycles/blender/blender_object.cpp
intern/cycles/blender/blender_shader.cpp
intern/cycles/blender/blender_sync.h
intern/cycles/device/device_cpu.cpp
intern/cycles/kernel/kernel.cl
intern/cycles/kernel/kernel.cpp
intern/cycles/kernel/kernel.cu
intern/cycles/kernel/kernel.h
intern/cycles/kernel/kernel_compat_cpu.h
intern/cycles/kernel/kernel_compat_cuda.h
intern/cycles/kernel/kernel_differential.h
intern/cycles/kernel/kernel_displace.h
intern/cycles/kernel/kernel_emission.h
intern/cycles/kernel/kernel_light.h
intern/cycles/kernel/kernel_montecarlo.h
intern/cycles/kernel/kernel_optimized.cpp
intern/cycles/kernel/kernel_path.h
intern/cycles/kernel/kernel_textures.h
intern/cycles/kernel/kernel_types.h
intern/cycles/kernel/svm/svm_image.h
intern/cycles/render/light.cpp
intern/cycles/render/light.h
intern/cycles/render/mesh_displace.cpp
intern/cycles/render/scene.h

index 0a3cffd5071a46ec227fe4739129c33939398236..46a49f52de0128a5830e2f27f343632cd996394c 100644 (file)
@@ -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):
index 70f38fa7e8ccc97235143c4c674a655c4f092c48..67ff79a20379de3b7ba3029c098de00605298420 100644 (file)
@@ -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"
index 608cb33eadde8bf694ff4f587732c1e39ac6b8de..bc7868e01921d0feb54e3fc21a6703bb3cc4d95e 100644 (file)
  * 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);
index f5ffc321476dd475ed80b52aa0a9ab48e97e1e4f..1ce134f309470b7782dee3fb9ae98fb818377d46 100644 (file)
@@ -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 */
index 7b65376bd846eb46d44d253a89e7ec2207db1810..5e76a0a0b1e858a349530504c5b4c228e7162e79 100644 (file)
@@ -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);
index f4b2b7a8269373b059e5fec64fb141d673b8d1f7..25da978edd8f053c870f582435d7b622caf603ca 100644 (file)
@@ -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;
index 479cf9b2e64c692c24d1657fe04b41c1f36fa93d..305d81339c2ada2151e0f49388983903fc306018 100644 (file)
@@ -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);
 
index e66ddd86cd65826dbc1f2fdc28059e93eebd8814..9e0d252772b9bb32202d6de6d237baed306bd4ea 100644 (file)
@@ -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);
 }
index c97aeb6754807eeb2f026eaa6d882c1116be517c..4e585fd563db8671c28e385f36277bf0a6f6f74e 100644 (file)
@@ -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;
 
index 20d43c911693cf8c909f46aa411b7df6060c06d1..df6b5ee92dab7596824e971be8c8a4edb500719c 100644 (file)
@@ -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
 
index 783ae5198455544c98f3ef9b7139007a7b235b01..79f894bfdace5c895a33e8b5557ebae1cab2259d 100644 (file)
@@ -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;
index 40129a2f68f5e7cf831771d28a10fd40a5a1d59c..cc719bfadbc87b33fab4da1947547acb9e57b6a8 100644 (file)
@@ -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;
index 4e2b1ea7d139148ea188ff05bc31483c7e9df2ad..5b4290a77226fb6540e5ed9d9e5cc0a4a0429c01 100644 (file)
@@ -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);
 
index c39e5e43dbbce4baef44668af0818874fc4fe13b..73666892cf3d7e6e4722b77334ea6b859ffbdf0c 100644 (file)
@@ -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
index b81db721eb3cf05695cc9599187a01d7ebff1a02..51698f3a9bd428bd4611d50737bd5096b2b87333 100644 (file)
@@ -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
 
index d5d47b28d59858b2b718b38b05fdfab5f1da9e16..4c2b69c27163e7d24e2e0e14907fb808f3ea492d 100644 (file)
@@ -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)
index df291b66b231746c193442c351efa953fed688ce..9776baf65e4dc984fc477fd2da7604a95e7024ba 100644 (file)
@@ -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__ */
index c437e06adfaaa1d4780a44718e9257a085ac6c07..50341021d9d8416c4fff7e4da87fbd1396473457 100644 (file)
@@ -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);
 }
index d27ad861c6add93c300574dfaeaa925eaddc6e2a..c80d2068506d41f3fe37d4f8cb0e94baceedec1b 100644 (file)
@@ -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
 
index 2bfb1b3b88e4470a9a2f1f4fef968ee683cc4719..ca7ae432efa314a596efe4440e428b7f5bf07dc2 100644 (file)
@@ -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)
index 2c03a34df1fe4e57fd032fd91d95a7309c9a9194..008ec0bdf2842a6aab155fb4fc3a0642affbcd7f 100644 (file)
@@ -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 */
index 62e24166970db8375f2214e5e31f0415ee31df98..073021bdd54f318b84aaf90ed6b6b1dacdcf7283 100644 (file)
@@ -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) {
index c8206365373e50ad8e83382ca53d8addeb574c22..eea5cfb013735b18d0c8c1fea7e742b83dc2ee1d 100644 (file)
 #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)
index 19cbcb55386c0760512a5db3ddf883605c66f838..0ed143f5ad1b688ca5b44b5ad0a92ea267ca8b89 100644 (file)
@@ -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
index c4f3e43bfba4e7ef109d6b894609cd8fc9c9c019..a6f8e3f6be88b3ff309c64028d68b0e2ab9fadc6 100644 (file)
@@ -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;
                        }
                }
        }
index 17bd7e20129642afa9965563e19a7d940091e5a1..4a5224f8a94f132c062a622002a0f111262eddb8 100644 (file)
@@ -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;