Differential D4383 Diff 13771 intern/cycles/kernel/kernels/opencl/kernel_base.cl

Changeset View

Standalone View

intern/cycles/kernel/kernels/opencl/kernel_base.cl

This file was moved from intern/cycles/kernel/kernels/opencl/kernel.cl.

/*		/*
* Copyright 2011-2013 Blender Foundation		* Copyright 2011-2013 Blender Foundation
*		*
* Licensed under the Apache License, Version 2.0 (the "License");		* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.		* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at		* You may obtain a copy of the License at
*		*
* http://www.apache.org/licenses/LICENSE-2.0		* http://www.apache.org/licenses/LICENSE-2.0
*		*
* Unless required by applicable law or agreed to in writing, software		* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,		* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.		* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and		* See the License for the specific language governing permissions and
* limitations under the License.		* limitations under the License.
*/		*/

/* OpenCL kernel entry points - unfinished */		/* OpenCL base kernels entry points */

#include "kernel/kernel_compat_opencl.h"		#include "kernel/kernel_compat_opencl.h"
#include "kernel/kernel_math.h"
#include "kernel/kernel_types.h"		#include "kernel/kernel_types.h"
#include "kernel/kernel_globals.h"		#include "kernel/kernel_globals.h"
#include "kernel/kernel_color.h"
#include "kernel/kernels/opencl/kernel_opencl_image.h"

#include "kernel/kernel_film.h"		#include "kernel/kernel_film.h"

#if defined(__COMPILE_ONLY_MEGAKERNEL__) \|\| !defined(__NO_BAKING__)
# include "kernel/kernel_path.h"
# include "kernel/kernel_path_branched.h"
#else /* __COMPILE_ONLY_MEGAKERNEL__ */
/* Include only actually used headers for the case
* when path tracing kernels are not needed.
*/
# include "kernel/kernel_random.h"
# include "kernel/kernel_differential.h"
# include "kernel/kernel_montecarlo.h"
# include "kernel/kernel_projection.h"
# include "kernel/geom/geom.h"
# include "kernel/bvh/bvh.h"

# include "kernel/kernel_accumulate.h"
# include "kernel/kernel_camera.h"
# include "kernel/kernel_shader.h"
#endif /* defined(__COMPILE_ONLY_MEGAKERNEL__) \|\| !defined(__NO_BAKING__) */

#include "kernel/kernel_bake.h"

#ifdef __COMPILE_ONLY_MEGAKERNEL__

__kernel void kernel_ocl_path_trace(
ccl_constant KernelData *data,
ccl_global float *buffer,

KERNEL_BUFFER_PARAMS,

int sample,
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);
bool thread_is_active = x < sx + sw && y < sy + sh;
if(thread_is_active) {
kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
}
if(kernel_data.film.cryptomatte_passes) {
/* Make sure no thread is writing to the buffers. */
ccl_barrier(CCL_LOCAL_MEM_FENCE);
if(thread_is_active) {
kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride);
}
}
}

#else /* __COMPILE_ONLY_MEGAKERNEL__ */

__kernel void kernel_ocl_convert_to_byte(		__kernel void kernel_ocl_convert_to_byte(
ccl_constant KernelData *data,		ccl_constant KernelData *data,
ccl_global uchar4 *rgba,		ccl_global uchar4 *rgba,
ccl_global float *buffer,		ccl_global float *buffer,

KERNEL_BUFFER_PARAMS,		KERNEL_BUFFER_PARAMS,

▲ Show 20 Lines • Show All 48 Lines • ▼ Show 20 Lines	__kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, uint64_t offset)
else if(i == size / sizeof(float4)) {		else if(i == size / sizeof(float4)) {
ccl_global uchar b = (ccl_global uchar)&buffer[i+offset/sizeof(float4)];		ccl_global uchar b = (ccl_global uchar)&buffer[i+offset/sizeof(float4)];

for(i = 0; i < size % sizeof(float4); i++) {		for(i = 0; i < size % sizeof(float4); i++) {
*(b++) = 0;		*(b++) = 0;
}		}
}		}
}		}

#endif /* __COMPILE_ONLY_MEGAKERNEL__ */