1 /*
2 * Copyright 2011-2013 Blender Foundation
3 *
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
7 *
8 * http://www.apache.org/licenses/LICENSE-2.0
9 *
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.
15 */
16
17 /* Constant Globals */
18
19 #ifndef __KERNEL_GLOBALS_H__
20 #define __KERNEL_GLOBALS_H__
21
22 #include "kernel/kernel_profiling.h"
23
24 #ifdef __KERNEL_CPU__
25 # include "util/util_map.h"
26 # include "util/util_vector.h"
27 #endif
28
29 #ifdef __KERNEL_OPENCL__
30 # include "util/util_atomic.h"
31 #endif
32
33 CCL_NAMESPACE_BEGIN
34
35 /* On the CPU, we pass along the struct KernelGlobals to nearly everywhere in
36 * the kernel, to access constant data. These are all stored as "textures", but
37 * these are really just standard arrays. We can't use actually globals because
38 * multiple renders may be running inside the same process. */
39
40 #ifdef __KERNEL_CPU__
41
42 # ifdef __OSL__
43 struct OSLGlobals;
44 struct OSLThreadData;
45 struct OSLShadingSystem;
46 # endif
47
48 typedef unordered_map<float, float> CoverageMap;
49
50 struct Intersection;
51 struct VolumeStep;
52
53 typedef struct KernelGlobals {
54 # define KERNEL_TEX(type, name) texture<type> name;
55 # include "kernel/kernel_textures.h"
56
57 KernelData __data;
58
59 # ifdef __OSL__
60 /* On the CPU, we also have the OSL globals here. Most data structures are shared
61 * with SVM, the difference is in the shaders and object/mesh attributes. */
62 OSLGlobals *osl;
63 OSLShadingSystem *osl_ss;
64 OSLThreadData *osl_tdata;
65 # endif
66
67 /* **** Run-time data **** */
68
69 /* Heap-allocated storage for transparent shadows intersections. */
70 Intersection *transparent_shadow_intersections;
71
72 /* Storage for decoupled volume steps. */
73 VolumeStep *decoupled_volume_steps[2];
74 int decoupled_volume_steps_index;
75
76 /* A buffer for storing per-pixel coverage for Cryptomatte. */
77 CoverageMap *coverage_object;
78 CoverageMap *coverage_material;
79 CoverageMap *coverage_asset;
80
81 /* split kernel */
82 SplitData split_data;
83 SplitParams split_param_data;
84
85 int2 global_size;
86 int2 global_id;
87
88 ProfilingState profiler;
89 } KernelGlobals;
90
91 #endif /* __KERNEL_CPU__ */
92
93 #ifdef __KERNEL_OPTIX__
94
95 typedef struct ShaderParams {
96 uint4 *input;
97 float4 *output;
98 int type;
99 int filter;
100 int sx;
101 int offset;
102 int sample;
103 } ShaderParams;
104
105 typedef struct KernelParams {
106 WorkTile tile;
107 KernelData data;
108 ShaderParams shader;
109 # define KERNEL_TEX(type, name) const type *name;
110 # include "kernel/kernel_textures.h"
111 } KernelParams;
112
113 typedef struct KernelGlobals {
114 # ifdef __VOLUME__
115 VolumeState volume_state;
116 # endif
117 Intersection hits_stack[64];
118 } KernelGlobals;
119
120 extern "C" __constant__ KernelParams __params;
121
122 #else /* __KERNEL_OPTIX__ */
123
124 /* For CUDA, constant memory textures must be globals, so we can't put them
125 * into a struct. As a result we don't actually use this struct and use actual
126 * globals and simply pass along a NULL pointer everywhere, which we hope gets
127 * optimized out. */
128
129 # ifdef __KERNEL_CUDA__
130
131 __constant__ KernelData __data;
132 typedef struct KernelGlobals {
133 /* NOTE: Keep the size in sync with SHADOW_STACK_MAX_HITS. */
134 Intersection hits_stack[64];
135 } KernelGlobals;
136
137 # define KERNEL_TEX(type, name) const __constant__ __device__ type *name;
138 # include "kernel/kernel_textures.h"
139
140 # endif /* __KERNEL_CUDA__ */
141
142 #endif /* __KERNEL_OPTIX__ */
143
144 /* OpenCL */
145
146 #ifdef __KERNEL_OPENCL__
147
148 # define KERNEL_TEX(type, name) typedef type name##_t;
149 # include "kernel/kernel_textures.h"
150
151 typedef ccl_addr_space struct KernelGlobals {
152 ccl_constant KernelData *data;
153 ccl_global char *buffers[8];
154
155 # define KERNEL_TEX(type, name) TextureInfo name;
156 # include "kernel/kernel_textures.h"
157
158 # ifdef __SPLIT_KERNEL__
159 SplitData split_data;
160 SplitParams split_param_data;
161 # endif
162 } KernelGlobals;
163
164 # define KERNEL_BUFFER_PARAMS \
165 ccl_global char *buffer0, ccl_global char *buffer1, ccl_global char *buffer2, \
166 ccl_global char *buffer3, ccl_global char *buffer4, ccl_global char *buffer5, \
167 ccl_global char *buffer6, ccl_global char *buffer7
168
169 # define KERNEL_BUFFER_ARGS buffer0, buffer1, buffer2, buffer3, buffer4, buffer5, buffer6, buffer7
170
kernel_set_buffer_pointers(KernelGlobals * kg,KERNEL_BUFFER_PARAMS)171 ccl_device_inline void kernel_set_buffer_pointers(KernelGlobals *kg, KERNEL_BUFFER_PARAMS)
172 {
173 # ifdef __SPLIT_KERNEL__
174 if (ccl_local_id(0) + ccl_local_id(1) == 0)
175 # endif
176 {
177 kg->buffers[0] = buffer0;
178 kg->buffers[1] = buffer1;
179 kg->buffers[2] = buffer2;
180 kg->buffers[3] = buffer3;
181 kg->buffers[4] = buffer4;
182 kg->buffers[5] = buffer5;
183 kg->buffers[6] = buffer6;
184 kg->buffers[7] = buffer7;
185 }
186
187 # ifdef __SPLIT_KERNEL__
188 ccl_barrier(CCL_LOCAL_MEM_FENCE);
189 # endif
190 }
191
kernel_set_buffer_info(KernelGlobals * kg)192 ccl_device_inline void kernel_set_buffer_info(KernelGlobals *kg)
193 {
194 # ifdef __SPLIT_KERNEL__
195 if (ccl_local_id(0) + ccl_local_id(1) == 0)
196 # endif
197 {
198 ccl_global TextureInfo *info = (ccl_global TextureInfo *)kg->buffers[0];
199
200 # define KERNEL_TEX(type, name) kg->name = *(info++);
201 # include "kernel/kernel_textures.h"
202 }
203
204 # ifdef __SPLIT_KERNEL__
205 ccl_barrier(CCL_LOCAL_MEM_FENCE);
206 # endif
207 }
208
209 #endif /* __KERNEL_OPENCL__ */
210
211 /* Interpolated lookup table access */
212
lookup_table_read(KernelGlobals * kg,float x,int offset,int size)213 ccl_device float lookup_table_read(KernelGlobals *kg, float x, int offset, int size)
214 {
215 x = saturate(x) * (size - 1);
216
217 int index = min(float_to_int(x), size - 1);
218 int nindex = min(index + 1, size - 1);
219 float t = x - index;
220
221 float data0 = kernel_tex_fetch(__lookup_table, index + offset);
222 if (t == 0.0f)
223 return data0;
224
225 float data1 = kernel_tex_fetch(__lookup_table, nindex + offset);
226 return (1.0f - t) * data0 + t * data1;
227 }
228
lookup_table_read_2D(KernelGlobals * kg,float x,float y,int offset,int xsize,int ysize)229 ccl_device float lookup_table_read_2D(
230 KernelGlobals *kg, float x, float y, int offset, int xsize, int ysize)
231 {
232 y = saturate(y) * (ysize - 1);
233
234 int index = min(float_to_int(y), ysize - 1);
235 int nindex = min(index + 1, ysize - 1);
236 float t = y - index;
237
238 float data0 = lookup_table_read(kg, x, offset + xsize * index, xsize);
239 if (t == 0.0f)
240 return data0;
241
242 float data1 = lookup_table_read(kg, x, offset + xsize * nindex, xsize);
243 return (1.0f - t) * data0 + t * data1;
244 }
245
246 CCL_NAMESPACE_END
247
248 #endif /* __KERNEL_GLOBALS_H__ */
249