diff options
-rw-r--r-- | intern/cycles/kernel/kernel_compat_cuda.h | 1 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_compat_opencl.h | 1 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_emission.h | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_path.h | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_path_surface.h | 24 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_path_volume.h | 18 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_volume.h | 6 | ||||
-rw-r--r-- | intern/cycles/util/util_defines.h | 1 |
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) |