diff options
Diffstat (limited to 'intern/cycles/kernel/kernels/opencl')
32 files changed, 0 insertions, 1598 deletions
diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl deleted file mode 100644 index 996bc27f71b..00000000000 --- a/intern/cycles/kernel/kernels/opencl/filter.cl +++ /dev/null @@ -1,321 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -/* OpenCL kernel entry points */ - -#include "kernel/kernel_compat_opencl.h" - -#include "kernel/filter/filter_kernel.h" - -/* kernels */ - -__kernel void kernel_ocl_filter_divide_shadow(int sample, - CCL_FILTER_TILE_INFO, - ccl_global float *unfilteredA, - ccl_global float *unfilteredB, - ccl_global float *sampleVariance, - ccl_global float *sampleVarianceV, - ccl_global float *bufferVariance, - int4 prefilter_rect, - int buffer_pass_stride, - int buffer_denoising_offset) -{ - int x = prefilter_rect.x + get_global_id(0); - int y = prefilter_rect.y + get_global_id(1); - if(x < prefilter_rect.z && y < prefilter_rect.w) { - kernel_filter_divide_shadow(sample, - CCL_FILTER_TILE_INFO_ARG, - x, y, - unfilteredA, - unfilteredB, - sampleVariance, - sampleVarianceV, - bufferVariance, - prefilter_rect, - buffer_pass_stride, - buffer_denoising_offset); - } -} - -__kernel void kernel_ocl_filter_get_feature(int sample, - CCL_FILTER_TILE_INFO, - int m_offset, - int v_offset, - ccl_global float *mean, - ccl_global float *variance, - float scale, - int4 prefilter_rect, - int buffer_pass_stride, - int buffer_denoising_offset) -{ - int x = prefilter_rect.x + get_global_id(0); - int y = prefilter_rect.y + get_global_id(1); - if(x < prefilter_rect.z && y < prefilter_rect.w) { - kernel_filter_get_feature(sample, - CCL_FILTER_TILE_INFO_ARG, - m_offset, v_offset, - x, y, - mean, variance, - scale, - prefilter_rect, - buffer_pass_stride, - buffer_denoising_offset); - } -} - -__kernel void kernel_ocl_filter_write_feature(int sample, - int4 buffer_params, - int4 filter_area, - ccl_global float *from, - ccl_global float *buffer, - int out_offset, - int4 prefilter_rect) -{ - int x = get_global_id(0); - int y = get_global_id(1); - if(x < filter_area.z && y < filter_area.w) { - kernel_filter_write_feature(sample, - x + filter_area.x, - y + filter_area.y, - buffer_params, - from, - buffer, - out_offset, - prefilter_rect); - } -} - -__kernel void kernel_ocl_filter_detect_outliers(ccl_global float *image, - ccl_global float *variance, - ccl_global float *depth, - ccl_global float *output, - int4 prefilter_rect, - int pass_stride) -{ - int x = prefilter_rect.x + get_global_id(0); - int y = prefilter_rect.y + get_global_id(1); - if(x < prefilter_rect.z && y < prefilter_rect.w) { - kernel_filter_detect_outliers(x, y, image, variance, depth, output, prefilter_rect, pass_stride); - } -} - -__kernel void kernel_ocl_filter_combine_halves(ccl_global float *mean, - ccl_global float *variance, - ccl_global float *a, - ccl_global float *b, - int4 prefilter_rect, - int r) -{ - int x = prefilter_rect.x + get_global_id(0); - int y = prefilter_rect.y + get_global_id(1); - if(x < prefilter_rect.z && y < prefilter_rect.w) { - kernel_filter_combine_halves(x, y, mean, variance, a, b, prefilter_rect, r); - } -} - -__kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_restrict buffer, - CCL_FILTER_TILE_INFO, - ccl_global float *transform, - ccl_global int *rank, - int4 filter_area, - int4 rect, - int pass_stride, - int frame_stride, - char use_time, - int radius, - float pca_threshold) -{ - int x = get_global_id(0); - int y = get_global_id(1); - if(x < filter_area.z && y < filter_area.w) { - ccl_global int *l_rank = rank + y*filter_area.z + x; - ccl_global float *l_transform = transform + y*filter_area.z + x; - kernel_filter_construct_transform(buffer, - CCL_FILTER_TILE_INFO_ARG, - x + filter_area.x, y + filter_area.y, - rect, - pass_stride, frame_stride, - use_time, - l_transform, l_rank, - radius, pca_threshold, - filter_area.z*filter_area.w, - get_local_id(1)*get_local_size(0) + get_local_id(0)); - } -} - -__kernel void kernel_ocl_filter_nlm_calc_difference(const ccl_global float *ccl_restrict weight_image, - const ccl_global float *ccl_restrict variance_image, - const ccl_global float *ccl_restrict scale_image, - ccl_global float *difference_image, - int w, - int h, - int stride, - int pass_stride, - int r, - int channel_offset, - int frame_offset, - float a, - float k_2) -{ - int4 co, rect; - int ofs; - if(get_nlm_coords(w, h, r, pass_stride, &rect, &co, &ofs)) { - kernel_filter_nlm_calc_difference(co.x, co.y, co.z, co.w, - weight_image, - variance_image, - scale_image, - difference_image + ofs, - rect, stride, - channel_offset, - frame_offset, - a, k_2); - } -} - -__kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict difference_image, - ccl_global float *out_image, - int w, - int h, - int stride, - int pass_stride, - int r, - int f) -{ - int4 co, rect; - int ofs; - if(get_nlm_coords(w, h, r, pass_stride, &rect, &co, &ofs)) { - kernel_filter_nlm_blur(co.x, co.y, - difference_image + ofs, - out_image + ofs, - rect, stride, f); - } -} - -__kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict difference_image, - ccl_global float *out_image, - int w, - int h, - int stride, - int pass_stride, - int r, - int f) -{ - int4 co, rect; - int ofs; - if(get_nlm_coords(w, h, r, pass_stride, &rect, &co, &ofs)) { - kernel_filter_nlm_calc_weight(co.x, co.y, - difference_image + ofs, - out_image + ofs, - rect, stride, f); - } -} - -__kernel void kernel_ocl_filter_nlm_update_output(const ccl_global float *ccl_restrict difference_image, - const ccl_global float *ccl_restrict image, - ccl_global float *out_image, - ccl_global float *accum_image, - int w, - int h, - int stride, - int pass_stride, - int channel_offset, - int r, - int f) -{ - int4 co, rect; - int ofs; - if(get_nlm_coords(w, h, r, pass_stride, &rect, &co, &ofs)) { - kernel_filter_nlm_update_output(co.x, co.y, co.z, co.w, - difference_image + ofs, - image, - out_image, - accum_image, - rect, - channel_offset, - stride, f); - } -} - -__kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *out_image, - const ccl_global float *ccl_restrict accum_image, - int w, - int h, - int stride) -{ - int x = get_global_id(0); - int y = get_global_id(1); - if(x < w && y < h) { - kernel_filter_nlm_normalize(x, y, out_image, accum_image, stride); - } -} - -__kernel void kernel_ocl_filter_nlm_construct_gramian(int t, - const ccl_global float *ccl_restrict difference_image, - const ccl_global float *ccl_restrict buffer, - const ccl_global float *ccl_restrict transform, - ccl_global int *rank, - ccl_global float *XtWX, - ccl_global float3 *XtWY, - int4 filter_window, - int w, - int h, - int stride, - int pass_stride, - int r, - int f, - int frame_offset, - char use_time) -{ - int4 co, rect; - int ofs; - if(get_nlm_coords_window(w, h, r, pass_stride, &rect, &co, &ofs, filter_window)) { - kernel_filter_nlm_construct_gramian(co.x, co.y, - co.z, co.w, - t, - difference_image + ofs, - buffer, - transform, rank, - XtWX, XtWY, - rect, filter_window, - stride, f, - pass_stride, - frame_offset, - use_time, - get_local_id(1)*get_local_size(0) + get_local_id(0)); - } -} - -__kernel void kernel_ocl_filter_finalize(ccl_global float *buffer, - ccl_global int *rank, - ccl_global float *XtWX, - ccl_global float3 *XtWY, - int4 filter_area, - int4 buffer_params, - int sample) -{ - int x = get_global_id(0); - int y = get_global_id(1); - if(x < filter_area.z && y < filter_area.w) { - int storage_ofs = y*filter_area.z+x; - rank += storage_ofs; - XtWX += storage_ofs; - XtWY += storage_ofs; - kernel_filter_finalize(x, y, buffer, rank, - filter_area.z*filter_area.w, - XtWX, XtWY, - buffer_params, sample); - } -} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl deleted file mode 100644 index ebdb99d4730..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_adaptive_adjust_samples.h" - -#define KERNEL_NAME adaptive_adjust_samples -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl deleted file mode 100644 index 76d82d4184e..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_adaptive_filter_x.h" - -#define KERNEL_NAME adaptive_filter_x -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl deleted file mode 100644 index 1e6d15ba0f2..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_adaptive_filter_y.h" - -#define KERNEL_NAME adaptive_filter_y -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl deleted file mode 100644 index 51de0059667..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright 2019 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_adaptive_stopping.h" - -#define KERNEL_NAME adaptive_stopping -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_background.cl deleted file mode 100644 index 0e600676e82..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_background.cl +++ /dev/null @@ -1,35 +0,0 @@ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/kernel_math.h" -#include "kernel/kernel_types.h" -#include "kernel/kernel_globals.h" -#include "kernel/kernel_color.h" -#include "kernel/kernels/opencl/kernel_opencl_image.h" - -#include "kernel/kernel_path.h" -#include "kernel/kernel_path_branched.h" - -#include "kernel/kernel_bake.h" - -__kernel void kernel_ocl_background( - ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, - - KERNEL_BUFFER_PARAMS, - - int type, int sx, int sw, int offset, int sample) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - - if(x < sx + sw) { - kernel_background_evaluate(kg, input, output, x); - } -} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_bake.cl b/intern/cycles/kernel/kernels/opencl/kernel_bake.cl deleted file mode 100644 index 7b81e387467..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_bake.cl +++ /dev/null @@ -1,36 +0,0 @@ -#include "kernel/kernel_compat_opencl.h" -#include "kernel/kernel_math.h" -#include "kernel/kernel_types.h" -#include "kernel/kernel_globals.h" -#include "kernel/kernel_color.h" -#include "kernel/kernels/opencl/kernel_opencl_image.h" - -#include "kernel/kernel_path.h" -#include "kernel/kernel_path_branched.h" - -#include "kernel/kernel_bake.h" - -__kernel void kernel_ocl_bake( - ccl_constant KernelData *data, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - int sx, int sy, int sw, int sh, int offset, int stride, int sample) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - - if(x < sx + sw && y < sy + sh) { -#ifndef __NO_BAKING__ - kernel_bake_evaluate(kg, buffer, sample, x, y, offset, stride); -#endif - } -} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_base.cl b/intern/cycles/kernel/kernels/opencl/kernel_base.cl deleted file mode 100644 index 1c2d89e8a92..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_base.cl +++ /dev/null @@ -1,88 +0,0 @@ -/* - * Copyright 2011-2013 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -/* OpenCL base kernels entry points */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/kernel_types.h" -#include "kernel/kernel_globals.h" - -#include "kernel/kernel_film.h" - - -__kernel void kernel_ocl_convert_to_byte( - ccl_constant KernelData *data, - ccl_global uchar4 *rgba, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - float sample_scale, - int sx, int sy, int sw, int sh, int offset, int stride) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - - if(x < sx + sw && y < sy + sh) - kernel_film_convert_to_byte(kg, rgba, buffer, sample_scale, x, y, offset, stride); -} - -__kernel void kernel_ocl_convert_to_half_float( - ccl_constant KernelData *data, - ccl_global uchar4 *rgba, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - float sample_scale, - int sx, int sy, int sw, int sh, int offset, int stride) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - - if(x < sx + sw && y < sy + sh) - kernel_film_convert_to_half_float(kg, rgba, buffer, sample_scale, x, y, offset, stride); -} - -__kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, uint64_t offset) -{ - size_t i = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0); - - if(i < size / sizeof(float4)) { - buffer[i+offset/sizeof(float4)] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - } - else if(i == size / sizeof(float4)) { - ccl_global uchar *b = (ccl_global uchar*)&buffer[i+offset/sizeof(float4)]; - - for(i = 0; i < size % sizeof(float4); i++) { - *(b++) = 0; - } - } -} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl deleted file mode 100644 index dcea2630aef..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_buffer_update.h" - -#define KERNEL_NAME buffer_update -#define LOCALS_TYPE unsigned int -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl deleted file mode 100644 index 7125348a49f..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl +++ /dev/null @@ -1,53 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_data_init.h" - -__kernel void kernel_ocl_path_trace_data_init( - ccl_global char *kg, - ccl_constant KernelData *data, - ccl_global void *split_data_buffer, - int num_elements, - ccl_global char *ray_state, - KERNEL_BUFFER_PARAMS, - int start_sample, - int end_sample, - int sx, int sy, int sw, int sh, int offset, int stride, - ccl_global int *Queue_index, /* Tracks the number of elements in queues */ - int queuesize, /* size (capacity) of the queue */ - ccl_global char *use_queues_flag, /* flag to decide if scene-intersect kernel should use queues to fetch ray index */ - ccl_global unsigned int *work_pool_wgs, /* Work pool for each work group */ - unsigned int num_samples, /* Total number of samples per pixel */ - ccl_global float *buffer) -{ - kernel_data_init((KernelGlobals*)kg, - data, - split_data_buffer, - num_elements, - ray_state, - KERNEL_BUFFER_ARGS, - start_sample, - end_sample, - sx, sy, sw, sh, offset, stride, - Queue_index, - queuesize, - use_queues_flag, - work_pool_wgs, - num_samples, - buffer); -} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl deleted file mode 100644 index ed64ae01aae..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_direct_lighting.h" - -#define KERNEL_NAME direct_lighting -#define LOCALS_TYPE unsigned int -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_displace.cl b/intern/cycles/kernel/kernels/opencl/kernel_displace.cl deleted file mode 100644 index 76cc36971f5..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_displace.cl +++ /dev/null @@ -1,36 +0,0 @@ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/kernel_math.h" -#include "kernel/kernel_types.h" -#include "kernel/kernel_globals.h" -#include "kernel/kernel_color.h" -#include "kernel/kernels/opencl/kernel_opencl_image.h" - -#include "kernel/kernel_path.h" -#include "kernel/kernel_path_branched.h" - -#include "kernel/kernel_bake.h" - -__kernel void kernel_ocl_displace( - ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, - - KERNEL_BUFFER_PARAMS, - - int type, int sx, int sw, int offset, int sample) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - - if(x < sx + sw) { - kernel_displace_evaluate(kg, input, output, x); - } -} - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl deleted file mode 100644 index 8afaa686e28..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_do_volume.h" - -#define KERNEL_NAME do_volume -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl b/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl deleted file mode 100644 index e68d4104a91..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_enqueue_inactive.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_enqueue_inactive.h" - -#define KERNEL_NAME enqueue_inactive -#define LOCALS_TYPE unsigned int -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl deleted file mode 100644 index 9e1e57beba6..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h" - -#define KERNEL_NAME holdout_emission_blurring_pathtermination_ao -#define LOCALS_TYPE BackgroundAOLocals -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl deleted file mode 100644 index 192d01444ba..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_indirect_background.h" - -#define KERNEL_NAME indirect_background -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl deleted file mode 100644 index 84938b889e5..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_indirect_subsurface.h" - -#define KERNEL_NAME indirect_subsurface -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl deleted file mode 100644 index c314dc96c33..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_lamp_emission.h" - -#define KERNEL_NAME lamp_emission -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl deleted file mode 100644 index 8b1332bf013..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_next_iteration_setup.h" - -#define KERNEL_NAME next_iteration_setup -#define LOCALS_TYPE unsigned int -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h b/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h deleted file mode 100644 index bb6b8a40e8e..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_opencl_image.h +++ /dev/null @@ -1,358 +0,0 @@ -/* - * Copyright 2016 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifdef WITH_NANOVDB -/* Data type to replace `double` used in the NanoVDB headers. Cycles don't need doubles, and is - * safer and more portable to never use double datatype on GPU. - * Use a special structure, so that the following is true: - * - No unnoticed implicit cast or mathematical operations used on scalar 64bit type - * (which rules out trick like using `uint64_t` as a drop-in replacement for double). - * - Padding rules are matching exactly `double` - * (which rules out array of `uint8_t`). */ -typedef struct ccl_vdb_double_t { - uint64_t i; -} ccl_vdb_double_t; - -# define double ccl_vdb_double_t -# include "nanovdb/CNanoVDB.h" -# undef double -#endif - -/* For OpenCL we do manual lookup and interpolation. */ - -ccl_device_inline ccl_global TextureInfo *kernel_tex_info(KernelGlobals *kg, uint id) -{ - const uint tex_offset = id -#define KERNEL_TEX(type, name) +1 -#include "kernel/kernel_textures.h" - ; - - return &((ccl_global TextureInfo *)kg->buffers[0])[tex_offset]; -} - -#define tex_fetch(type, info, index) \ - ((ccl_global type *)(kg->buffers[info->cl_buffer] + info->data))[(index)] - -ccl_device_inline int svm_image_texture_wrap_periodic(int x, int width) -{ - x %= width; - if (x < 0) - x += width; - return x; -} - -ccl_device_inline int svm_image_texture_wrap_clamp(int x, int width) -{ - return clamp(x, 0, width - 1); -} - -ccl_device_inline float4 svm_image_texture_read( - KernelGlobals *kg, const ccl_global TextureInfo *info, void *acc, int x, int y, int z) -{ - const int data_offset = x + info->width * y + info->width * info->height * z; - const int texture_type = info->data_type; - - /* Float4 */ - if (texture_type == IMAGE_DATA_TYPE_FLOAT4) { - return tex_fetch(float4, info, data_offset); - } - /* Byte4 */ - else if (texture_type == IMAGE_DATA_TYPE_BYTE4) { - uchar4 r = tex_fetch(uchar4, info, data_offset); - float f = 1.0f / 255.0f; - return make_float4(r.x * f, r.y * f, r.z * f, r.w * f); - } - /* Ushort4 */ - else if (texture_type == IMAGE_DATA_TYPE_USHORT4) { - ushort4 r = tex_fetch(ushort4, info, data_offset); - float f = 1.0f / 65535.f; - return make_float4(r.x * f, r.y * f, r.z * f, r.w * f); - } - /* Float */ - else if (texture_type == IMAGE_DATA_TYPE_FLOAT) { - float f = tex_fetch(float, info, data_offset); - return make_float4(f, f, f, 1.0f); - } - /* UShort */ - else if (texture_type == IMAGE_DATA_TYPE_USHORT) { - ushort r = tex_fetch(ushort, info, data_offset); - float f = r * (1.0f / 65535.0f); - return make_float4(f, f, f, 1.0f); - } -#ifdef WITH_NANOVDB - /* NanoVDB Float */ - else if (texture_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT) { - cnanovdb_coord coord; - coord.mVec[0] = x; - coord.mVec[1] = y; - coord.mVec[2] = z; - float f = cnanovdb_readaccessor_getValueF((cnanovdb_readaccessor *)acc, &coord); - return make_float4(f, f, f, 1.0f); - } - /* NanoVDB Float3 */ - else if (texture_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT3) { - cnanovdb_coord coord; - coord.mVec[0] = x; - coord.mVec[1] = y; - coord.mVec[2] = z; - cnanovdb_Vec3F f = cnanovdb_readaccessor_getValueF3((cnanovdb_readaccessor *)acc, &coord); - return make_float4(f.mVec[0], f.mVec[1], f.mVec[2], 1.0f); - } -#endif -#ifdef __KERNEL_CL_KHR_FP16__ - /* Half and Half4 are optional in OpenCL */ - else if (texture_type == IMAGE_DATA_TYPE_HALF) { - float f = tex_fetch(half, info, data_offset); - return make_float4(f, f, f, 1.0f); - } - else if (texture_type == IMAGE_DATA_TYPE_HALF4) { - half4 r = tex_fetch(half4, info, data_offset); - return make_float4(r.x, r.y, r.z, r.w); - } -#endif - /* Byte */ - else { - uchar r = tex_fetch(uchar, info, data_offset); - float f = r * (1.0f / 255.0f); - return make_float4(f, f, f, 1.0f); - } -} - -ccl_device_inline float4 -svm_image_texture_read_2d(KernelGlobals *kg, int id, void *acc, int x, int y) -{ - const ccl_global TextureInfo *info = kernel_tex_info(kg, id); - -#ifdef WITH_NANOVDB - if (info->data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT && - info->data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) { -#endif - /* Wrap */ - if (info->extension == EXTENSION_REPEAT) { - x = svm_image_texture_wrap_periodic(x, info->width); - y = svm_image_texture_wrap_periodic(y, info->height); - } - else { - x = svm_image_texture_wrap_clamp(x, info->width); - y = svm_image_texture_wrap_clamp(y, info->height); - } -#ifdef WITH_NANOVDB - } -#endif - - return svm_image_texture_read(kg, info, acc, x, y, 0); -} - -ccl_device_inline float4 -svm_image_texture_read_3d(KernelGlobals *kg, int id, void *acc, int x, int y, int z) -{ - const ccl_global TextureInfo *info = kernel_tex_info(kg, id); - -#ifdef WITH_NANOVDB - if (info->data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT && - info->data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) { -#endif - /* Wrap */ - if (info->extension == EXTENSION_REPEAT) { - x = svm_image_texture_wrap_periodic(x, info->width); - y = svm_image_texture_wrap_periodic(y, info->height); - z = svm_image_texture_wrap_periodic(z, info->depth); - } - else { - x = svm_image_texture_wrap_clamp(x, info->width); - y = svm_image_texture_wrap_clamp(y, info->height); - z = svm_image_texture_wrap_clamp(z, info->depth); - } -#ifdef WITH_NANOVDB - } -#endif - - return svm_image_texture_read(kg, info, acc, x, y, z); -} - -ccl_device_inline float svm_image_texture_frac(float x, int *ix) -{ - int i = float_to_int(x) - ((x < 0.0f) ? 1 : 0); - *ix = i; - return x - (float)i; -} - -#define SET_CUBIC_SPLINE_WEIGHTS(u, t) \ - { \ - u[0] = (((-1.0f / 6.0f) * t + 0.5f) * t - 0.5f) * t + (1.0f / 6.0f); \ - u[1] = ((0.5f * t - 1.0f) * t) * t + (2.0f / 3.0f); \ - u[2] = ((-0.5f * t + 0.5f) * t + 0.5f) * t + (1.0f / 6.0f); \ - u[3] = (1.0f / 6.0f) * t * t * t; \ - } \ - (void)0 - -ccl_device float4 kernel_tex_image_interp(KernelGlobals *kg, int id, float x, float y) -{ - const ccl_global TextureInfo *info = kernel_tex_info(kg, id); - - if (info->extension == EXTENSION_CLIP) { - if (x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) { - return make_float4(0.0f, 0.0f, 0.0f, 0.0f); - } - } - - if (info->interpolation == INTERPOLATION_CLOSEST) { - /* Closest interpolation. */ - int ix, iy; - svm_image_texture_frac(x * info->width, &ix); - svm_image_texture_frac(y * info->height, &iy); - - return svm_image_texture_read_2d(kg, id, NULL, ix, iy); - } - else if (info->interpolation == INTERPOLATION_LINEAR) { - /* Bilinear interpolation. */ - int ix, iy; - float tx = svm_image_texture_frac(x * info->width - 0.5f, &ix); - float ty = svm_image_texture_frac(y * info->height - 0.5f, &iy); - - float4 r; - r = (1.0f - ty) * (1.0f - tx) * svm_image_texture_read_2d(kg, id, NULL, ix, iy); - r += (1.0f - ty) * tx * svm_image_texture_read_2d(kg, id, NULL, ix + 1, iy); - r += ty * (1.0f - tx) * svm_image_texture_read_2d(kg, id, NULL, ix, iy + 1); - r += ty * tx * svm_image_texture_read_2d(kg, id, NULL, ix + 1, iy + 1); - return r; - } - else { - /* Bicubic interpolation. */ - int ix, iy; - float tx = svm_image_texture_frac(x * info->width - 0.5f, &ix); - float ty = svm_image_texture_frac(y * info->height - 0.5f, &iy); - - float u[4], v[4]; - SET_CUBIC_SPLINE_WEIGHTS(u, tx); - SET_CUBIC_SPLINE_WEIGHTS(v, ty); - - float4 r = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - - for (int y = 0; y < 4; y++) { - for (int x = 0; x < 4; x++) { - float weight = u[x] * v[y]; - r += weight * svm_image_texture_read_2d(kg, id, NULL, ix + x - 1, iy + y - 1); - } - } - return r; - } -} - -ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals *kg, int id, float3 P, int interp) -{ - const ccl_global TextureInfo *info = kernel_tex_info(kg, id); - - if (info->use_transform_3d) { - Transform tfm = info->transform_3d; - P = transform_point(&tfm, P); - } - - float x = P.x; - float y = P.y; - float z = P.z; - - uint interpolation = (interp == INTERPOLATION_NONE) ? info->interpolation : interp; - -#ifdef WITH_NANOVDB - cnanovdb_readaccessor acc; - if (info->data_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT || - info->data_type == IMAGE_DATA_TYPE_NANOVDB_FLOAT3) { - ccl_global cnanovdb_griddata *grid = - (ccl_global cnanovdb_griddata *)(kg->buffers[info->cl_buffer] + info->data); - cnanovdb_readaccessor_init(&acc, cnanovdb_treedata_rootF(cnanovdb_griddata_tree(grid))); - } - else { - if (info->extension == EXTENSION_CLIP) { - if (x < 0.0f || y < 0.0f || z < 0.0f || x > 1.0f || y > 1.0f || z > 1.0f) { - return make_float4(0.0f, 0.0f, 0.0f, 0.0f); - } - } - - x *= info->width; - y *= info->height; - z *= info->depth; - } -# define NANOVDB_ACCESS_POINTER &acc -#else -# define NANOVDB_ACCESS_POINTER NULL -#endif - - if (interpolation == INTERPOLATION_CLOSEST) { - /* Closest interpolation. */ - int ix, iy, iz; - svm_image_texture_frac(x, &ix); - svm_image_texture_frac(y, &iy); - svm_image_texture_frac(z, &iz); - - return svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix, iy, iz); - } - else if (interpolation == INTERPOLATION_LINEAR) { - /* Trilinear interpolation. */ - int ix, iy, iz; - float tx = svm_image_texture_frac(x - 0.5f, &ix); - float ty = svm_image_texture_frac(y - 0.5f, &iy); - float tz = svm_image_texture_frac(z - 0.5f, &iz); - - float4 r; - r = (1.0f - tz) * (1.0f - ty) * (1.0f - tx) * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix, iy, iz); - r += (1.0f - tz) * (1.0f - ty) * tx * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix + 1, iy, iz); - r += (1.0f - tz) * ty * (1.0f - tx) * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix, iy + 1, iz); - r += (1.0f - tz) * ty * tx * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix + 1, iy + 1, iz); - - r += tz * (1.0f - ty) * (1.0f - tx) * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix, iy, iz + 1); - r += tz * (1.0f - ty) * tx * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix + 1, iy, iz + 1); - r += tz * ty * (1.0f - tx) * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix, iy + 1, iz + 1); - r += tz * ty * tx * - svm_image_texture_read_3d(kg, id, NANOVDB_ACCESS_POINTER, ix + 1, iy + 1, iz + 1); - return r; - } - else { - /* Tricubic interpolation. */ - int ix, iy, iz; - float tx = svm_image_texture_frac(x - 0.5f, &ix); - float ty = svm_image_texture_frac(y - 0.5f, &iy); - float tz = svm_image_texture_frac(z - 0.5f, &iz); - - float u[4], v[4], w[4]; - SET_CUBIC_SPLINE_WEIGHTS(u, tx); - SET_CUBIC_SPLINE_WEIGHTS(v, ty); - SET_CUBIC_SPLINE_WEIGHTS(w, tz); - - float4 r = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - - for (int z = 0; z < 4; z++) { - for (int y = 0; y < 4; y++) { - for (int x = 0; x < 4; x++) { - float weight = u[x] * v[y] * w[z]; - r += weight * svm_image_texture_read_3d( - kg, id, NANOVDB_ACCESS_POINTER, ix + x - 1, iy + y - 1, iz + z - 1); - } - } - } - return r; - } -#undef NANOVDB_ACCESS_POINTER -} - -#undef SET_CUBIC_SPLINE_WEIGHTS diff --git a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl deleted file mode 100644 index fa210e747c0..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_path_init.h" - -#define KERNEL_NAME path_init -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl deleted file mode 100644 index 68ee6f1d536..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_queue_enqueue.h" - -#define KERNEL_NAME queue_enqueue -#define LOCALS_TYPE QueueEnqueueLocals -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl deleted file mode 100644 index 10d09377ba9..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_scene_intersect.h" - -#define KERNEL_NAME scene_intersect -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl deleted file mode 100644 index 40eaa561863..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_shader_eval.h" - -#define KERNEL_NAME shader_eval -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl deleted file mode 100644 index 8c36100f762..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_setup.cl +++ /dev/null @@ -1,26 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_shader_setup.h" - -#define KERNEL_NAME shader_setup -#define LOCALS_TYPE unsigned int -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl deleted file mode 100644 index bcacaa4a054..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl +++ /dev/null @@ -1,27 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_shader_sort.h" - -__attribute__((reqd_work_group_size(64, 1, 1))) -#define KERNEL_NAME shader_sort -#define LOCALS_TYPE ShaderSortLocals -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME -#undef LOCALS_TYPE - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl deleted file mode 100644 index 8de250a375c..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_shadow_blocked_ao.h" - -#define KERNEL_NAME shadow_blocked_ao -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl deleted file mode 100644 index 29da77022ed..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2015 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_shadow_blocked_dl.h" - -#define KERNEL_NAME shadow_blocked_dl -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl b/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl deleted file mode 100644 index c3b7b09460a..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl +++ /dev/null @@ -1,34 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" // PRECOMPILED -#include "kernel/split/kernel_split_common.h" // PRECOMPILED - -#include "kernel/kernels/opencl/kernel_data_init.cl" -#include "kernel/kernels/opencl/kernel_path_init.cl" -#include "kernel/kernels/opencl/kernel_state_buffer_size.cl" -#include "kernel/kernels/opencl/kernel_scene_intersect.cl" -#include "kernel/kernels/opencl/kernel_queue_enqueue.cl" -#include "kernel/kernels/opencl/kernel_shader_setup.cl" -#include "kernel/kernels/opencl/kernel_shader_sort.cl" -#include "kernel/kernels/opencl/kernel_enqueue_inactive.cl" -#include "kernel/kernels/opencl/kernel_next_iteration_setup.cl" -#include "kernel/kernels/opencl/kernel_indirect_subsurface.cl" -#include "kernel/kernels/opencl/kernel_buffer_update.cl" -#include "kernel/kernels/opencl/kernel_adaptive_stopping.cl" -#include "kernel/kernels/opencl/kernel_adaptive_filter_x.cl" -#include "kernel/kernels/opencl/kernel_adaptive_filter_y.cl" -#include "kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl" diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h b/intern/cycles/kernel/kernels/opencl/kernel_split_function.h deleted file mode 100644 index e123b4cd6ec..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_split_function.h +++ /dev/null @@ -1,67 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#define KERNEL_NAME_JOIN(a, b) a##_##b -#define KERNEL_NAME_EVAL(a, b) KERNEL_NAME_JOIN(a, b) - -__kernel void KERNEL_NAME_EVAL(kernel_ocl_path_trace, - KERNEL_NAME)(ccl_global char *kg_global, - ccl_constant KernelData *data, - - ccl_global void *split_data_buffer, - ccl_global char *ray_state, - - KERNEL_BUFFER_PARAMS, - - ccl_global int *queue_index, - ccl_global char *use_queues_flag, - ccl_global unsigned int *work_pools, - ccl_global float *buffer) -{ -#ifdef LOCALS_TYPE - ccl_local LOCALS_TYPE locals; -#endif - - KernelGlobals *kg = (KernelGlobals *)kg_global; - - if (ccl_local_id(0) + ccl_local_id(1) == 0) { - kg->data = data; - - kernel_split_params.queue_index = queue_index; - kernel_split_params.use_queues_flag = use_queues_flag; - kernel_split_params.work_pools = work_pools; - kernel_split_params.tile.buffer = buffer; - - split_data_init(kg, - &kernel_split_state, - ccl_global_size(0) * ccl_global_size(1), - split_data_buffer, - ray_state); - } - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - - KERNEL_NAME_EVAL(kernel, KERNEL_NAME) - (kg -#ifdef LOCALS_TYPE - , - &locals -#endif - ); -} - -#undef KERNEL_NAME_JOIN -#undef KERNEL_NAME_EVAL diff --git a/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl b/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl deleted file mode 100644 index c10ecc426c6..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl +++ /dev/null @@ -1,29 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" - -__kernel void kernel_ocl_path_trace_state_buffer_size( - ccl_global char *kg, - ccl_constant KernelData *data, - uint num_threads, - ccl_global uint64_t *size) -{ - ((KernelGlobals*)kg)->data = data; - *size = split_data_buffer_size((KernelGlobals*)kg, num_threads); -} - diff --git a/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl b/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl deleted file mode 100644 index 2b3be38df84..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright 2011-2017 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/split/kernel_split_common.h" -#include "kernel/split/kernel_subsurface_scatter.h" - -#define KERNEL_NAME subsurface_scatter -#include "kernel/kernels/opencl/kernel_split_function.h" -#undef KERNEL_NAME - |