diff options
author | Thomas Dinges <blender@dingto.org> | 2016-05-07 00:11:41 +0300 |
---|---|---|
committer | Thomas Dinges <blender@dingto.org> | 2016-05-07 00:13:33 +0300 |
commit | 4422b3f9199cdd13c162ebc16c9e1d1b18f76bae (patch) | |
tree | e1126c08d85fc4ce51be98b132ce5832f2e227cd /intern | |
parent | 734d1aec3f93b8757533284330afc2ac651442bd (diff) |
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.
Diffstat (limited to 'intern')
-rw-r--r-- | intern/cycles/device/device.h | 8 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_bake.h | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/kernel_types.h | 11 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/cuda/kernel.cu | 2 | ||||
-rw-r--r-- | intern/cycles/render/shader.cpp | 4 |
5 files changed, 26 insertions, 3 deletions
diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h index 30d0003b940..4c1b7224837 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 d0ca256f323..60110fb980a 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 c9a895d9aec..02e69c7d75f 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 259b634f939..37fae54faf0 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 797f31ce92d..c07c6a04d03 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; } } |