archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity

Cycles: more code refactoring to rename things internally as well. Also change

Browse Source 
property name back so we keep compatibility.
This commit is contained in:
Brecht Van Lommel
2013-08-23 14:34:34 +00:00
parent 22f4874877
commit 01e22d1b9f
17 changed files with 71 additions and 62 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
intern/cycles/kernel/kernel.cpp
View File

@@ -88,12 +88,12 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
#ifdef __NON_PROGRESSIVE__
if(!kernel_data.integrator.progressive)
kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
#ifdef __BRANCHED_PATH__
if(kernel_data.integrator.branched)
kernel_branched_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
else
#endif
kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */

10
intern/cycles/kernel/kernel.cu
View File

@@ -24,23 +24,23 @@
#include "kernel_path.h"
#include "kernel_displace.h"
extern "C" __global__ void kernel_cuda_path_trace_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
extern "C" __global__ void kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
if(x < sx + sw && y < sy + sh)
kernel_path_trace_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
}
#ifdef __NON_PROGRESSIVE__
extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
#ifdef __BRANCHED_PATH__
extern "C" __global__ void kernel_cuda_branched_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
if(x < sx + sw && y < sy + sh)
kernel_path_trace_non_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride);
kernel_branched_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
}
#endif

2
intern/cycles/kernel/kernel_emission.h
View File

