[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