bartvdbraak/blender
1
0
Fork 4
Code Issues 4 Pull Requests 1 Actions 1 Packages Projects Releases Wiki Activity

Cycles: code to optionally zero initialize some structs in the kernel

Browse Source 
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
This commit is contained in:
Brecht Van Lommel 2019-08-26 14:41:15 +02:00
parent 39439a3afe
commit d133934ea4
8 changed files with 29 additions and 26 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
intern/cycles/kernel/kernel_compat_cuda.h
View File

@ -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

1
intern/cycles/kernel/kernel_compat_opencl.h
View File

@ -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

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

@ -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;

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

@ -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;

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

@ -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;

18
intern/cycles/kernel/kernel_path_volume.h
View File

@ -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);

6
intern/cycles/kernel/kernel_volume.h
View File

@ -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)) {

1
intern/cycles/util/util_defines.h
View File

@ -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)