diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2016-08-03 12:41:58 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2016-08-03 12:41:58 +0300 |
commit | 960db4c961da6c7faaa0b5fcdbba4d38c90ef298 (patch) | |
tree | b10df6c5fc858d892ec3a3d35648e61a86f0b0aa /intern/cycles/kernel/kernel_compat_cuda.h | |
parent | 70100b4ec7ba299912bf999f15914c0a29125ffd (diff) |
Cycles: Revert recent inline changes for CUDA 8 and sm_50+
This changes actually lead to 2x slowdown. It's getting a bit annoying
because those are the changes to make pre-maxwell cards render with the
same speed.
Diffstat (limited to 'intern/cycles/kernel/kernel_compat_cuda.h')
-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 fb5812ebcb8..a039b414006 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -36,7 +36,11 @@ /* Qualifier wrappers for different names on different devices */ #define ccl_device __device__ __inline__ -#define ccl_device_inline __device__ __forceinline__ +#if (__KERNEL_CUDA_VERSION__ == 80) && (__CUDA_ARCH__ < 500) +# define ccl_device_inline __device__ __forceinline__ +#else +# define ccl_device_inline __device__ __inline__ +#endif #define ccl_device_noinline __device__ __noinline__ #define ccl_global #define ccl_constant |