Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--intern/cycles/kernel/kernel_compat_cuda.h1
-rw-r--r--intern/cycles/kernel/kernel_compat_opencl.h1
-rw-r--r--intern/cycles/kernel/kernel_emission.h2
-rw-r--r--intern/cycles/kernel/kernel_path.h2
-rw-r--r--intern/cycles/kernel/kernel_path_surface.h24
-rw-r--r--intern/cycles/kernel/kernel_path_volume.h18
-rw-r--r--intern/cycles/kernel/kernel_volume.h6
-rw-r--r--intern/cycles/util/util_defines.h1
8 files changed, 29 insertions, 26 deletions
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)
generated by cgit v1.2.3 (git 2.25.1) at 2024-08-08 21:59:45 +0300