griefith/test
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Cycles: Attempt to work around compilation of sm_20 and sm_21

Browse Source 
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.
This commit is contained in:
Sergey Sharybin
2017-09-08 18:37:54 +02:00
parent 5d1a5001f4
commit ae41a08288
1 changed files with 5 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
intern/cycles/kernel/kernel_compat_cuda.h
View File

@@ -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