diff options
Diffstat (limited to 'intern/cycles/kernel/split/kernel_split_data_types.h')
| -rw-r--r-- | intern/cycles/kernel/split/kernel_split_data_types.h | 14 |
1 files changed, 11 insertions, 3 deletions
diff --git a/intern/cycles/kernel/split/kernel_split_data_types.h b/intern/cycles/kernel/split/kernel_split_data_types.h index 5c2aadcf4ec..9ac3f904819 100644 --- a/intern/cycles/kernel/split/kernel_split_data_types.h +++ b/intern/cycles/kernel/split/kernel_split_data_types.h @@ -51,7 +51,6 @@ typedef ccl_global struct SplitBranchedState { float3 throughput; Ray ray; - struct ShaderData sd; Intersection isect; char ray_state; @@ -77,6 +76,9 @@ typedef ccl_global struct SplitBranchedState { int shared_sample_count; /* number of branched samples shared with other threads */ int original_ray; /* index of original ray when sharing branched samples */ bool waiting_on_shared_samples; + + /* Must be last in to allow for dynamic size of closures */ + struct ShaderData sd; } SplitBranchedState; #define SPLIT_DATA_BRANCHED_ENTRIES \ @@ -110,11 +112,11 @@ typedef ccl_global struct SplitBranchedState { SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \ SPLIT_DATA_ENTRY(ccl_global uint, buffer_offset, 1) \ - SPLIT_DATA_ENTRY(ShaderData, sd, 1) \ SPLIT_DATA_ENTRY(ShaderDataTinyStorage, sd_DL_shadow, 1) \ SPLIT_DATA_SUBSURFACE_ENTRIES \ SPLIT_DATA_VOLUME_ENTRIES \ SPLIT_DATA_BRANCHED_ENTRIES \ + SPLIT_DATA_ENTRY(ShaderData, _sd, 0) /* entries to be copied to inactive rays when sharing branched samples (TODO: which are actually needed?) */ #define SPLIT_DATA_ENTRIES_BRANCHED_SHARED \ @@ -126,11 +128,11 @@ typedef ccl_global struct SplitBranchedState { SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ - SPLIT_DATA_ENTRY(ShaderData, sd, 1) \ SPLIT_DATA_ENTRY(ShaderDataTinyStorage, sd_DL_shadow, 1) \ SPLIT_DATA_SUBSURFACE_ENTRIES \ SPLIT_DATA_VOLUME_ENTRIES \ SPLIT_DATA_BRANCHED_ENTRIES \ + SPLIT_DATA_ENTRY(ShaderData, _sd, 0) /* struct that holds pointers to data in the shared state buffer */ typedef struct SplitData { @@ -154,6 +156,12 @@ __device__ SplitParams __split_param_data; # define kernel_split_params (__split_param_data) #endif /* __KERNEL_CUDA__ */ +#define kernel_split_sd(sd, ray_index) ((ShaderData*) \ + ( \ + ((ccl_global char*)kernel_split_state._##sd) + \ + (sizeof(ShaderData) + sizeof(ShaderClosure)*(kernel_data.integrator.max_closures-1)) * (ray_index) \ + )) + /* Local storage for queue_enqueue kernel. */ typedef struct QueueEnqueueLocals { uint queue_atomics[2]; |