blender
/
blender
129
398
Fork 572
Code Issues 6.2k Pull Requests 745 Projects 19 Wiki Activity

Fix T40228: cycles CUDA multi GPU + world MIS giving error.

This commit is contained in:
Brecht Van Lommel 2014-06-05 18:10:06 +02:00
parent 2305e3289b
commit 865dfa8a7e
Notes: blender-bot 2023-02-14 10:37:51 +01:00 
Referenced by issue #40228, World multiple importance sample CUDA error
2 changed files with 11 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
intern/cycles/device/device_cuda.cpp
View File

@ -762,6 +762,8 @@ public:
if(task.get_cancel())
break;
int shader_w = min(shader_chunk_size, end - shader_x);
/* pass in parameters */
int offset = 0;
@ -780,13 +782,15 @@ public:
cuda_assert(cuParamSeti(cuShader, offset, shader_x));
offset += sizeof(shader_x);
cuda_assert(cuParamSeti(cuShader, offset, shader_w));
offset += sizeof(shader_w);
cuda_assert(cuParamSetSize(cuShader, offset));
/* launch kernel */
int threads_per_block;
cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuShader));
int shader_w = min(shader_chunk_size, end - shader_x);
int xblocks = (shader_w + threads_per_block - 1)/threads_per_block;
cuda_assert(cuFuncSetCacheConfig(cuShader, CU_FUNC_CACHE_PREFER_L1));

10
intern/cycles/kernel/kernel.cu
View File

@ -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