[Bf-blender-cvs] [d133934ea46] master: Cycles: code to optionally zero initialize some structs in the kernel
Brecht Van Lommel
noreply at git.blender.org
Mon Aug 26 16:12:19 CEST 2019
Commit: d133934ea461995fe959f130115652ff803d4269
Author: Brecht Van Lommel
Date: Mon Aug 26 14:41:15 2019 +0200
Branches: master
https://developer.blender.org/rBd133934ea461995fe959f130115652ff803d4269
Cycles: code to optionally zero initialize some structs in the kernel
This will be used by Optix to help the compiler figure out scoping. It is not
used by other devices currently, but worth testing if it helps there too.
Ref D5363
===================================================================
M intern/cycles/kernel/kernel_compat_cuda.h
M intern/cycles/kernel/kernel_compat_opencl.h
M intern/cycles/kernel/kernel_emission.h
M intern/cycles/kernel/kernel_path.h
M intern/cycles/kernel/kernel_path_surface.h
M intern/cycles/kernel/kernel_path_volume.h
M intern/cycles/kernel/kernel_volume.h
M intern/cycles/util/util_defines.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index 5075c434b10..4f508d7cdaa 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -73,6 +73,7 @@ __device__ half __float2half(const float f)
*/
#define ccl_ref
#define ccl_align(n) __align__(n)
+#define ccl_optional_struct_init
#define ATTR_FALLTHROUGH
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index 1fe52c51ab0..552734cc361 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -46,6 +46,7 @@
#define ccl_restrict restrict
#define ccl_ref
#define ccl_align(n) __attribute__((aligned(n)))
+#define ccl_optional_struct_init
#ifdef __SPLIT_KERNEL__
# define ccl_addr_space __global
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index 16d52b0c733..9761dbfcc6d 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -245,7 +245,7 @@ ccl_device_noinline_cpu bool indirect_lamp_emission(KernelGlobals *kg,
*emission = make_float3(0.0f, 0.0f, 0.0f);
for (int lamp = 0; lamp < kernel_data.integrator.num_all_lights; lamp++) {
- LightSample ls;
+ LightSample ls ccl_optional_struct_init;
if (!lamp_light_eval(kg, lamp, ray->P, ray->D, ray->t, &ls))
continue;
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 63be0a7f505..0dc55eba14a 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -92,7 +92,7 @@ ccl_device_forceinline void kernel_path_lamp_emission(KernelGlobals *kg,
#ifdef __LAMP_MIS__
if (kernel_data.integrator.use_lamp_mis && !(state->flag & PATH_RAY_CAMERA)) {
/* ray starting from previous non-transparent bounce */
- Ray light_ray;
+ Ray light_ray ccl_optional_struct_init;
light_ray.P = ray->P - state->ray_t * ray->D;
state->ray_t += isect->t;
diff --git a/intern/cycles/kernel/kernel_path_surface.h b/intern/cycles/kernel/kernel_path_surface.h
index d299106ea96..a32690d51eb 100644
--- a/intern/cycles/kernel/kernel_path_surface.h
+++ b/intern/cycles/kernel/kernel_path_surface.h
@@ -32,7 +32,7 @@ ccl_device_noinline_cpu void kernel_branched_path_surface_connect_light(
{
# ifdef __EMISSION__
/* sample illumination from lights to find path contribution */
- BsdfEval L_light;
+ BsdfEval L_light ccl_optional_struct_init;
int num_lights = 0;
if (kernel_data.integrator.use_direct_light) {
@@ -79,7 +79,7 @@ ccl_device_noinline_cpu void kernel_branched_path_surface_connect_light(
float num_samples_inv = num_samples_adjust / (num_samples * num_all_lights);
for (int j = 0; j < num_samples; j++) {
- Ray light_ray;
+ Ray light_ray ccl_optional_struct_init;
light_ray.t = 0.0f; /* reset ray */
# ifdef __OBJECT_MOTION__
light_ray.time = sd->time;
@@ -98,7 +98,7 @@ ccl_device_noinline_cpu void kernel_branched_path_surface_connect_light(
light_u = 0.5f * light_u;
}
- LightSample ls;
+ LightSample ls ccl_optional_struct_init;
const int lamp = is_lamp ? i : -1;
if (light_sample(kg, lamp, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) {
/* The sampling probability returned by lamp_light_sample assumes that all lights were
@@ -147,9 +147,9 @@ ccl_device bool kernel_branched_path_surface_bounce(KernelGlobals *kg,
{
/* sample BSDF */
float bsdf_pdf;
- BsdfEval bsdf_eval;
- float3 bsdf_omega_in;
- differential3 bsdf_domega_in;
+ BsdfEval bsdf_eval ccl_optional_struct_init;
+ float3 bsdf_omega_in ccl_optional_struct_init;
+ differential3 bsdf_domega_in ccl_optional_struct_init;
float bsdf_u, bsdf_v;
path_branched_rng_2D(
kg, state->rng_hash, state, sample, num_samples, PRNG_BSDF_U, &bsdf_u, &bsdf_v);
@@ -220,8 +220,8 @@ ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg,
kernel_branched_path_surface_connect_light(kg, sd, emission_sd, state, throughput, 1.0f, L, all);
# else
/* sample illumination from lights to find path contribution */
- Ray light_ray;
- BsdfEval L_light;
+ Ray light_ray ccl_optional_struct_init;
+ BsdfEval L_light ccl_optional_struct_init;
bool is_lamp = false;
bool has_emission = false;
@@ -234,7 +234,7 @@ ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg,
float light_u, light_v;
path_state_rng_2D(kg, state, PRNG_LIGHT_U, &light_u, &light_v);
- LightSample ls;
+ LightSample ls ccl_optional_struct_init;
if (light_sample(kg, -1, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) {
float terminate = path_state_rng_light_termination(kg, state);
has_emission = direct_emission(
@@ -274,9 +274,9 @@ ccl_device bool kernel_path_surface_bounce(KernelGlobals *kg,
if (sd->flag & SD_BSDF) {
/* sample BSDF */
float bsdf_pdf;
- BsdfEval bsdf_eval;
- float3 bsdf_omega_in;
- differential3 bsdf_domega_in;
+ BsdfEval bsdf_eval ccl_optional_struct_init;
+ float3 bsdf_omega_in ccl_optional_struct_init;
+ differential3 bsdf_domega_in ccl_optional_struct_init;
float bsdf_u, bsdf_v;
path_state_rng_2D(kg, state, PRNG_BSDF_U, &bsdf_u, &bsdf_v);
int label;
diff --git a/intern/cycles/kernel/kernel_path_volume.h b/intern/cycles/kernel/kernel_path_volume.h
index 6b62005d19a..db10629ee9f 100644
--- a/intern/cycles/kernel/kernel_path_volume.h
+++ b/intern/cycles/kernel/kernel_path_volume.h
@@ -27,8 +27,8 @@ ccl_device_inline void kernel_path_volume_connect_light(KernelGlobals *kg,
{
# ifdef __EMISSION__
/* sample illumination from lights to find path contribution */
- Ray light_ray;
- BsdfEval L_light;
+ Ray light_ray ccl_optional_struct_init;
+ BsdfEval L_light ccl_optional_struct_init;
bool is_lamp = false;
bool has_emission = false;
@@ -42,7 +42,7 @@ ccl_device_inline void kernel_path_volume_connect_light(KernelGlobals *kg,
float light_u, light_v;
path_state_rng_2D(kg, state, PRNG_LIGHT_U, &light_u, &light_v);
- LightSample ls;
+ LightSample ls ccl_optional_struct_init;
if (light_sample(kg, -1, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) {
float terminate = path_state_rng_light_termination(kg, state);
has_emission = direct_emission(
@@ -71,9 +71,9 @@ ccl_device_noinline_cpu bool kernel_path_volume_bounce(KernelGlobals *kg,
{
/* sample phase function */
float phase_pdf;
- BsdfEval phase_eval;
- float3 phase_omega_in;
- differential3 phase_domega_in;
+ BsdfEval phase_eval ccl_optional_struct_init;
+ float3 phase_omega_in ccl_optional_struct_init;
+ differential3 phase_domega_in ccl_optional_struct_init;
float phase_u, phase_v;
path_state_rng_2D(kg, state, PRNG_BSDF_U, &phase_u, &phase_v);
int label;
@@ -139,7 +139,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg,
const VolumeSegment *segment)
{
# ifdef __EMISSION__
- BsdfEval L_light;
+ BsdfEval L_light ccl_optional_struct_init;
int num_lights = 1;
if (sample_all_lights) {
@@ -181,7 +181,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg,
float num_samples_inv = 1.0f / (num_samples * num_all_lights);
for (int j = 0; j < num_samples; j++) {
- Ray light_ray;
+ Ray light_ray ccl_optional_struct_init;
light_ray.t = 0.0f; /* reset ray */
# ifdef __OBJECT_MOTION__
light_ray.time = sd->time;
@@ -201,7 +201,7 @@ ccl_device void kernel_branched_path_volume_connect_light(KernelGlobals *kg,
light_u = 0.5f * light_u;
}
- LightSample ls;
+ LightSample ls ccl_optional_struct_init;
const int lamp = is_lamp ? i : -1;
light_sample(kg, lamp, light_u, light_v, sd->time, ray->P, state->bounce, &ls);
diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h
index 2705526abe4..4ddcbec0fef 100644
--- a/intern/cycles/kernel/kernel_volume.h
+++ b/intern/cycles/kernel/kernel_volume.h
@@ -428,7 +428,7 @@ kernel_volume_integrate_homogeneous(KernelGlobals *kg,
ccl_addr_space float3 *throughput,
bool probalistic_scatter)
{
- VolumeShaderCoefficients coeff;
+ VolumeShaderCoefficients coeff ccl_optional_struct_init;
if (!volume_shader_sample(kg, sd, state, ray->P, &coeff))
return VOLUME_PATH_MISSED;
@@ -565,7 +565,7 @@ kernel_volume_integrate_heterogeneous_distance(KernelGlobals *kg,
}
float3 new_P = ray->P + ray->D * (t + step_offset);
- VolumeShaderCoefficients coeff;
+ VolumeShaderCoefficients coeff ccl_optional_struct_init;
/* compute segment */
if (volume_shader_sample(kg, sd, state, new_P, &coeff)) {
@@ -801,7 +801,7 @@ ccl_device void kernel_volume_decoupled_record(KernelGlobals *kg,
}
float3 new_P = ray->P + ray->D * (t + step_offset);
- VolumeShaderCoefficients coeff;
+ VolumeShaderCoefficients coeff ccl_optional_struct_init;
/* compute segment */
if (volume_shader_sample(kg, sd, state, new_P, &coeff)) {
diff --git a/intern/cycles/util/util_defines.h b/intern/cycles/util/util_defines.h
index 760985447a8..2778cffba3a 100644
--- a/intern/cycles/util/util_defines.h
+++ b/intern/cycles/util/util_defines.h
@@ -39,6 +39,7 @@
# define ccl_private
# define ccl_restrict __restrict
# define ccl_ref &
+# define ccl_optional_struct_init
# define __KERNEL_WITH_SSE_ALIGN__
# if defined(_WIN32) && !defined(FREE_WINDOWS)
More information about the Bf-blender-cvs
mailing list