28 #define __device__ __attribute__((device))  macro
94 #define GET(name, reg) static inline __device__ uint3 name() {\
123 #define TEX2D(type, ret) static inline __device__ void conv(type* out, unsigned a, unsigned b, unsi…
137 inline __device__ T tex2D(cudaTextureObject_t texObject, float x, float y)
149 inline __device__ float4 tex2D<float4>(cudaTextureObject_t texObject, float x, float y)
159 inline __device__ float tex2D<float>(cudaTextureObject_t texObject, float x, float y)
165 inline __device__ float2 tex2D<float2>(cudaTextureObject_t texObject, float x, float y)
172 static inline __device__ float floorf(float a) { return __builtin_floorf(a); }  in floorf()
173 static inline __device__ float floor(float a) { return __builtin_floorf(a); }  in floor()
174 static inline __device__ double floor(double a) { return __builtin_floor(a); }  in floor()
175 static inline __device__ float ceilf(float a) { return __builtin_ceilf(a); }  in ceilf()
176 static inline __device__ float ceil(float a) { return __builtin_ceilf(a); }  in ceil()
177 static inline __device__ double ceil(double a) { return __builtin_ceil(a); }  in ceil()
178 static inline __device__ float truncf(float a) { return __builtin_truncf(a); }  in truncf()
179 static inline __device__ float trunc(float a) { return __builtin_truncf(a); }  in trunc()
180 static inline __device__ double trunc(double a) { return __builtin_trunc(a); }  in trunc()
181 static inline __device__ float fabsf(float a) { return __builtin_fabsf(a); }  in fabsf()
182 static inline __device__ float fabs(float a) { return __builtin_fabsf(a); }  in fabs()
183 static inline __device__ double fabs(double a) { return __builtin_fabs(a); }  in fabs()
185 static inline __device__ float __sinf(float a) { return __nvvm_sin_approx_f(a); }  in __sinf()
186 static inline __device__ float __cosf(float a) { return __nvvm_cos_approx_f(a); }  in __cosf()