diff options
author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2014-05-27 15:20:07 +0400 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2014-05-27 17:11:32 +0400 |
commit | 69c7522b2463245ef16ebcf2806645c78e83b4df (patch) | |
tree | 96721c9d2bc755d68cf9eeded77255f4c58018f3 /intern/cycles/kernel | |
parent | bc9e66f0830627e807f720dca4b9d5d8d39e732a (diff) |
Fix T40379: world MIS causing too much CUDA memory usage.
The kernel for baking the world texture was the same as the one used for
baking. Now that's separate which allows the kernel to reserve much less
memory.
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/kernel.cl | 25 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.cpp | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel.cu | 9 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_avx.cpp | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_bake.h | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_sse2.cpp | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_sse3.cpp | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_sse41.cpp | 5 |
8 files changed, 54 insertions, 10 deletions
diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl index 1dc0793a7bc..d7d3438036e 100644 --- a/intern/cycles/kernel/kernel.cl +++ b/intern/cycles/kernel/kernel.cl @@ -131,3 +131,28 @@ __kernel void kernel_ocl_shader( kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, x); } +__kernel void kernel_ocl_bake( + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, + +#define KERNEL_TEX(type, ttype, name) \ + ccl_global type *name, +#include "kernel_textures.h" + + int type, int sx, int sw) +{ + KernelGlobals kglobals, *kg = &kglobals; + + kg->data = data; + +#define KERNEL_TEX(type, ttype, name) \ + kg->name = name; +#include "kernel_textures.h" + + int x = sx + get_global_id(0); + + if(x < sx + sw) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, x); +} + diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index a0b6b8e13d0..a535659b3b1 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -122,7 +122,10 @@ void kernel_cpu_convert_to_half_float(KernelGlobals *kg, uchar4 *rgba, float *bu void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { - kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); + if(type >= SHADER_EVAL_BAKE) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, i); + else + kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index b9b41f755ba..bb20819f6fc 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -153,5 +153,14 @@ kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx) 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) +{ + int x = sx + blockDim.x*blockIdx.x + threadIdx.x; + + kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, x); +} + #endif diff --git a/intern/cycles/kernel/kernel_avx.cpp b/intern/cycles/kernel/kernel_avx.cpp index c572fcd3df6..7d354de16d2 100644 --- a/intern/cycles/kernel/kernel_avx.cpp +++ b/intern/cycles/kernel/kernel_avx.cpp @@ -69,7 +69,10 @@ void kernel_cpu_avx_convert_to_half_float(KernelGlobals *kg, uchar4 *rgba, float void kernel_cpu_avx_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { - kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); + if(type >= SHADER_EVAL_BAKE) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, i); + else + kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h index 6e735517763..c3ae2b6a54e 100644 --- a/intern/cycles/kernel/kernel_bake.h +++ b/intern/cycles/kernel/kernel_bake.h @@ -356,11 +356,6 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_device void kernel_shader_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output, ShaderEvalType type, int i) { - if(type >= SHADER_EVAL_BAKE) { - kernel_bake_evaluate(kg, input, output, type, i); - return; - } - ShaderData sd; uint4 in = input[i]; float3 out; diff --git a/intern/cycles/kernel/kernel_sse2.cpp b/intern/cycles/kernel/kernel_sse2.cpp index 455cac046a3..3b5faea2994 100644 --- a/intern/cycles/kernel/kernel_sse2.cpp +++ b/intern/cycles/kernel/kernel_sse2.cpp @@ -66,7 +66,10 @@ void kernel_cpu_sse2_convert_to_half_float(KernelGlobals *kg, uchar4 *rgba, floa void kernel_cpu_sse2_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { - kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); + if(type >= SHADER_EVAL_BAKE) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, i); + else + kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_sse3.cpp b/intern/cycles/kernel/kernel_sse3.cpp index 29aca52890e..3b18b164ffd 100644 --- a/intern/cycles/kernel/kernel_sse3.cpp +++ b/intern/cycles/kernel/kernel_sse3.cpp @@ -68,7 +68,10 @@ void kernel_cpu_sse3_convert_to_half_float(KernelGlobals *kg, uchar4 *rgba, floa void kernel_cpu_sse3_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { - kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); + if(type >= SHADER_EVAL_BAKE) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, i); + else + kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_sse41.cpp b/intern/cycles/kernel/kernel_sse41.cpp index 0ece67e6d2b..a3731d790f4 100644 --- a/intern/cycles/kernel/kernel_sse41.cpp +++ b/intern/cycles/kernel/kernel_sse41.cpp @@ -69,7 +69,10 @@ void kernel_cpu_sse41_convert_to_half_float(KernelGlobals *kg, uchar4 *rgba, flo void kernel_cpu_sse41_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i) { - kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); + if(type >= SHADER_EVAL_BAKE) + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, i); + else + kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i); } CCL_NAMESPACE_END |