[Bf-blender-cvs] [11b8d04] cycles_kernel_split: Cycles kernel split: Move address space specification to PathRadiance typedef
Sergey Sharybin
noreply at git.blender.org
Thu May 7 11:57:47 CEST 2015
Commit: 11b8d0495b3aabf2b744e7bc0656fd6a8d4f2c24
Author: Sergey Sharybin
Date: Thu May 7 14:26:36 2015 +0500
Branches: cycles_kernel_split
https://developer.blender.org/rB11b8d0495b3aabf2b744e7bc0656fd6a8d4f2c24
Cycles kernel split: Move address space specification to PathRadiance typedef
===================================================================
M intern/cycles/kernel/kernel_Background_BufferUpdate.cl
M intern/cycles/kernel/kernel_DataInit.cl
M intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
M intern/cycles/kernel/kernel_LampEmission.cl
M intern/cycles/kernel/kernel_NextIterationSetUp.cl
M intern/cycles/kernel/kernel_accumulate.h
M intern/cycles/kernel/kernel_passes.h
M intern/cycles/kernel/kernel_path_surface.h
M intern/cycles/kernel/kernel_types.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
index ce9e1a8..9634e60 100644
--- a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
+++ b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
@@ -103,7 +103,7 @@ __kernel void kernel_ocl_path_trace_Background_BufferUpdate_SPLIT_KERNEL(
ccl_global uint *rng_state,
ccl_global uint *rng_coop, /* Required for buffer Update */
ccl_global float3 *throughput_coop, /* Required for background hit processing */
- ccl_global PathRadiance *PathRadiance_coop, /* Required for background hit processing and buffer Update */
+ PathRadiance *PathRadiance_coop, /* Required for background hit processing and buffer Update */
ccl_global Ray *Ray_coop, /* Required for background hit processing */
ccl_global PathState *PathState_coop, /* Required for background hit processing */
ccl_global float *L_transparent_coop, /* Required for background hit processing and buffer Update */
@@ -163,7 +163,7 @@ __kernel void kernel_ocl_path_trace_Background_BufferUpdate_SPLIT_KERNEL(
ccl_global DebugData *debug_data = &debugdata_coop[ray_index];
#endif
ccl_global PathState *state = &PathState_coop[ray_index];
- ccl_global PathRadiance *L = L = &PathRadiance_coop[ray_index];
+ PathRadiance *L = L = &PathRadiance_coop[ray_index];
ccl_global Ray *ray = &Ray_coop[ray_index];
ccl_global float3 *throughput = &throughput_coop[ray_index];
ccl_global float *L_transparent = &L_transparent_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_DataInit.cl b/intern/cycles/kernel/kernel_DataInit.cl
index e4e39de..8019f3d 100644
--- a/intern/cycles/kernel/kernel_DataInit.cl
+++ b/intern/cycles/kernel/kernel_DataInit.cl
@@ -167,7 +167,7 @@ __kernel void kernel_ocl_path_trace_data_initialization_SPLIT_KERNEL(
ccl_global uint *rng_coop, /* rng array to store rng values for all rays */
ccl_global float3 *throughput_coop, /* throughput array to store throughput values for all rays */
ccl_global float *L_transparent_coop, /* L_transparent array to store L_transparent values for all rays */
- ccl_global PathRadiance *PathRadiance_coop, /* PathRadiance array to store PathRadiance values for all rays */
+ PathRadiance *PathRadiance_coop, /* PathRadiance array to store PathRadiance values for all rays */
ccl_global Ray *Ray_coop, /* Ray array to store Ray information for all rays */
ccl_global PathState *PathState_coop, /* PathState array to store PathState information for all rays */
ccl_global char *ray_state, /* Stores information on current state of a ray */
diff --git a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
index 693211d..b419f6e 100644
--- a/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
+++ b/intern/cycles/kernel/kernel_Holdout_Emission_Blurring_Pathtermination_AO.cl
@@ -80,7 +80,7 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_AO
ccl_global uint *rng_coop, /* Required for "kernel_write_data_passes" and AO */
ccl_global float3 *throughput_coop, /* Required for handling holdout material and AO */
ccl_global float *L_transparent_coop, /* Required for handling holdout material */
- ccl_global PathRadiance *PathRadiance_coop, /* Required for "kernel_write_data_passes" and indirect primitive emission */
+ PathRadiance *PathRadiance_coop, /* Required for "kernel_write_data_passes" and indirect primitive emission */
ccl_global PathState *PathState_coop, /* Required throughout the kernel and AO */
Intersection *Intersection_coop, /* Required for indirect primitive emission */
ccl_global float3 *AOAlpha_coop, /* Required for AO */
@@ -187,7 +187,7 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_AO
if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) {
- ccl_global PathRadiance *L = &PathRadiance_coop[ray_index];
+ PathRadiance *L = &PathRadiance_coop[ray_index];
/* holdout mask objects do not write data passes */
kernel_write_data_passes(kg, per_sample_output_buffers, L, sd, sample, state, throughput);
diff --git a/intern/cycles/kernel/kernel_LampEmission.cl b/intern/cycles/kernel/kernel_LampEmission.cl
index 26e2ade..2098a16 100644
--- a/intern/cycles/kernel/kernel_LampEmission.cl
+++ b/intern/cycles/kernel/kernel_LampEmission.cl
@@ -45,7 +45,7 @@ __kernel void kernel_ocl_path_trace_LampEmission_SPLIT_KERNEL(
ccl_constant KernelData *data,
ccl_global char *shader_data, /* Required for lamp emission */
ccl_global float3 *throughput_coop, /* Required for lamp emission */
- ccl_global PathRadiance *PathRadiance_coop, /* Required for lamp emission */
+ PathRadiance *PathRadiance_coop, /* Required for lamp emission */
ccl_global Ray *Ray_coop, /* Required for lamp emission */
ccl_global PathState *PathState_coop, /* Required for lamp emission */
Intersection *Intersection_coop, /* Required for lamp emission */
@@ -92,7 +92,7 @@ __kernel void kernel_ocl_path_trace_LampEmission_SPLIT_KERNEL(
if(IS_STATE(ray_state, ray_index, RAY_ACTIVE) || IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) {
KernelGlobals *kg = (KernelGlobals *)globals;
ShaderData *sd = (ShaderData *)shader_data;
- ccl_global PathRadiance *L = &PathRadiance_coop[ray_index];
+ PathRadiance *L = &PathRadiance_coop[ray_index];
float3 throughput = throughput_coop[ray_index];
Ray ray = Ray_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_NextIterationSetUp.cl b/intern/cycles/kernel/kernel_NextIterationSetUp.cl
index 88e3a6a..693bce6 100644
--- a/intern/cycles/kernel/kernel_NextIterationSetUp.cl
+++ b/intern/cycles/kernel/kernel_NextIterationSetUp.cl
@@ -67,7 +67,7 @@ __kernel void kernel_ocl_path_trace_SetupNextIteration_SPLIT_KERNEL(
ccl_global char *shader_data, /* Required for setting up ray for next iteration */
ccl_global uint *rng_coop, /* Required for setting up ray for next iteration */
ccl_global float3 *throughput_coop, /* Required for setting up ray for next iteration */
- ccl_global PathRadiance *PathRadiance_coop, /* Required for setting up ray for next iteration */
+ PathRadiance *PathRadiance_coop, /* Required for setting up ray for next iteration */
ccl_global Ray *Ray_coop, /* Required for setting up ray for next iteration */
ccl_global PathState *PathState_coop, /* Required for setting up ray for next iteration */
ccl_global Ray *LightRay_dl_coop, /* Required for radiance update - direct lighting */
@@ -124,7 +124,7 @@ __kernel void kernel_ocl_path_trace_SetupNextIteration_SPLIT_KERNEL(
/* Load kernel globals structure and ShaderData structure */
KernelGlobals *kg = (KernelGlobals *)globals;
ShaderData *sd = (ShaderData *)shader_data;
- ccl_global PathRadiance *L = 0x0;
+ PathRadiance *L = 0x0;
ccl_global PathState *state = 0x0;
/* Path radiance update for AO/Direct_lighting's shadow blocked */
diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h
index 0530729..850ceac 100644
--- a/intern/cycles/kernel/kernel_accumulate.h
+++ b/intern/cycles/kernel/kernel_accumulate.h
@@ -128,7 +128,7 @@ ccl_device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value)
* visible as the first non-transparent hit, while indirectly visible are the
* bounces after that. */
-ccl_device_inline void path_radiance_init(ccl_addr_space PathRadiance *L, int use_light_pass)
+ccl_device_inline void path_radiance_init(PathRadiance *L, int use_light_pass)
{
/* clear all */
#ifdef __PASSES__
@@ -177,7 +177,7 @@ ccl_device_inline void path_radiance_init(ccl_addr_space PathRadiance *L, int us
}
-ccl_device_inline void path_radiance_bsdf_bounce(ccl_addr_space PathRadiance *L, ccl_addr_space float3 *throughput,
+ccl_device_inline void path_radiance_bsdf_bounce(PathRadiance *L, ccl_addr_space float3 *throughput,
BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label)
{
float inverse_pdf = 1.0f/bsdf_pdf;
@@ -213,7 +213,7 @@ ccl_device_inline void path_radiance_bsdf_bounce(ccl_addr_space PathRadiance *L,
}
-ccl_device_inline void path_radiance_accum_emission(ccl_addr_space PathRadiance *L, float3 throughput, float3 value, int bounce)
+ccl_device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce)
{
#ifdef __PASSES__
if(L->use_light_pass) {
@@ -232,7 +232,7 @@ ccl_device_inline void path_radiance_accum_emission(ccl_addr_space PathRadiance
}
-ccl_device_inline void path_radiance_accum_ao(ccl_addr_space PathRadiance *L, float3 throughput, float3 alpha, float3 bsdf, float3 ao, int bounce)
+ccl_device_inline void path_radiance_accum_ao(PathRadiance *L, float3 throughput, float3 alpha, float3 bsdf, float3 ao, int bounce)
{
#ifdef __PASSES__
if(L->use_light_pass) {
@@ -254,7 +254,7 @@ ccl_device_inline void path_radiance_accum_ao(ccl_addr_space PathRadiance *L, fl
}
-ccl_device_inline void path_radiance_accum_light(ccl_addr_space PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, float shadow_fac, int bounce, bool is_lamp)
+ccl_device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, float shadow_fac, int bounce, bool is_lamp)
{
#ifdef __PASSES__
if(L->use_light_pass) {
@@ -286,7 +286,7 @@ ccl_device_inline void path_radiance_accum_light(ccl_addr_space PathRadiance *L,
}
-ccl_device_inline void path_radiance_accum_background(ccl_addr_space PathRadiance *L, float3 throughput, float3 value, int bounce)
+ccl_device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce)
{
#ifdef __PASSES__
if(L->use_light_pass) {
@@ -305,7 +305,7 @@ ccl_device_inline void path_radiance_accum_background(ccl_addr_space PathRadianc
}
-ccl_
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list