From 743a7a4a4b2b6999b911704b6367962fc365474a Mon Sep 17 00:00:00 2001 From: Thomas Dinges Date: Fri, 9 Aug 2013 20:03:49 +0000 Subject: Cycles: * GPU kernel can now be compiled without __NON_PROGRESSIVE__ again, was broken after my last commit. Also add a check for have_error(), in case the GPU kernel comes without Non-Progressive, to avoid a crash. * Don't compile progressive kernel twice on CPU, if __NON_PROGRESSIVE__ would be disabled there. --- intern/cycles/device/device_cuda.cpp | 5 ++++- intern/cycles/kernel/kernel.cpp | 8 +++++--- intern/cycles/kernel/kernel.cu | 2 ++ intern/cycles/kernel/kernel_path.h | 6 ++---- intern/cycles/kernel/kernel_sse2.cpp | 8 +++++--- intern/cycles/kernel/kernel_sse3.cpp | 8 +++++--- 6 files changed, 23 insertions(+), 14 deletions(-) diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index 8f11782b168..ec7157aa083 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -572,8 +572,11 @@ public: /* get kernel function */ if(progressive) cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_progressive")) - else + else { cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace_non_progressive")) + if(have_error()) + return; + } /* pass in parameters */ int offset = 0; diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp index 88a8b2aa569..589c9e1e073 100644 --- a/intern/cycles/kernel/kernel.cpp +++ b/intern/cycles/kernel/kernel.cpp @@ -90,10 +90,12 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) { - if(kernel_data.integrator.progressive) - kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); - else +#ifdef __NON_PROGRESSIVE__ + if(!kernel_data.integrator.progressive) kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); + else +#endif + kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); } /* Tonemapping */ diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu index e3fe23d3a37..df865aa1b7e 100644 --- a/intern/cycles/kernel/kernel.cu +++ b/intern/cycles/kernel/kernel.cu @@ -35,6 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace_progressive(float *buffer, uin kernel_path_trace_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride); } +#ifdef __NON_PROGRESSIVE__ extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; @@ -43,6 +44,7 @@ extern "C" __global__ void kernel_cuda_path_trace_non_progressive(float *buffer, if(x < sx + sw && y < sy + sh) kernel_path_trace_non_progressive(NULL, buffer, rng_state, sample, x, y, offset, stride); } +#endif extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int sx, int sy, int sw, int sh, int offset, int stride) { diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 48df60162b1..d6863af8462 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -1194,6 +1194,7 @@ __device void kernel_path_trace_progressive(KernelGlobals *kg, path_rng_end(kg, rng_state, rng); } +#ifdef __NON_PROGRESSIVE__ __device void kernel_path_trace_non_progressive(KernelGlobals *kg, __global float *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride) @@ -1215,11 +1216,7 @@ __device void kernel_path_trace_non_progressive(KernelGlobals *kg, float4 L; if (ray.t != 0.0f) -#ifdef __NON_PROGRESSIVE__ L = kernel_path_non_progressive(kg, &rng, sample, ray, buffer); -#else - L = kernel_path_progressive(kg, &rng, sample, ray, buffer); -#endif else L = make_float4(0.0f, 0.0f, 0.0f, 0.0f); @@ -1228,6 +1225,7 @@ __device void kernel_path_trace_non_progressive(KernelGlobals *kg, path_rng_end(kg, rng_state, rng); } +#endif CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_sse2.cpp b/intern/cycles/kernel/kernel_sse2.cpp index 8321c679f07..344d530396b 100644 --- a/intern/cycles/kernel/kernel_sse2.cpp +++ b/intern/cycles/kernel/kernel_sse2.cpp @@ -39,10 +39,12 @@ CCL_NAMESPACE_BEGIN void kernel_cpu_sse2_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) { - if(kernel_data.integrator.progressive) - kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); - else +#ifdef __NON_PROGRESSIVE__ + if(!kernel_data.integrator.progressive) kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); + else +#endif + kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); } /* Tonemapping */ diff --git a/intern/cycles/kernel/kernel_sse3.cpp b/intern/cycles/kernel/kernel_sse3.cpp index 35a816f7da9..a25c0bca93a 100644 --- a/intern/cycles/kernel/kernel_sse3.cpp +++ b/intern/cycles/kernel/kernel_sse3.cpp @@ -41,10 +41,12 @@ CCL_NAMESPACE_BEGIN void kernel_cpu_sse3_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride) { - if(kernel_data.integrator.progressive) - kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); - else +#ifdef __NON_PROGRESSIVE__ + if(!kernel_data.integrator.progressive) kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); + else +#endif + kernel_path_trace_progressive(kg, buffer, rng_state, sample, x, y, offset, stride); } /* Tonemapping */ -- cgit v1.2.3