2 * Copyright 2011-2013 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 /* OpenCL kernel entry points - unfinished */
19 #include "kernel/kernel_compat_opencl.h"
20 #include "kernel/kernel_math.h"
21 #include "kernel/kernel_types.h"
22 #include "kernel/kernel_globals.h"
23 #include "kernel/kernel_color.h"
24 #include "kernel/kernels/opencl/kernel_opencl_image.h"
26 #include "kernel/kernel_film.h"
28 #if defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__)
29 # include "kernel/kernel_path.h"
30 # include "kernel/kernel_path_branched.h"
31 #else /* __COMPILE_ONLY_MEGAKERNEL__ */
32 /* Include only actually used headers for the case
33 * when path tracing kernels are not needed.
35 # include "kernel/kernel_random.h"
36 # include "kernel/kernel_differential.h"
37 # include "kernel/kernel_montecarlo.h"
38 # include "kernel/kernel_projection.h"
39 # include "kernel/geom/geom.h"
40 # include "kernel/bvh/bvh.h"
42 # include "kernel/kernel_accumulate.h"
43 # include "kernel/kernel_camera.h"
44 # include "kernel/kernel_shader.h"
45 #endif /* defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__) */
47 #include "kernel/kernel_bake.h"
49 #ifdef __COMPILE_ONLY_MEGAKERNEL__
51 __kernel void kernel_ocl_path_trace(
52 ccl_constant KernelData *data,
53 ccl_global float *buffer,
58 int sx, int sy, int sw, int sh, int offset, int stride)
60 KernelGlobals kglobals, *kg = &kglobals;
64 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
65 kernel_set_buffer_info(kg);
67 int x = sx + ccl_global_id(0);
68 int y = sy + ccl_global_id(1);
69 bool thread_is_active = x < sx + sw && y < sy + sh;
70 if(thread_is_active) {
71 kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
73 if(kernel_data.film.cryptomatte_passes) {
74 /* Make sure no thread is writing to the buffers. */
75 ccl_barrier(CCL_LOCAL_MEM_FENCE);
76 if(thread_is_active) {
77 kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride);
82 #else /* __COMPILE_ONLY_MEGAKERNEL__ */
84 __kernel void kernel_ocl_displace(
85 ccl_constant KernelData *data,
86 ccl_global uint4 *input,
87 ccl_global float4 *output,
91 int type, int sx, int sw, int offset, int sample)
93 KernelGlobals kglobals, *kg = &kglobals;
97 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
98 kernel_set_buffer_info(kg);
100 int x = sx + ccl_global_id(0);
103 kernel_displace_evaluate(kg, input, output, x);
106 __kernel void kernel_ocl_background(
107 ccl_constant KernelData *data,
108 ccl_global uint4 *input,
109 ccl_global float4 *output,
111 KERNEL_BUFFER_PARAMS,
113 int type, int sx, int sw, int offset, int sample)
115 KernelGlobals kglobals, *kg = &kglobals;
119 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
120 kernel_set_buffer_info(kg);
122 int x = sx + ccl_global_id(0);
125 kernel_background_evaluate(kg, input, output, x);
129 __kernel void kernel_ocl_bake(
130 ccl_constant KernelData *data,
131 ccl_global uint4 *input,
132 ccl_global float4 *output,
134 KERNEL_BUFFER_PARAMS,
136 int type, int filter, int sx, int sw, int offset, int sample)
138 KernelGlobals kglobals, *kg = &kglobals;
142 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
143 kernel_set_buffer_info(kg);
145 int x = sx + ccl_global_id(0);
149 output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
151 kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, x, offset, sample);
156 __kernel void kernel_ocl_convert_to_byte(
157 ccl_constant KernelData *data,
158 ccl_global uchar4 *rgba,
159 ccl_global float *buffer,
161 KERNEL_BUFFER_PARAMS,
164 int sx, int sy, int sw, int sh, int offset, int stride)
166 KernelGlobals kglobals, *kg = &kglobals;
170 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
171 kernel_set_buffer_info(kg);
173 int x = sx + ccl_global_id(0);
174 int y = sy + ccl_global_id(1);
176 if(x < sx + sw && y < sy + sh)
177 kernel_film_convert_to_byte(kg, rgba, buffer, sample_scale, x, y, offset, stride);
180 __kernel void kernel_ocl_convert_to_half_float(
181 ccl_constant KernelData *data,
182 ccl_global uchar4 *rgba,
183 ccl_global float *buffer,
185 KERNEL_BUFFER_PARAMS,
188 int sx, int sy, int sw, int sh, int offset, int stride)
190 KernelGlobals kglobals, *kg = &kglobals;
194 kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
195 kernel_set_buffer_info(kg);
197 int x = sx + ccl_global_id(0);
198 int y = sy + ccl_global_id(1);
200 if(x < sx + sw && y < sy + sh)
201 kernel_film_convert_to_half_float(kg, rgba, buffer, sample_scale, x, y, offset, stride);
204 __kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, uint64_t offset)
206 size_t i = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0);
208 if(i < size / sizeof(float4)) {
209 buffer[i+offset/sizeof(float4)] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
211 else if(i == size / sizeof(float4)) {
212 ccl_global uchar *b = (ccl_global uchar*)&buffer[i+offset/sizeof(float4)];
214 for(i = 0; i < size % sizeof(float4); i++) {
220 #endif /* __COMPILE_ONLY_MEGAKERNEL__ */