Changeset View
Changeset View
Standalone View
Standalone View
intern/cycles/kernel/kernels/cuda/kernel.cu
| Show First 20 Lines • Show All 123 Lines • ▼ Show 20 Lines | if(x < sx + sw) { | ||||
| KernelGlobals kg; | KernelGlobals kg; | ||||
| kernel_background_evaluate(&kg, input, output, x); | kernel_background_evaluate(&kg, input, output, x); | ||||
| } | } | ||||
| } | } | ||||
| #ifdef __BAKING__ | #ifdef __BAKING__ | ||||
| extern "C" __global__ void | extern "C" __global__ void | ||||
| CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) | 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; | 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 | #endif | ||||
| #endif | #endif | ||||