diff options
Diffstat (limited to 'intern/cycles/kernel/device/cuda/compat.h')
| -rw-r--r-- | intern/cycles/kernel/device/cuda/compat.h | 5 |
1 files changed, 3 insertions, 2 deletions
diff --git a/intern/cycles/kernel/device/cuda/compat.h b/intern/cycles/kernel/device/cuda/compat.h index 51e1381d552..3a950779c11 100644 --- a/intern/cycles/kernel/device/cuda/compat.h +++ b/intern/cycles/kernel/device/cuda/compat.h @@ -30,6 +30,7 @@ typedef unsigned long long uint64_t; /* Qualifiers */ #define ccl_device __device__ __inline__ +#define ccl_device_extern extern "C" __device__ #if __CUDA_ARCH__ < 500 # define ccl_device_inline __device__ __forceinline__ # define ccl_device_forceinline __device__ __forceinline__ @@ -109,14 +110,14 @@ ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D typedef unsigned short half; -__device__ half __float2half(const float f) +ccl_device_forceinline half __float2half(const float f) { half val; asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f)); return val; } -__device__ float __half2float(const half h) +ccl_device_forceinline float __half2float(const half h) { float val; asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h)); |