@@ -76,7 +76,7 @@ __device_noinline bool direct_emission(KernelGlobals *kg, ShaderData *sd, int li
{
LightSample ls;
#ifdef __NON_PROGRESSIVE__
#ifdef __BRANCHED_PATH__
if(lindex != -1) {
/* sample position on a specified light */
light_select(kg, lindex, randu, randv, sd->P, &ls);

24
intern/cycles/kernel/kernel_path.h
View File

@@ -231,7 +231,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra
return result;
}
__device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
__device float4 kernel_path_integrator(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
{
/* initialize */
PathRadiance L;
@@ -543,7 +543,7 @@ __device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample,
return make_float4(L_sum.x, L_sum.y, L_sum.z, 1.0f - L_transparent);
}
#ifdef __NON_PROGRESSIVE__
#ifdef __BRANCHED_PATH__
__device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer,
float3 throughput, int num_samples, int num_total_samples,
@@ -778,7 +778,7 @@ __device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray
}
}
__device_noinline void kernel_path_non_progressive_lighting(KernelGlobals *kg, RNG *rng, int sample,
__device_noinline void kernel_branched_path_integrate_lighting(KernelGlobals *kg, RNG *rng, int sample,
ShaderData *sd, float3 throughput, float num_samples_adjust,
float min_ray_pdf, float ray_pdf, PathState state,
int rng_offset, PathRadiance *L, __global float *buffer)
@@ -971,7 +971,7 @@ __device_noinline void kernel_path_non_progressive_lighting(KernelGlobals *kg, R
}
}
__device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
__device float4 kernel_branched_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
{
/* initialize */
PathRadiance L;
@@ -1114,7 +1114,7 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
old_subsurface_scatter_step(kg, &bssrdf_sd, state.flag, sc, &lcg_state, true);
/* compute lighting with the BSDF closure */
kernel_path_non_progressive_lighting(kg, rng, sample*num_samples + j,
kernel_branched_path_integrate_lighting(kg, rng, sample*num_samples + j,
&bssrdf_sd, throughput, num_samples_inv,
ray_pdf, ray_pdf, state, rng_offset, &L, buffer);
}
@@ -1126,7 +1126,7 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
/* compute lighting with the BSDF closure */
for(int hit = 0; hit < num_hits; hit++)
kernel_path_non_progressive_lighting(kg, rng, sample*num_samples + j,
kernel_branched_path_integrate_lighting(kg, rng, sample*num_samples + j,
&bssrdf_sd[hit], throughput, num_samples_inv,
ray_pdf, ray_pdf, state, rng_offset, &L, buffer);
}
@@ -1136,7 +1136,7 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
#endif
/* lighting */
kernel_path_non_progressive_lighting(kg, rng, sample, &sd, throughput,
kernel_branched_path_integrate_lighting(kg, rng, sample, &sd, throughput,
1.0f, ray_pdf, ray_pdf, state, rng_offset, &L, buffer);
/* continue in case of transparency */
@@ -1192,7 +1192,7 @@ __device_inline void kernel_path_trace_setup(KernelGlobals *kg, __global uint *r
camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, time, ray);
}
__device void kernel_path_trace_progressive(KernelGlobals *kg,
__device void kernel_path_trace(KernelGlobals *kg,
__global float *buffer, __global uint *rng_state,
int sample, int x, int y, int offset, int stride)
{
@@ -1213,7 +1213,7 @@ __device void kernel_path_trace_progressive(KernelGlobals *kg,
float4 L;
if (ray.t != 0.0f)
L = kernel_path_progressive(kg, &rng, sample, ray, buffer);
L = kernel_path_integrator(kg, &rng, sample, ray, buffer);
else
L = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
@@ -1223,8 +1223,8 @@ __device void kernel_path_trace_progressive(KernelGlobals *kg,
path_rng_end(kg, rng_state, rng);
}
#ifdef __NON_PROGRESSIVE__
__device void kernel_path_trace_non_progressive(KernelGlobals *kg,
#ifdef __BRANCHED_PATH__
__device void kernel_branched_path_trace(KernelGlobals *kg,
__global float *buffer, __global uint *rng_state,
int sample, int x, int y, int offset, int stride)
{
@@ -1245,7 +1245,7 @@ __device void kernel_path_trace_non_progressive(KernelGlobals *kg,
float4 L;
if (ray.t != 0.0f)
L = kernel_path_non_progressive(kg, &rng, sample, ray, buffer);
L = kernel_branched_path_integrate(kg, &rng, sample, ray, buffer);
else
L = make_float4(0.0f, 0.0f, 0.0f, 0.0f);

2
intern/cycles/kernel/kernel_shader.h
View File

@@ -959,7 +959,7 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect)
/* Merging */
#ifdef __NON_PROGRESSIVE__
#ifdef __BRANCHED_PATH__
__device void shader_merge_closures(KernelGlobals *kg, ShaderData *sd)
{
/* merge identical closures, better when we sample a single closure at a time */

8
intern/cycles/kernel/kernel_sse2.cpp
View File

@@ -37,12 +37,12 @@ CCL_NAMESPACE_BEGIN
void kernel_cpu_sse2_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
#ifdef __NON_PROGRESSIVE__
if(!kernel_data.integrator.progressive)
kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
#ifdef __BRANCHED_PATH__
if(kernel_data.integrator.branched)
kernel_branched_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
else
#endif
kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */

8
intern/cycles/kernel/kernel_sse3.cpp
View File

@@ -39,12 +39,12 @@ CCL_NAMESPACE_BEGIN
void kernel_cpu_sse3_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
#ifdef __NON_PROGRESSIVE__
if(!kernel_data.integrator.progressive)
kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
#ifdef __BRANCHED_PATH__
if(kernel_data.integrator.branched)
kernel_branched_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
else
#endif
kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */

8
intern/cycles/kernel/kernel_types.h
View File

@@ -55,7 +55,7 @@ CCL_NAMESPACE_BEGIN
#ifdef __KERNEL_CPU__
#define __KERNEL_SHADING__
#define __KERNEL_ADV_SHADING__
#define __NON_PROGRESSIVE__
#define __BRANCHED_PATH__
#ifdef WITH_OSL
#define __OSL__
#endif
@@ -67,7 +67,7 @@ CCL_NAMESPACE_BEGIN
#define __KERNEL_SHADING__
#if __CUDA_ARCH__ >= 200
#define __KERNEL_ADV_SHADING__
#define __NON_PROGRESSIVE__
#define __BRANCHED_PATH__
#endif
#endif
@@ -763,8 +763,8 @@ typedef struct KernelIntegrator {
/* clamp */
float sample_clamp;
/* non-progressive */
int progressive;
/* branched path */
int branched;
int aa_samples;
int diffuse_samples;
int glossy_samples;