diff options
author | Martijn Versteegh <blender@aaltjegron.nl> | 2022-11-11 21:44:16 +0300 |
---|---|---|
committer | Martijn Versteegh <blender@aaltjegron.nl> | 2022-11-11 21:44:16 +0300 |
commit | b4757873bf1c2fb5e3f858a085af62359d3c1b3f (patch) | |
tree | 98d4407723ef2816334d06cd79388c0d7b2d708c /intern/cycles/kernel/device/cuda/compat.h | |
parent | 62effd6791858221392ff7b92104a22ed44ebf7d (diff) | |
parent | 2c596319a4888aa40bfdf41f9ea5d446179141d0 (diff) |
Merge branch 'master' into refactor-mesh-uv-map-generic
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)); |