diff options
Diffstat (limited to 'intern/cycles/kernel/kernels')
| -rw-r--r-- | intern/cycles/kernel/kernels/cpu/kernel_cpu.h | 3 | ||||
| -rw-r--r-- | intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h | 19 | ||||
| -rw-r--r-- | intern/cycles/kernel/kernels/cuda/kernel.cu | 11 |
3 files changed, 23 insertions, 10 deletions
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h index 683f4b88d79..ea3103f12c3 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h @@ -46,6 +46,9 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, int offset, int sample); +void KERNEL_FUNCTION_FULL_NAME(bake)( + KernelGlobals *kg, float *buffer, int sample, int x, int y, int offset, int stride); + /* Split kernels */ void KERNEL_FUNCTION_FULL_NAME(data_init)(KernelGlobals *kg, diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h index 091e53cfd83..5aa3fb14318 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h @@ -132,6 +132,18 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg, # endif /* KERNEL_STUB */ } +/* Bake */ + +void KERNEL_FUNCTION_FULL_NAME(bake)( + KernelGlobals *kg, float *buffer, int sample, int x, int y, int offset, int stride) +{ +# ifdef KERNEL_STUB + STUB_ASSERT(KERNEL_ARCH, bake); +# else + kernel_bake_evaluate(kg, buffer, sample, x, y, offset, stride); +# endif /* KERNEL_STUB */ +} + /* Shader Evaluate */ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, @@ -146,12 +158,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg, # ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, shader); # else - if (type >= SHADER_EVAL_BAKE) { -# ifdef __BAKING__ - kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, i, offset, sample); -# endif - } - else if (type == SHADER_EVAL_DISPLACE) { + if (type == SHADER_EVAL_DISPLACE) { kernel_displace_evaluate(kg, input, output, i); } else { diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu index c4c810c6a82..d4f41132a11 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel.cu @@ -214,13 +214,16 @@ kernel_cuda_background(uint4 *input, #ifdef __BAKING__ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_bake(uint4 *input, float4 *output, int type, int filter, int sx, int sw, int offset, int sample) +kernel_cuda_bake(WorkTile *tile, uint total_work_size) { - int x = sx + blockDim.x*blockIdx.x + threadIdx.x; + int work_index = ccl_global_id(0); + + if(work_index < total_work_size) { + uint x, y, sample; + get_work_pixel(tile, work_index, &x, &y, &sample); - if(x < sx + sw) { KernelGlobals kg; - kernel_bake_evaluate(&kg, input, output, (ShaderEvalType)type, filter, x, offset, sample); + kernel_bake_evaluate(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride); } } #endif |