[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [59039] trunk/blender/intern/cycles: Cycles:

Thomas Dinges blender at dingto.org
Fri Aug 9 22:03:51 CEST 2013


Revision: 59039
          http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=59039
Author:   dingto
Date:     2013-08-09 20:03:49 +0000 (Fri, 09 Aug 2013)
Log Message:
-----------
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.

Modified Paths:
--------------
    trunk/blender/intern/cycles/device/device_cuda.cpp
    trunk/blender/intern/cycles/kernel/kernel.cpp
    trunk/blender/intern/cycles/kernel/kernel.cu
    trunk/blender/intern/cycles/kernel/kernel_path.h
    trunk/blender/intern/cycles/kernel/kernel_sse2.cpp
    trunk/blender/intern/cycles/kernel/kernel_sse3.cpp

Modified: trunk/blender/intern/cycles/device/device_cuda.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_cuda.cpp	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/device/device_cuda.cpp	2013-08-09 20:03:49 UTC (rev 59039)
@@ -572,8 +572,11 @@
 		/* 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;

Modified: trunk/blender/intern/cycles/kernel/kernel.cpp
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cpp	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/kernel/kernel.cpp	2013-08-09 20:03:49 UTC (rev 59039)
@@ -90,10 +90,12 @@
 
 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)
+#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);
-	else
-		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */

Modified: trunk/blender/intern/cycles/kernel/kernel.cu
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cu	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/kernel/kernel.cu	2013-08-09 20:03:49 UTC (rev 59039)
@@ -35,6 +35,7 @@
 		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 @@
 	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)
 {

Modified: trunk/blender/intern/cycles/kernel/kernel_path.h
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_path.h	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/kernel/kernel_path.h	2013-08-09 20:03:49 UTC (rev 59039)
@@ -1194,6 +1194,7 @@
 	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 @@
 	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 @@
 
 	path_rng_end(kg, rng_state, rng);
 }
+#endif
 
 CCL_NAMESPACE_END
 

Modified: trunk/blender/intern/cycles/kernel/kernel_sse2.cpp
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_sse2.cpp	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/kernel/kernel_sse2.cpp	2013-08-09 20:03:49 UTC (rev 59039)
@@ -39,10 +39,12 @@
 
 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)
+#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);
-	else
-		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */

Modified: trunk/blender/intern/cycles/kernel/kernel_sse3.cpp
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_sse3.cpp	2013-08-09 19:56:20 UTC (rev 59038)
+++ trunk/blender/intern/cycles/kernel/kernel_sse3.cpp	2013-08-09 20:03:49 UTC (rev 59039)
@@ -41,10 +41,12 @@
 
 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)
+#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);
-	else
-		kernel_path_trace_non_progressive(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */




More information about the Bf-blender-cvs mailing list