2 * Copyright 2011-2015 Blender Foundation
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
8 * http://www.apache.org/licenses/LICENSE-2.0
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
17 #include "split/kernel_data_init.h"
19 __kernel void kernel_ocl_path_trace_data_init(
22 ccl_global char *sd_DL_shadow,
24 ccl_global float3 *P_sd,
25 ccl_global float3 *P_sd_DL_shadow,
27 ccl_global float3 *N_sd,
28 ccl_global float3 *N_sd_DL_shadow,
30 ccl_global float3 *Ng_sd,
31 ccl_global float3 *Ng_sd_DL_shadow,
33 ccl_global float3 *I_sd,
34 ccl_global float3 *I_sd_DL_shadow,
36 ccl_global int *shader_sd,
37 ccl_global int *shader_sd_DL_shadow,
39 ccl_global int *flag_sd,
40 ccl_global int *flag_sd_DL_shadow,
42 ccl_global int *prim_sd,
43 ccl_global int *prim_sd_DL_shadow,
45 ccl_global int *type_sd,
46 ccl_global int *type_sd_DL_shadow,
48 ccl_global float *u_sd,
49 ccl_global float *u_sd_DL_shadow,
51 ccl_global float *v_sd,
52 ccl_global float *v_sd_DL_shadow,
54 ccl_global int *object_sd,
55 ccl_global int *object_sd_DL_shadow,
57 ccl_global float *time_sd,
58 ccl_global float *time_sd_DL_shadow,
60 ccl_global float *ray_length_sd,
61 ccl_global float *ray_length_sd_DL_shadow,
63 /* Ray differentials. */
64 ccl_global differential3 *dP_sd,
65 ccl_global differential3 *dP_sd_DL_shadow,
67 ccl_global differential3 *dI_sd,
68 ccl_global differential3 *dI_sd_DL_shadow,
70 ccl_global differential *du_sd,
71 ccl_global differential *du_sd_DL_shadow,
73 ccl_global differential *dv_sd,
74 ccl_global differential *dv_sd_DL_shadow,
77 ccl_global float3 *dPdu_sd,
78 ccl_global float3 *dPdu_sd_DL_shadow,
80 ccl_global float3 *dPdv_sd,
81 ccl_global float3 *dPdv_sd_DL_shadow,
84 ccl_global Transform *ob_tfm_sd,
85 ccl_global Transform *ob_tfm_sd_DL_shadow,
87 ccl_global Transform *ob_itfm_sd,
88 ccl_global Transform *ob_itfm_sd_DL_shadow,
90 ShaderClosure *closure_sd,
91 ShaderClosure *closure_sd_DL_shadow,
93 ccl_global int *num_closure_sd,
94 ccl_global int *num_closure_sd_DL_shadow,
96 ccl_global float *randb_closure_sd,
97 ccl_global float *randb_closure_sd_DL_shadow,
99 ccl_global float3 *ray_P_sd,
100 ccl_global float3 *ray_P_sd_DL_shadow,
102 ccl_global differential3 *ray_dP_sd,
103 ccl_global differential3 *ray_dP_sd_DL_shadow,
105 ccl_constant KernelData *data,
106 ccl_global float *per_sample_output_buffers,
107 ccl_global uint *rng_state,
108 ccl_global uint *rng_coop, /* rng array to store rng values for all rays */
109 ccl_global float3 *throughput_coop, /* throughput array to store throughput values for all rays */
110 ccl_global float *L_transparent_coop, /* L_transparent array to store L_transparent values for all rays */
111 PathRadiance *PathRadiance_coop, /* PathRadiance array to store PathRadiance values for all rays */
112 ccl_global Ray *Ray_coop, /* Ray array to store Ray information for all rays */
113 ccl_global PathState *PathState_coop, /* PathState array to store PathState information for all rays */
114 ccl_global char *ray_state, /* Stores information on current state of a ray */
116 #define KERNEL_TEX(type, ttype, name) \
117 ccl_global type *name,
118 #include "../../kernel_textures.h"
120 int start_sample, int sx, int sy, int sw, int sh, int offset, int stride,
121 int rng_state_offset_x,
122 int rng_state_offset_y,
123 int rng_state_stride,
124 ccl_global int *Queue_data, /* Memory for queues */
125 ccl_global int *Queue_index, /* Tracks the number of elements in queues */
126 int queuesize, /* size (capacity) of the queue */
127 ccl_global char *use_queues_flag, /* flag to decide if scene-intersect kernel should use queues to fetch ray index */
128 ccl_global unsigned int *work_array, /* work array to store which work each ray belongs to */
129 #ifdef __WORK_STEALING__
130 ccl_global unsigned int *work_pool_wgs, /* Work pool for each work group */
131 unsigned int num_samples, /* Total number of samples per pixel */
133 #ifdef __KERNEL_DEBUG__
134 DebugData *debugdata_coop,
136 int parallel_samples) /* Number of samples to be processed in parallel */
138 kernel_data_init((KernelGlobals *)kg,
140 (ShaderData *)sd_DL_shadow,
166 ray_length_sd_DL_shadow,
168 /* Ray differentials. */
188 ob_itfm_sd_DL_shadow,
191 closure_sd_DL_shadow,
193 num_closure_sd_DL_shadow,
195 randb_closure_sd_DL_shadow,
201 per_sample_output_buffers,
211 #define KERNEL_TEX(type, ttype, name) name,
212 #include "../../kernel_textures.h"
214 start_sample, sx, sy, sw, sh, offset, stride,
223 #ifdef __WORK_STEALING__
227 #ifdef __KERNEL_DEBUG__