diff options
| author | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-09-08 19:37:54 +0300 |
|---|---|---|
| committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-09-08 19:37:54 +0300 |
| commit | ae41a082889d45c696ab9426cfd921f61540c2ea (patch) | |
| tree | 489042f04d331c501cc1d72f677920e64305418b /intern | |
| parent | 5d1a5001f4d039b421757419472287331ab03bd3 (diff) | |
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.
Diffstat (limited to 'intern')
| -rw-r--r-- | intern/cycles/kernel/kernel_compat_cuda.h | 6 |
1 files changed, 5 insertions, 1 deletions
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 |