[Bf-blender-cvs] [c9e8888] master: Cycles: Disable bake OpenCL kernel for NVidia devices prior to sm_30

Sergey Sharybin noreply at git.blender.org
Fri May 15 10:23:55 CEST 2015


Commit: c9e8888f877c7c4332d4d1a1dbb372a3e0b7f28a
Author: Sergey Sharybin
Date:   Fri May 15 12:30:04 2015 +0500
Branches: master
https://developer.blender.org/rBc9e8888f877c7c4332d4d1a1dbb372a3e0b7f28a

Cycles: Disable bake OpenCL kernel for NVidia devices prior to sm_30

Driver fails to compile kernel in reasonable time for those devices here,
so for easier testing of the OpenCL split kernel work disabling bake kernel
for now.

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

M	intern/cycles/kernel/kernel.cl

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

diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index 19e3949..cbc0592 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -102,8 +102,19 @@ __kernel void kernel_ocl_bake(
 
 	int x = sx + get_global_id(0);
 
-	if(x < sx + sw)
+	if(x < sx + sw) {
+#if defined(__KERNEL_OPENCL_NVIDIA__) && __COMPUTE_CAPABILITY__ < 300
+		/* NVidia compiler is spending infinite amount of time trying
+		 * to deal with kernel_bake_evaluate() on architectures prior
+		 * to sm_30.
+		 * For now we disable baking kernel for those devices, so at
+		 * least rendering with split kernel could be compiled.
+		 */
+		output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+#else
 		kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, x, offset, sample);
+#endif
+	}
 }
 
 __kernel void kernel_ocl_convert_to_byte(




More information about the Bf-blender-cvs mailing list