cudax

#pragma once

#include <cuda.h>
#ifdef __CUDACC__ 
	#include <cuda_runtime.h>
	#include <math_functions.h>
#else
	#include <host_defines.h>
	#include <builtin_types.h>
#endif

#ifndef MATHX_API
	#define MATHX_API __host__ __device__
#endif
#include "mathx"

namespace cudax
{
	
// Half-precision floating-point bit representation vectors
typedef glm::vec<unsigned short, 2> hbvec2;
typedef glm::vec<unsigned short, 3> hbvec3;
typedef glm::vec<unsigned short, 4> hbvec4;

template <class T>
inline __host__ __device__ T* asRuntimePtr(CUdeviceptr ptr)
{
	return reinterpret_cast<T*>( static_cast<size_t>(ptr) );
}

template <class T>
inline __host__ __device__ CUdeviceptr asDriverPtr(T* ptr)
{
	return static_cast<CUdeviceptr>( reinterpret_cast<size_t>(ptr) );
}

template <class T, class X>
inline __host__ __device__ T as_(X const& x)
{
	static_assert(sizeof(T) == sizeof(X), "size mismatch");
	union U { X x; T t; char c[sizeof(T)]; } u;
	u.x = x;
	return u.t;
//	return reinterpret_cast<U const&>(x).t;
}

template <class T, size_t S>
inline __device__ __host__ char (&arraylen_helper(T const volatile (&a)[S]))[S];
#ifndef arraylen
	#define arraylen(x) sizeof(::cudax::arraylen_helper(x))
#endif

#ifdef __CUDACC__ 

template <class Type, class SurfaceType>
__device__ inline Type surf2Dread(SurfaceType surf, int2 coord, cudaSurfaceBoundaryMode boundaryMode = cudaBoundaryModeTrap)
{
	Type x;
	::surf2Dread(&x, surf, coord.x * (int) sizeof(Type), coord.y, boundaryMode);
	return x;
}
template <class Type, class SurfaceType>
__device__ inline void surf2Dwrite(Type const& x, SurfaceType surf, int2 coord, cudaSurfaceBoundaryMode boundaryMode = cudaBoundaryModeTrap)
{
	using glm::to_cuda;
	::surf2Dwrite(to_cuda(x), surf, coord.x * (int) sizeof(Type), coord.y, boundaryMode);
}


template <class Type, class SurfaceType>
__device__ inline Type surf3Dread(SurfaceType surf, int3 coord, cudaSurfaceBoundaryMode boundaryMode = cudaBoundaryModeTrap)
{
	Type x;
	::surf3Dread(&x, surf, coord.x * (int) sizeof(Type), coord.y, coord.z, boundaryMode);
	return x;
}
template <class Type, class SurfaceType>
__device__ inline void surf3Dwrite(Type const& x, SurfaceType surf, int3 coord, cudaSurfaceBoundaryMode boundaryMode = cudaBoundaryModeTrap)
{
	using glm::to_cuda;
	::surf3Dwrite(to_cuda(x), surf, coord.x * (int) sizeof(Type), coord.y, coord.z, boundaryMode);
}

template <class Type, class Vec>
__device__ inline Type tex2D(cudaTextureObject_t surf, Vec coord)
{
	return ::tex2D<Type>(surf, (float) coord.x, (float) coord.y);
}
template <class Type, class SurfaceType, class Vec>
__device__ inline Type tex2D(SurfaceType surf, Vec coord)
{
	return ::tex2D(surf, (float) coord.x, (float) coord.y);
}

template <class Type, class Vec>
__device__ inline Type tex3D(cudaTextureObject_t surf, Vec coord)
{
	return ::tex3D<Type>(surf, (float) coord.x, (float) coord.y, (float) coord.z);
}
template <class Type, class SurfaceType, class Vec>
__device__ inline Type tex3D(SurfaceType surf, Vec coord)
{
	return ::tex3D(surf, (float) coord.x, (float) coord.y, (float) coord.z);
}

__device__ inline hbvec3 toHalf(glm::vec3 const& v)
{
	return hbvec3( __float2half_rn(v.x), __float2half_rn(v.y), __float2half_rn(v.z) );
}

__device__ inline hbvec4 toHalf(glm::vec4 const& v)
{
	return hbvec4( __float2half_rn(v.x), __float2half_rn(v.y), __float2half_rn(v.z), __float2half_rn(v.w) );
}

__device__ inline glm::vec3 fromHalf(hbvec3 const& v)
{
	return glm::vec3( __half2float(v.x), __half2float(v.y), __half2float(v.z) );
}

__device__ inline glm::vec4 fromHalf(hbvec4 const& v)
{
	return glm::vec4( __half2float(v.x), __half2float(v.y), __half2float(v.z), __half2float(v.w) );
}

__device__ inline float fast_sqrt(float x)
{
	float r;
	asm("sqrt.approx.f32 %0, %1;" : "=f"(r) : "f"(x));
	return r;
}

__device__ inline float fast_rsqrt(float x)
{
	float r;
	asm("rsqrt.approx.f32 %0, %1;" : "=f"(r) : "f"(x));
	return r;
}

__device__ inline float fast_rcp(float x)
{
	float r;
	asm("rcp.approx.f32 %0, %1;" : "=f"(r) : "f"(x));
	return r;
}

#endif

} // namespace