From ae41a082889d45c696ab9426cfd921f61540c2ea Mon Sep 17 00:00:00 2001 From: Sergey Sharybin Date: Fri, 8 Sep 2017 18:37:54 +0200 Subject: [PATCH] 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. --- intern/cycles/kernel/kernel_compat_cuda.h | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) 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