Sample as Lamp option for world shaders, to enable multiple importance sampling.

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, 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;