Differential D15902 Diff 57407 intern/cycles/kernel/device/cuda/compat.h

Changeset View

Standalone View

intern/cycles/kernel/device/cuda/compat.h

	Show All 24 Lines
	# define FLT_MIN 1.175494350822287507969e-38f			# define FLT_MIN 1.175494350822287507969e-38f
	# define FLT_MAX 340282346638528859811704183484516925440.0f			# define FLT_MAX 340282346638528859811704183484516925440.0f
	# define FLT_EPSILON 1.192092896e-07F			# define FLT_EPSILON 1.192092896e-07F
	#endif			#endif

	/* Qualifiers */			/* Qualifiers */

	#define ccl_device __device__ __inline__			#define ccl_device __device__ __inline__
				#define ccl_device_extern extern "C" __device__
	#if __CUDA_ARCH__ < 500			#if __CUDA_ARCH__ < 500
	# define ccl_device_inline __device__ __forceinline__			# define ccl_device_inline __device__ __forceinline__
	# define ccl_device_forceinline __device__ __forceinline__			# define ccl_device_forceinline __device__ __forceinline__
	#else			#else
	# define ccl_device_inline __device__ __inline__			# define ccl_device_inline __device__ __inline__
	# define ccl_device_forceinline __device__ __forceinline__			# define ccl_device_forceinline __device__ __forceinline__
	#endif			#endif
	#define ccl_device_noinline __device__ __noinline__			#define ccl_device_noinline __device__ __noinline__
	▲ Show 20 Lines • Show All 63 Lines • ▼ Show 20 Lines
	#define tanf(x) __tanf(((float)(x)))			#define tanf(x) __tanf(((float)(x)))
	#define logf(x) __logf(((float)(x)))			#define logf(x) __logf(((float)(x)))
	#define expf(x) __expf(((float)(x)))			#define expf(x) __expf(((float)(x)))

	/* Half */			/* Half */

	typedef unsigned short half;			typedef unsigned short half;

	__device__ half __float2half(const float f)			ccl_device_forceinline half __float2half(const float f)
	{			{
	half val;			half val;
	asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f));			asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f));
	return val;			return val;
	}			}

	__device__ float __half2float(const half h)			ccl_device_forceinline float __half2float(const half h)
	{			{
	float val;			float val;
	asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h));			asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h));
	return val;			return val;
	}			}

	/* Types */			/* Types */

	#include "util/half.h"			#include "util/half.h"
	#include "util/types.h"			#include "util/types.h"