stdin

#define __global__ __attribute__((global))

typedef __attribute__((device_builtin)) unsigned long long cudaTextureObject_t;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(2))) uchar2
{
    unsigned char x, y;
} uchar2;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(4))) ushort2
{
    unsigned short x, y;
} ushort2;

typedef struct __attribute__((device_builtin)) uint3
{
    unsigned int x, y, z;
} uint3;

typedef struct uint3 dim3;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(8))) int2
{
    int x, y;
} int2;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(4))) uchar4
{
    unsigned char x, y, z, w;
} uchar4;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(8))) ushort4
{
    unsigned char x, y, z, w;
} ushort4;

typedef struct __attribute__((device_builtin)) __attribute__((aligned(16))) int4
{
    int x, y, z, w;
} int4;

uint3 __attribute__((device_builtin)) extern const __attribute__((device)) blockIdx;
dim3 __attribute__((device_builtin)) extern const __attribute__((device)) blockDim;
uint3 __attribute__((device_builtin)) extern const __attribute__((device)) threadIdx;

#define max(a, b) ((a) > (b) ? (a) : (b))
#define min(a, b) ((a) < (b) ? (a) : (b))

#define TEX2D(type, name) __attribute__((device)) __attribute__((cudart_builtin)) void tex2D(type *, cudaTextureObject_t, float, float) asm("__itex2D_" #name);

TEX2D(unsigned char, uchar)
TEX2D(unsigned short, ushort)
TEX2D(uchar2, uchar2)
TEX2D(ushort2, ushort2)
TEX2D(uchar4, uchar4)
TEX2D(ushort4, ushort4)

#define make_uchar2(a, b) ((uchar2){.x = a, .y = b})
#define make_ushort2(a, b) ((ushort2){.x = a, .y = b})
#define make_uchar4(a, b, c, d) ((uchar4){.x = a, .y = b, .z = c, .w = d})
#define make_ushort4(a, b, c, d) ((ushort4){.x = a, .y = b, .z = c, .w = d})

template <class T>
static __attribute__((device)) T tex2D(cudaTextureObject_t texObject, float x, float y)
{
  T ret;
  tex2D(&ret, texObject, x, y);
  return ret;
}