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.

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 */