diff options
| author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2014-06-05 20:10:06 +0400 |
|---|---|---|
| committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2014-06-05 20:10:32 +0400 |
| commit | 865dfa8a7e8c29b2a7f0aa3e6bcdfb3353e39052 (patch) | |
| tree | 63d3f2a4ab983be993ec06a2fd1907f4fcbb791d /intern/cycles/kernel/kernel.cu | |
| parent | 2305e3289bcfced0e56d4d0fac8d525cb847d46e (diff) | |
Fix T40228: cycles CUDA multi GPU + world MIS giving error.
Diffstat (limited to 'intern/cycles/kernel/kernel.cu')
| -rw-r--r-- | intern/cycles/kernel/kernel.cu | 10 |
1 files changed, 6 insertions, 4 deletions
diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index 69c2600ea98..12273cc2853 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -146,20 +146,22 @@ kernel_cuda_convert_to_half_float(uchar4 *rgba, float *buffer, float sample_scal extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx) +kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx, int sw) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; - kernel_shader_evaluate(NULL, input, output, (ShaderEvalType)type, x); + if(x < sx + sw) + kernel_shader_evaluate(NULL, input, output, (ShaderEvalType)type, x); } 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 sx) +kernel_cuda_bake(uint4 *input, float4 *output, int type, int sx, int sw) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; - kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, x); + if(x < sx + sw) + kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, x); } #endif |