Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

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

Browse Source 
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!
This commit is contained in:
Brecht Van Lommel
2012-01-20 17:49:17 +00:00
parent bddc01a7e1
commit 5873301257
26 changed files with 505 additions and 66 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
intern/cycles/blender/addon/properties.py
View 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):

28
intern/cycles/blender/addon/ui.py
View 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"

34
intern/cycles/blender/blender_object.cpp
View File

@@ -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);

3
intern/cycles/blender/blender_shader.cpp
View 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 */

1
intern/cycles/blender/blender_sync.h
View 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);

4
intern/cycles/device/device_cpu.cpp
View 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;

2
intern/cycles/kernel/kernel.cl
View 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);

2
intern/cycles/kernel/kernel.cpp
View 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);
}

2
intern/cycles/kernel/kernel.cu
View 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;

4
intern/cycles/kernel/kernel.h
View 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

1
intern/cycles/kernel/kernel_compat_cpu.h
View 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;

1
intern/cycles/kernel/kernel_compat_cuda.h
View 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;

4
intern/cycles/kernel/kernel_differential.h
View 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);

4
intern/cycles/kernel/kernel_displace.h
View 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

84
intern/cycles/kernel/kernel_emission.h
View 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

133
intern/cycles/kernel/kernel_light.h
View 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)

34
intern/cycles/kernel/kernel_montecarlo.h
View 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__ */

2
intern/cycles/kernel/kernel_optimized.cpp
View 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);
}

13
intern/cycles/kernel/kernel_path.h
View 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

2
intern/cycles/kernel/kernel_textures.h
View 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)

12
intern/cycles/kernel/kernel_types.h
View 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 */

5
intern/cycles/kernel/svm/svm_image.h
View 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) {

178
intern/cycles/render/light.cpp
View File

@@ -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)

3
intern/cycles/render/light.h
View 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

9
intern/cycles/render/mesh_displace.cpp
View 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;
}
}
}

2
intern/cycles/render/scene.h
View 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;