diff options
author | Campbell Barton <ideasman42@gmail.com> | 2017-10-09 12:49:27 +0300 |
---|---|---|
committer | Campbell Barton <ideasman42@gmail.com> | 2017-10-09 12:49:27 +0300 |
commit | a5b4b0f21c1ae8c96e4fea9abdcfac2fab1cf300 (patch) | |
tree | 0658d8bdfb8ec03652aa04f82ee8a4d243ec6370 /intern/cycles/kernel/kernel_compat_cuda.h | |
parent | d68f698cf0321477c0734474150eb4bc43c4e85f (diff) | |
parent | abcda06934aba054de8540b66b13c2bbc5f8f515 (diff) |
Merge branch '28' into custom-manipulatorscustom-manipulators
Diffstat (limited to 'intern/cycles/kernel/kernel_compat_cuda.h')
-rw-r--r-- | intern/cycles/kernel/kernel_compat_cuda.h | 13 |
1 files changed, 11 insertions, 2 deletions
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index 38708f7ff0b..2e8ca48c413 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -38,21 +38,30 @@ /* 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 -#define ccl_constant +#define ccl_static_constant __constant__ +#define ccl_constant const #define ccl_local __shared__ #define ccl_local_param #define ccl_private #define ccl_may_alias #define ccl_addr_space #define ccl_restrict __restrict__ +/* TODO(sergey): In theory we might use references with CUDA, however + * performance impact yet to be investigated. + */ +#define ccl_ref #define ccl_align(n) __align__(n) #define ATTR_FALLTHROUGH |