Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern
diff options
context:
space:
mode:
authorSergey Sharybin <sergey.vfx@gmail.com>2016-08-03 12:41:58 +0300
committerSergey Sharybin <sergey.vfx@gmail.com>2016-08-03 12:41:58 +0300
commit960db4c961da6c7faaa0b5fcdbba4d38c90ef298 (patch)
treeb10df6c5fc858d892ec3a3d35648e61a86f0b0aa /intern
parent70100b4ec7ba299912bf999f15914c0a29125ffd (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')
-rw-r--r--intern/cycles/kernel/kernel_compat_cuda.h6
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