[Bf-blender-cvs] [4422b3f] master: Some fixes for CUDA runtime compile: * When Baking wasn't used we got an error. * On top of Volume Nodes (NODES_FEATURE_VOLUME), we now also check if we need volume sampling code, so we can disable that as well and save some further compilation time.

[Thomas Dinges]

Commit: 4422b3f9199cdd13c162ebc16c9e1d1b18f76bae
Author: Thomas Dinges
Date:   Fri May 6 23:11:41 2016 +0200
Branches: master
https://developer.blender.org/rB4422b3f9199cdd13c162ebc16c9e1d1b18f76bae

Some fixes for CUDA runtime compile:
* When Baking wasn't used we got an error.
* On top of Volume Nodes (NODES_FEATURE_VOLUME), we now also check if we need volume sampling code,
so we can disable that as well and save some further compilation time.

===================================================================

M	intern/cycles/device/device.h
M	intern/cycles/kernel/kernel_bake.h
M	intern/cycles/kernel/kernel_types.h
M	intern/cycles/kernel/kernels/cuda/kernel.cu
M	intern/cycles/render/shader.cpp

===================================================================

diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h
index 30d0003..4c1b722 100644
--- a/intern/cycles/device/device.h
+++ b/intern/cycles/device/device.h
@@ -103,6 +103,9 @@ public:
 	/* Use subsurface scattering materials. */
 	bool use_subsurface;
 
+	/* Use volume materials. */
+	bool use_volume;
+
 	/* Use branched integrator. */
 	bool use_integrator_branched;
 
@@ -118,6 +121,7 @@ public:
 		use_camera_motion = false;
 		use_baking = false;
 		use_subsurface = false;
+		use_volume = false;
 		use_integrator_branched = false;
 	}
 
@@ -132,6 +136,7 @@ public:
 		         use_camera_motion == requested_features.use_camera_motion &&
 		         use_baking == requested_features.use_baking &&
 		         use_subsurface == requested_features.use_subsurface &&
+		         use_volume == requested_features.use_volume &&
 		         use_integrator_branched == requested_features.use_integrator_branched);
 	}
 
@@ -161,6 +166,9 @@ public:
 		if(!use_baking) {
 			build_options += " -D__NO_BAKING__";
 		}
+		if(!use_volume) {
+			build_options += " -D__NO_VOLUME__";
+		}
 		if(!use_subsurface) {
 			build_options += " -D__NO_SUBSURFACE__";
 		}
diff --git a/intern/cycles/kernel/kernel_bake.h b/intern/cycles/kernel/kernel_bake.h
index d0ca256..60110fb 100644
--- a/intern/cycles/kernel/kernel_bake.h
+++ b/intern/cycles/kernel/kernel_bake.h
@@ -16,7 +16,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#ifndef __NO_BAKING__
+#ifdef __BAKING__
 
 ccl_device void compute_light_pass(KernelGlobals *kg, ShaderData *sd, PathRadiance *L, RNG rng,
                                    int pass_filter, int sample)
@@ -483,7 +483,7 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
 		output[i] += make_float4(out.x, out.y, out.z, 1.0f) * output_fac;
 }
 
-#endif  /* __NO_BAKING__ */
+#endif  /* __BAKING__ */
 
 ccl_device void kernel_shader_evaluate(KernelGlobals *kg,
                                        ccl_global uint4 *input,
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index c9a895d..02e69c7 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -120,6 +120,7 @@ CCL_NAMESPACE_BEGIN
 #    define __CAMERA_MOTION__
 #    define __OBJECT_MOTION__
 #    define __HAIR__
+#    define __BAKING__
 #    ifdef __KERNEL_EXPERIMENTAL__
 #      define __TRANSPARENT_SHADOWS__
 #    endif
@@ -167,13 +168,14 @@ CCL_NAMESPACE_BEGIN
 #  define __CAMERA_MOTION__
 #  define __OBJECT_MOTION__
 #  define __HAIR__
+#  define __BAKING__
 #endif
 
 #ifdef WITH_CYCLES_DEBUG
 #  define __KERNEL_DEBUG__
 #endif
 
-/* Scene-based selective featrues compilation. */
+/* Scene-based selective features compilation. */
 #ifdef __NO_CAMERA_MOTION__
 #  undef __CAMERA_MOTION__
 #endif
@@ -183,9 +185,16 @@ CCL_NAMESPACE_BEGIN
 #ifdef __NO_HAIR__
 #  undef __HAIR__
 #endif
+#ifdef __NO_VOLUME__
+#  undef __VOLUME__
+#  undef __VOLUME_SCATTER__
+#endif
 #ifdef __NO_SUBSURFACE__
 #  undef __SUBSURFACE__
 #endif
+#ifdef __NO_BAKING__
+#  undef __BAKING__
+#endif
 #ifdef __NO_BRANCHED_PATH__
 #  undef __BRANCHED_PATH__
 #endif
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index 259b634..37fae54 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -193,6 +193,7 @@ kernel_cuda_shader(uint4 *input,
 	}
 }
 
+#ifdef __BAKING__
 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 filter, int sx, int sw, int offset, int sample)
@@ -202,6 +203,7 @@ kernel_cuda_bake(uint4 *input, float4 *output, int type, int filter, int sx, int
 	if(x < sx + sw)
 		kernel_bake_evaluate(NULL, input, output, (ShaderEvalType)type, filter, x, offset, sample);
 }
+#endif
 
 #endif
 
diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp
index 797f31c..c07c6a0 100644
--- a/intern/cycles/render/shader.cpp
+++ b/intern/cycles/render/shader.cpp
@@ -528,6 +528,10 @@ void ShaderManager::get_requested_features(Scene *scene,
 		if(output_node->input("Displacement")->link != NULL) {
 			requested_features->nodes_features |= NODE_FEATURE_BUMP;
 		}
+		/* On top of volume nodes, also check if we need volume sampling because
+		 * e.g. an Emission node would slip through the NODE_FEATURE_VOLUME check */
+		if(shader->has_volume)
+			requested_features->use_volume |= true;
 	}
 }