[Bf-blender-cvs] [ae41a082889] master: Cycles: Attempt to work around compilation of sm_20 and sm_21

[Sergey Sharybin]

Commit: ae41a082889d45c696ab9426cfd921f61540c2ea
Author: Sergey Sharybin
Date:   Fri Sep 8 18:37:54 2017 +0200
Branches: master
https://developer.blender.org/rBae41a082889d45c696ab9426cfd921f61540c2ea

Cycles: Attempt to work around compilation of sm_20 and sm_21

Disabled forceinline for those architectures, which seems to be compiling
successfully more often.

There might be ~3% slowdown based on quick tests, but better be rendering
something rather than failing to compile kernels again and again.

Those architectures will be doomed for abandon once we'll switch to toolkit 9.

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

M	intern/cycles/kernel/kernel_compat_cuda.h

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

diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index 6d1cf055f2c..1e2af9de8b3 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -38,11 +38,15 @@
 /* Qualifier wrappers for different names on different devices */
 
 #define ccl_device  __device__ __inline__
+#if __CUDA_ARCH__ < 300
+#  define ccl_device_inline  __device__ __inline__
 #  define ccl_device_forceinline  __device__ __forceinline__
-#if __CUDA_ARCH__ < 500
+#elif __CUDA_ARCH__ < 500
 #  define ccl_device_inline  __device__ __forceinline__
+#  define ccl_device_forceinline  __device__ __forceinline__
 #else
 #  define ccl_device_inline  __device__ __inline__
+#  define ccl_device_forceinline  __device__ __forceinline__
 #endif
 #define ccl_device_noinline  __device__ __noinline__
 #define ccl_global