From 93fa17490091ddcd04d9dc8e6fa5f2b0fe38eddb Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 16 May 2017 21:35:40 -0500 Subject: [PATCH] changed vector types to make sure it generate proper llvm vector types Change-Id: I6c4616dae137dc4eac35e5827dc5b7f3251e0247 --- hipamd/include/hip/hcc_detail/hip_fp16.h | 125 +- .../include/hip/hcc_detail/hip_vector_types.h | 4067 +---------------- hipamd/src/hip_fp16.cpp | 442 +- hipamd/src/hip_hc_gfx803.ll | 147 +- 4 files changed, 266 insertions(+), 4515 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h index 0a861b64af..f1f52e4122 100644 --- a/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -25,17 +25,6 @@ THE SOFTWARE. #include "hip/hcc_detail/hip_vector_types.h" -#if __clang_major__ > 3 - -typedef __fp16 __half; - -typedef struct __attribute__((aligned(4))){ - union { - __half p[2]; - unsigned int q; - }; -} __half2; - typedef __half half; typedef __half2 half2; @@ -214,10 +203,10 @@ __device__ __half __ushort2half_ru(unsigned short int i); __device__ __half __ushort2half_rz(unsigned short int i); __device__ __half __ushort_as_half(const unsigned short int i); -extern "C" int __hip_hc_ir_hadd2_int(int, int); -extern "C" int __hip_hc_ir_hfma2_int(int, int, int); -extern "C" int __hip_hc_ir_hmul2_int(int, int); -extern "C" int __hip_hc_ir_hsub2_int(int, int); +extern "C" __half2 __hip_hc_ir_hadd2_int(__half2, __half2); +extern "C" __half2 __hip_hc_ir_hfma2_int(__half2, __half2, __half2); +extern "C" __half2 __hip_hc_ir_hmul2_int(__half2, __half2); +extern "C" __half2 __hip_hc_ir_hsub2_int(__half2, __half2); extern "C" __half __hip_hc_ir_hceil_half(__half) __asm("llvm.ceil.f16"); extern "C" __half __hip_hc_ir_hcos_half(__half) __asm("llvm.cos.f16"); @@ -231,16 +220,16 @@ extern "C" __half __hip_hc_ir_hsin_half(__half) __asm("llvm.sin.f16"); extern "C" __half __hip_hc_ir_hsqrt_half(__half) __asm("llvm.sqrt.f16"); extern "C" __half __hip_hc_ir_htrunc_half(__half) __asm("llvm.trunc.f16"); -extern "C" int __hip_hc_ir_h2ceil_int(int); -extern "C" int __hip_hc_ir_h2cos_int(int); -extern "C" int __hip_hc_ir_h2exp2_int(int); -extern "C" int __hip_hc_ir_h2floor_int(int); -extern "C" int __hip_hc_ir_h2log2_int(int); -extern "C" int __hip_hc_ir_h2rcp_int(int); -extern "C" int __hip_hc_ir_h2rsqrt_int(int); -extern "C" int __hip_hc_ir_h2sin_int(int); -extern "C" int __hip_hc_ir_h2sqrt_int(int); -extern "C" int __hip_hc_ir_h2trunc_int(int); +extern "C" __half2 __hip_hc_ir_h2ceil_int(__half2); +extern "C" __half2 __hip_hc_ir_h2cos_int(__half2); +extern "C" __half2 __hip_hc_ir_h2exp2_int(__half2); +extern "C" __half2 __hip_hc_ir_h2floor_int(__half2); +extern "C" __half2 __hip_hc_ir_h2log2_int(__half2); +extern "C" __half2 __hip_hc_ir_h2rcp_int(__half2); +extern "C" __half2 __hip_hc_ir_h2rsqrt_int(__half2); +extern "C" __half2 __hip_hc_ir_h2sin_int(__half2); +extern "C" __half2 __hip_hc_ir_h2sqrt_int(__half2); +extern "C" __half2 __hip_hc_ir_h2trunc_int(__half2); /* Half2 Arithmetic Functions @@ -248,63 +237,63 @@ extern "C" int __hip_hc_ir_h2trunc_int(int); __device__ static inline __half2 __hadd2(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hadd2_int(a.q, b.q); + c.xy = __hip_hc_ir_hadd2_int(a.xy, b.xy); return c; } __device__ static inline __half2 __hadd2_sat(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hadd2_int(a.q, b.q); + c.xy = __hip_hc_ir_hadd2_int(a.xy, b.xy); return c; } __device__ static inline __half2 __hfma2(__half2 a, __half2 b, __half2 c) { __half2 d; - d.q = __hip_hc_ir_hfma2_int(a.q, b.q, c.q); + d.xy = __hip_hc_ir_hfma2_int(a.xy, b.xy, c.xy); return d; } __device__ static inline __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c) { __half2 d; - d.q = __hip_hc_ir_hfma2_int(a.q, b.q, c.q); + d.xy = __hip_hc_ir_hfma2_int(a.xy, b.xy, c.xy); return d; } __device__ static inline __half2 __hmul2(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hmul2_int(a.q, b.q); + c.xy = __hip_hc_ir_hmul2_int(a.xy, b.xy); return c; } __device__ static inline __half2 __hmul2_sat(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hmul2_int(a.q, b.q); + c.xy = __hip_hc_ir_hmul2_int(a.xy, b.xy); return c; } __device__ static inline __half2 __hsub2(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hsub2_int(a.q, b.q); + c.xy = __hip_hc_ir_hsub2_int(a.xy, b.xy); return c; } __device__ static inline __half2 __hneg2(__half2 a) { __half2 c; - c.p[0] = - a.p[0]; - c.p[1] = - a.p[1]; + c.x = - a.x; + c.y = - a.y; return c; } __device__ static inline __half2 __hsub2_sat(__half2 a, __half2 b) { __half2 c; - c.q = __hip_hc_ir_hsub2_int(a.q, b.q); + c.xy = __hip_hc_ir_hsub2_int(a.xy, b.xy); return c; } __device__ static inline __half2 h2div(__half2 a, __half2 b) { __half2 c; - c.p[0] = a.p[0] / b.p[0]; - c.p[1] = a.p[1] / b.p[1]; + c.x = a.x / b.x; + c.y = a.y / b.y; return c; } @@ -375,112 +364,94 @@ Half2 Math Operations __device__ static inline __half2 h2ceil(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2ceil_int(h.q); + a.xy = __hip_hc_ir_h2ceil_int(h.xy); return a; } __device__ static inline __half2 h2cos(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2cos_int(h.q); + a.xy = __hip_hc_ir_h2cos_int(h.xy); return a; } __device__ static inline __half2 h2exp(const __half2 h) { __half2 factor; - factor.p[0] = 1.442694; - factor.p[1] = 1.442694; - factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q)); + factor.x = 1.442694; + factor.y = 1.442694; + factor.xy = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.xy, factor.xy)); return factor; } __device__ static inline __half2 h2exp10(const __half2 h) { __half2 factor; - factor.p[0] = 3.3219281; - factor.p[1] = 3.3219281; - factor.q = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.q, factor.q)); + factor.x = 3.3219281; + factor.y = 3.3219281; + factor.xy = __hip_hc_ir_h2exp2_int(__hip_hc_ir_hmul2_int(h.xy, factor.xy)); return factor; } __device__ static inline __half2 h2exp2(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2exp2_int(h.q); + a.xy = __hip_hc_ir_h2exp2_int(h.xy); return a; } __device__ static inline __half2 h2floor(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2floor_int(h.q); + a.xy = __hip_hc_ir_h2floor_int(h.xy); return a; } __device__ static inline __half2 h2log(const __half2 h) { __half2 factor; - factor.p[0] = 0.693147; - factor.p[1] = 0.693147; - factor. q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q); + factor.x = 0.693147; + factor.y = 0.693147; + factor.xy = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.xy), factor.xy); return factor; } __device__ static inline __half2 h2log10(const __half2 h) { __half2 factor; - factor.p[0] = 0.301029; - factor.p[1] = 0.301029; - factor.q = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.q), factor.q); + factor.x = 0.301029; + factor.y = 0.301029; + factor.xy = __hip_hc_ir_hmul2_int(__hip_hc_ir_h2log2_int(h.xy), factor.xy); return factor; } __device__ static inline __half2 h2log2(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2log2_int(h.q); + a.xy = __hip_hc_ir_h2log2_int(h.xy); return a; } __device__ static inline __half2 h2rcp(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2rcp_int(h.q); + a.xy = __hip_hc_ir_h2rcp_int(h.xy); return a; } __device__ static inline __half2 h2rsqrt(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2rsqrt_int(h.q); + a.xy = __hip_hc_ir_h2rsqrt_int(h.xy); return a; } __device__ static inline __half2 h2sin(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2sin_int(h.q); + a.xy = __hip_hc_ir_h2sin_int(h.xy); return a; } __device__ static inline __half2 h2sqrt(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2sqrt_int(h.q); + a.xy = __hip_hc_ir_h2sqrt_int(h.xy); return a; } __device__ static inline __half2 h2trunc(const __half2 h) { __half2 a; - a.q = __hip_hc_ir_h2trunc_int(h.q); + a.xy = __hip_hc_ir_h2trunc_int(h.xy); return a; } -#endif - -#if __clang_major__ == 3 - -typedef struct { - unsigned x: 16; -} __half; - -typedef struct __attribute__((aligned(4))){ - union { - __half p[2]; - unsigned int q; - }; -} __half2; - - -#endif - #endif diff --git a/hipamd/include/hip/hcc_detail/hip_vector_types.h b/hipamd/include/hip/hcc_detail/hip_vector_types.h index 35c6c23548..251da504ab 100644 --- a/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -34,1120 +34,93 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -#define MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(type) \ -__device__ __host__ type() {} \ -__device__ __host__ type(type& val) : x(val.x) { } \ -__device__ __host__ type(const type& val) : x(val.x) { } - -#define MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(type) \ -__device__ __host__ type() {} \ -__device__ __host__ type(type& val) : x(val.x), y(val.y) { } \ -__device__ __host__ type(const type& val) : x(val.x), y(val.y) { } - -#define MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(type) \ -__device__ __host__ type() {} \ -__device__ __host__ type(type& val) : x(val.x), y(val.y), z(val.z) { } \ -__device__ __host__ type(const type& val) : x(val.x), y(val.y), z(val.z) { } - -#define MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(type) \ -__device__ __host__ type() {} \ -__device__ __host__ type(type& val) : x(val.x), y(val.y), z(val.z), w(val.w) { } \ -__device__ __host__ type(const type& val) : x(val.x), y(val.y), z(val.z), w(val.w) { } - - -#define MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(type, type1) \ -__device__ __host__ type(type1 val) : x(val) {} \ - -#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(type, type1) \ -__device__ __host__ type(type1 val) : x(val), y(val) {} \ -__device__ __host__ type(type1 val1, type1 val2) : x(val1), y(val2) {} - -#define MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(type, type1) \ -__device__ __host__ type(type1 val) : x(val), y(val), z(val) {} \ -__device__ __host__ type(type1 val1, type1 val2, type1 val3) : x(val1), y(val2), z(val3) {} - -#define MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(type, type1) \ -__device__ __host__ type(type1 val) : x(val), y(val), z(val), w(val) {} \ -__device__ __host__ type(type1 val1, type1 val2, type1 val3, type1 val4) : x(val1), y(val2), z(val3), w(val4) {} - -struct uchar1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(uchar1) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uchar1, signed long long) - - #endif - unsigned char x; - -} __attribute__((aligned(1))); - -struct uchar2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(uchar2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uchar2, signed long long) - #endif - union { - struct { - unsigned char x, y; - }; - unsigned short a; - }; -} __attribute__((aligned(2))); - -struct uchar3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(uchar3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uchar3, signed long long) - #endif - unsigned char x, y, z; -}; - -struct uchar4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(uchar4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uchar4, signed long long) - #endif - union { - struct { - unsigned char x, y, z, w; - }; - unsigned int a; - }; -} __attribute__((aligned(4))); - - -struct char1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(char1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(char1, signed long long) - #endif - signed char x; -} __attribute__((aligned(1))); - -struct char2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(char2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(char2, signed long long) - #endif - union { - struct { - signed char x, y; - }; - unsigned short a; - }; -} __attribute__((aligned(2))); - -struct char3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(char3) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(char3, signed long long) - #endif - signed char x, y, z; -}; - -struct char4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(char4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(char4, signed long long) - #endif - union { - struct { - signed char x, y, z, w; - }; - unsigned int a; - }; -} __attribute__((aligned(4))); - - - -struct ushort1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(ushort1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ushort1, signed long long) - #endif - unsigned short x; -} __attribute__((aligned(2))); - -struct ushort2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(ushort2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ushort2, signed long long) - #endif - union { - struct { - unsigned short x, y; - }; - unsigned int a; - }; -} __attribute__((aligned(4))); - -struct ushort3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(ushort3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ushort3, signed long long) - #endif - unsigned short x, y, z; -}; - -struct ushort4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(ushort4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ushort4, signed long long) - #endif - union { - struct { - unsigned short x, y, z, w; - }; - unsigned int a, b; - }; -} __attribute__((aligned(8))); - -struct short1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(short1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(short1, signed long long) - #endif - signed short x; -} __attribute__((aligned(2))); - -struct short2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(short2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(short2, signed long long) - #endif - union { - struct { - signed short x, y; - }; - unsigned int a; - }; - -} __attribute__((aligned(4))); - -struct short3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(short3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(short3, signed long long) - #endif - signed short x, y, z; -}; - -struct short4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(short4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(short4, signed long long) - #endif - union { - struct { - signed short x, y, z, w; - }; - unsigned int a, b; - }; -} __attribute__((aligned(8))); - - -struct uint1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(uint1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(uint1, signed long long) - #endif - unsigned int x; -} __attribute__((aligned(4))); - -struct uint2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(uint2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(uint2, signed long long) - #endif - unsigned int x, y; -} __attribute__((aligned(8))); - -struct uint3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(uint3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(uint3, signed long long) - #endif - unsigned int x, y, z; -}; - -struct uint4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(uint4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(uint4, signed long long) - #endif - unsigned int x, y, z, w; -} __attribute__((aligned(16))); - -struct int1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(int1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(int1, signed long long) - #endif - signed int x; -} __attribute__((aligned(4))); - -struct int2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(int2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(int2, signed long long) - #endif - signed int x, y; -} __attribute__((aligned(8))); - -struct int3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(int3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(int3, signed long long) - #endif - signed int x, y, z; -}; - -struct int4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(int4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(int4, signed long long) - #endif - signed int x, y, z, w; -} __attribute__((aligned(16))); - - -struct float1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(float1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(float1, signed long long) - #endif - float x; -} __attribute__((aligned(4))); - -struct float2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(float2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(float2, signed long long) - #endif - float x, y; -} __attribute__((aligned(8))); - -struct float3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(float3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(float3, signed long long) - #endif - float x, y, z; -}; - -struct float4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(float4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(float4, signed long long) - #endif - float x, y, z, w; -} __attribute__((aligned(16))); - - - -struct double1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(double1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(double1, signed long long) - #endif - double x; -} __attribute__((aligned(8))); - -struct double2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(double2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(double2, signed long long) - #endif - double x, y; -} __attribute__((aligned(16))); - -struct double3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(double3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(double3, signed long long) - #endif - double x, y, z; -}; - -struct double4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(double4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(double4, signed long long) - #endif - double x, y, z, w; -} __attribute__((aligned(32))); - - -struct ulong1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(ulong1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulong1, signed long long) - #endif - unsigned long x; -} __attribute__((aligned(8))); - -struct ulong2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(ulong2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulong2, signed long long) - #endif - unsigned long x, y; -} __attribute__((aligned(16))); - -struct ulong3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(ulong3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulong3, signed long long) - #endif - unsigned long x, y, z; -}; - -struct ulong4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(ulong4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulong4, signed long long) - #endif - unsigned long x, y, z, w; -} __attribute__((aligned(32))); - - -struct long1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(long1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(long1, signed long long) - #endif - signed long x; -} __attribute__((aligned(8))); - -struct long2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(long2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(long2, signed long long) - #endif - signed long x, y; -} __attribute__((aligned(16))); - -struct long3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(long3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(long3, signed long long) - #endif - signed long x, y, z; -}; - -struct long4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(long4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(long4, signed long long) - #endif - signed long x, y, z, w; -} __attribute__((aligned(32))); - - -struct ulonglong1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(ulonglong1, signed long long) - #endif - unsigned long long x; -} __attribute__((aligned(8))); - -struct ulonglong2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ulonglong2, signed long long) - #endif - unsigned long long x, y; -} __attribute__((aligned(16))); - -struct ulonglong3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(ulonglong3, signed long long) - #endif - unsigned long long x, y, z; -}; - -struct ulonglong4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(ulonglong4, signed long long) - #endif - unsigned long long x, y, z, w; -} __attribute__((aligned(32))); - - -struct longlong1 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(longlong1) - - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, signed char) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, signed short) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, signed int) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, float) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, double) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, signed long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_ONE_COMPONENT(longlong1, signed long long) - #endif - signed long long x; -} __attribute__((aligned(8))); - -struct longlong2 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_TWO_COMPONENT(longlong2) - - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, signed char) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, signed short) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, signed int) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, float) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, double) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, signed long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(longlong2, signed long long) - #endif - signed long long x, y; -} __attribute__((aligned(16))); - -struct longlong3 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_THREE_COMPONENT(longlong3) - - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, signed char) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, signed short) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, signed int) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, float) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, double) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, signed long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_THREE_COMPONENT(longlong3, signed long long) - #endif - signed long long x, y, z; -}; - -struct longlong4 { - #ifdef __cplusplus - public: - MAKE_DEFAULT_CONSTRUCTOR_FOUR_COMPONENT(longlong4) - - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, unsigned char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, signed char) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, unsigned short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, signed short) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, unsigned int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, signed int) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, float) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, double) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, unsigned long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, signed long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, unsigned long long) - MAKE_COMPONENT_CONSTRUCTOR_FOUR_COMPONENT(longlong4, signed long long) - #endif - signed long x, y, z, w; -} __attribute__((aligned(32))); +#if __cplusplus + +typedef unsigned char uchar1 __attribute__((ext_vector_type(1))); +typedef unsigned char uchar2 __attribute__((ext_vector_type(2))); +typedef unsigned char uchar3 __attribute__((ext_vector_type(3))); +typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); + +typedef signed char char1 __attribute__((ext_vector_type(1))); +typedef signed char char2 __attribute__((ext_vector_type(2))); +typedef signed char char3 __attribute__((ext_vector_type(3))); +typedef signed char char4 __attribute__((ext_vector_type(4))); + +typedef unsigned short ushort1 __attribute__((ext_vector_type(1))); +typedef unsigned short ushort2 __attribute__((ext_vector_type(2))); +typedef unsigned short ushort3 __attribute__((ext_vector_type(3))); +typedef unsigned short ushort4 __attribute__((ext_vector_type(4))); + +typedef signed short short1 __attribute__((ext_vector_type(1))); +typedef signed short short2 __attribute__((ext_vector_type(2))); +typedef signed short short3 __attribute__((ext_vector_type(3))); +typedef signed short short4 __attribute__((ext_vector_type(4))); + +typedef __fp16 __half; + +typedef __fp16 __half1 __attribute__((ext_vector_type(1))); +typedef __fp16 __half2 __attribute__((ext_vector_type(2))); +typedef __fp16 __half3 __attribute__((ext_vector_type(3))); +typedef __fp16 __half4 __attribute__((ext_vector_type(4))); + +typedef unsigned int uint1 __attribute__((ext_vector_type(1))); +typedef unsigned int uint2 __attribute__((ext_vector_type(2))); +typedef unsigned int uint3 __attribute__((ext_vector_type(3))); +typedef unsigned int uint4 __attribute__((ext_vector_type(4))); + +typedef signed int int1 __attribute__((ext_vector_type(1))); +typedef signed int int2 __attribute__((ext_vector_type(2))); +typedef signed int int3 __attribute__((ext_vector_type(3))); +typedef signed int int4 __attribute__((ext_vector_type(4))); + +typedef float float1 __attribute__((ext_vector_type(1))); +typedef float float2 __attribute__((ext_vector_type(2))); +typedef float float3 __attribute__((ext_vector_type(3))); +typedef float float4 __attribute__((ext_vector_type(4))); + +typedef unsigned long ulong1 __attribute__((ext_vector_type(1))); +typedef unsigned long ulong2 __attribute__((ext_vector_type(2))); +typedef unsigned long ulong3 __attribute__((ext_vector_type(3))); +typedef unsigned long ulong4 __attribute__((ext_vector_type(4))); + +typedef signed long long1 __attribute__((ext_vector_type(1))); +typedef signed long long2 __attribute__((ext_vector_type(2))); +typedef signed long long3 __attribute__((ext_vector_type(3))); +typedef signed long long4 __attribute__((ext_vector_type(4))); + +typedef double double1 __attribute__((ext_vector_type(1))); +typedef double double2 __attribute__((ext_vector_type(2))); +typedef double double3 __attribute__((ext_vector_type(3))); +typedef double double4 __attribute__((ext_vector_type(4))); + +typedef unsigned long long ulonglong1 __attribute__((ext_vector_type(1))); +typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2))); +typedef unsigned long long ulonglong3 __attribute__((ext_vector_type(3))); +typedef unsigned long long ulonglong4 __attribute__((ext_vector_type(4))); + +typedef signed long long longlong1 __attribute__((ext_vector_type(1))); +typedef signed long long longlong2 __attribute__((ext_vector_type(2))); +typedef signed long long longlong3 __attribute__((ext_vector_type(3))); +typedef signed long long longlong4 __attribute__((ext_vector_type(4))); #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ -__device__ __host__ static inline struct type make_##type(comp x) { \ - struct type ret; \ +__device__ __host__ static inline type make_##type(comp x) { \ + type ret; \ ret.x = x; \ return ret; \ } #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ -__device__ __host__ static inline struct type make_##type(comp x, comp y) { \ - struct type ret; \ +__device__ __host__ static inline type make_##type(comp x, comp y) { \ + type ret; \ ret.x = x; \ ret.y = y; \ return ret; \ } #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ -__device__ __host__ static inline struct type make_##type(comp x, comp y, comp z) { \ - struct type ret; \ +__device__ __host__ static inline type make_##type(comp x, comp y, comp z) { \ + type ret; \ ret.x = x; \ ret.y = y; \ ret.z = z; \ @@ -1155,8 +128,8 @@ __device__ __host__ static inline struct type make_##type(comp x, comp y, comp z } #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ -__device__ __host__ static inline struct type make_##type(comp x, comp y, comp z, comp w) { \ - struct type ret; \ +__device__ __host__ static inline type make_##type(comp x, comp y, comp z, comp w) { \ + type ret; \ ret.x = x; \ ret.y = y; \ ret.z = z; \ @@ -1164,6 +137,7 @@ __device__ __host__ static inline struct type make_##type(comp x, comp y, comp z return ret; \ } + DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1); DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2); DECLOP_MAKE_THREE_COMPONENT(unsigned char, uchar3); @@ -1225,2894 +199,9 @@ DECLOP_MAKE_THREE_COMPONENT(signed long, longlong3); DECLOP_MAKE_FOUR_COMPONENT(signed long, longlong4); -#if __cplusplus - -#define DECLOP_1VAR_2IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs.x op rhs.x; \ - return ret; \ -} - -#define DECLOP_1VAR_SCALE_PRODUCT(type, type1) \ -__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \ - type ret; \ - ret.x = lhs.x * rhs; \ - return ret; \ -} \ -\ -__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs * rhs.x; \ - return ret; \ -} - -#define DECLOP_1VAR_ASSIGN(type, op) \ -__device__ __host__ static inline type& operator op ( type& lhs, const type& rhs) { \ - lhs.x op rhs.x; \ - return lhs; \ -} - -#define DECLOP_1VAR_PREOP(type, op) \ -__device__ __host__ static inline type& operator op (type& val) { \ - op val.x; \ - return val; \ -} - -#define DECLOP_1VAR_POSTOP(type, op) \ -__device__ __host__ static inline type operator op (type& val, int) { \ - type ret; \ - ret.x = val.x; \ - val.x op; \ - return ret; \ -} - -#define DECLOP_1VAR_COMP(type, op) \ -__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \ - return lhs.x op rhs.x; \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \ - return lhs.x op rhs.x; \ -} \ -__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \ - return lhs.x op rhs.x ; \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \ - return lhs.x op rhs.x ; \ -} - -#define DECLOP_1VAR_1IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op(type& rhs) { \ - type ret; \ - ret.x = op rhs.x; \ - return ret; \ -} - -#define DECLOP_1VAR_1IN_BOOLOUT(type, op) \ -__device__ __host__ static inline bool operator op (type& rhs) { \ - return op rhs.x; \ -} - -/* - Two Element Access -*/ - -#define DECLOP_2VAR_2IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs.x op rhs.x; \ - ret.y = lhs.y op rhs.y; \ - return ret; \ -} - -#define DECLOP_2VAR_SCALE_PRODUCT(type, type1) \ -__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \ - type ret; \ - ret.x = lhs.x * rhs; \ - ret.y = lhs.y * rhs; \ - return ret; \ -} \ -\ -__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs * rhs.x; \ - ret.y = lhs * rhs.y; \ - return ret; \ -} - -#define DECLOP_2VAR_ASSIGN(type, op) \ -__device__ __host__ static inline type& operator op ( type& lhs, const type& rhs) { \ - lhs.x op rhs.x; \ - lhs.y op rhs.y; \ - return lhs; \ -} - -#define DECLOP_2VAR_PREOP(type, op) \ -__device__ __host__ static inline type& operator op (type& val) { \ - op val.x; \ - op val.y; \ - return val; \ -} - -#define DECLOP_2VAR_POSTOP(type, op) \ -__device__ __host__ static inline type operator op (type& val, int) { \ - type ret; \ - ret.x = val.x; \ - ret.y = val.y; \ - val.x op; \ - val.y op; \ - return ret; \ -} - -#define DECLOP_2VAR_COMP(type, op) \ -__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y); \ -} \ -__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y); \ -} - -#define DECLOP_2VAR_1IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op(type &rhs) { \ - type ret; \ - ret.x = op rhs.x; \ - ret.y = op rhs.y; \ - return ret; \ -} - -#define DECLOP_2VAR_1IN_BOOLOUT(type, op) \ -__device__ __host__ static inline bool operator op (type &rhs) { \ - return (op rhs.x) && (op rhs.y); \ -} - - -/* - Three Element Access -*/ - -#define DECLOP_3VAR_2IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs.x op rhs.x; \ - ret.y = lhs.y op rhs.y; \ - ret.z = lhs.z op rhs.z; \ - return ret; \ -} - -#define DECLOP_3VAR_SCALE_PRODUCT(type, type1) \ -__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \ - type ret; \ - ret.x = lhs.x * rhs; \ - ret.y = lhs.y * rhs; \ - ret.z = lhs.z * rhs; \ - return ret; \ -} \ -\ -__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs * rhs.x; \ - ret.y = lhs * rhs.y; \ - ret.z = lhs * rhs.z; \ - return ret; \ -} - -#define DECLOP_3VAR_ASSIGN(type, op) \ -__device__ __host__ static inline type& operator op ( type& lhs, const type& rhs) { \ - lhs.x op rhs.x; \ - lhs.y op rhs.y; \ - lhs.z op rhs.z; \ - return lhs; \ -} - -#define DECLOP_3VAR_PREOP(type, op) \ -__device__ __host__ static inline type& operator op (type& val) { \ - op val.x; \ - op val.y; \ - op val.z; \ - return val; \ -} - -#define DECLOP_3VAR_POSTOP(type, op) \ -__device__ __host__ static inline type operator op (type& val, int) { \ - type ret; \ - ret.x = val.x; \ - ret.y = val.y; \ - ret.z = val.z; \ - val.x op; \ - val.y op; \ - val.z op; \ - return ret; \ -} - -#define DECLOP_3VAR_COMP(type, op) \ -__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \ -} \ -__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z); \ -} \ - -#define DECLOP_3VAR_1IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op(type &rhs) { \ - type ret; \ - ret.x = op rhs.x; \ - ret.y = op rhs.y; \ - ret.z = op rhs.z; \ - return ret; \ -} - -#define DECLOP_3VAR_1IN_BOOLOUT(type, op) \ -__device__ __host__ static inline bool operator op (type &rhs) { \ - return (op rhs.x) && (op rhs.y) && (op rhs.z); \ -} - - -/* - Four Element Access -*/ - -#define DECLOP_4VAR_2IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op ( const type& lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs.x op rhs.x; \ - ret.y = lhs.y op rhs.y; \ - ret.z = lhs.z op rhs.z; \ - ret.w = lhs.w op rhs.w; \ - return ret; \ -} - -#define DECLOP_4VAR_SCALE_PRODUCT(type, type1) \ -__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \ - type ret; \ - ret.x = lhs.x * rhs; \ - ret.y = lhs.y * rhs; \ - ret.z = lhs.z * rhs; \ - ret.w = lhs.w * rhs; \ - return ret; \ -} \ -\ -__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \ - type ret; \ - ret.x = lhs * rhs.x; \ - ret.y = lhs * rhs.y; \ - ret.z = lhs * rhs.z; \ - ret.w = lhs * rhs.w; \ - return ret; \ -} - -#define DECLOP_4VAR_ASSIGN(type, op) \ -__device__ __host__ static inline type& operator op ( type& lhs, const type& rhs) { \ - lhs.x op rhs.x; \ - lhs.y op rhs.y; \ - lhs.z op rhs.z; \ - lhs.w op rhs.w; \ - return lhs; \ -} - -#define DECLOP_4VAR_PREOP(type, op) \ -__device__ __host__ static inline type& operator op (type& val) { \ - op val.x; \ - op val.y; \ - op val.z; \ - op val.w; \ - return val; \ -} - -#define DECLOP_4VAR_POSTOP(type, op) \ -__device__ __host__ static inline type operator op (type& val, int) { \ - type ret; \ - ret.x = val.x; \ - ret.y = val.y; \ - ret.z = val.z; \ - ret.w = val.w; \ - val.x op; \ - val.y op; \ - val.z op; \ - val.w op; \ - return ret; \ -} - -#define DECLOP_4VAR_COMP(type, op) \ -__device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \ -} \ -__device__ __host__ static inline bool operator op (type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \ -} \ -__device__ __host__ static inline bool operator op (const type& lhs, const type& rhs) { \ - return (lhs.x op rhs.x) && (lhs.y op rhs.y) && (lhs.z op rhs.z) && (lhs.w op rhs.w); \ -} - -#define DECLOP_4VAR_1IN_1OUT(type, op) \ -__device__ __host__ static inline type operator op(type &rhs) { \ - type ret; \ - ret.x = op rhs.x; \ - ret.y = op rhs.y; \ - ret.z = op rhs.z; \ - ret.w = op rhs.w; \ - return ret; \ -} - -#define DECLOP_4VAR_1IN_BOOLOUT(type, op) \ -__device__ __host__ static inline bool operator op (type &rhs) { \ - return (op rhs.x) && (op rhs.y) && (op rhs.z) && (op rhs.w); \ -} - - -/* -Overloading operators -*/ - -// UNSIGNED CHAR1 - -DECLOP_1VAR_2IN_1OUT(uchar1, +) -DECLOP_1VAR_2IN_1OUT(uchar1, -) -DECLOP_1VAR_2IN_1OUT(uchar1, *) -DECLOP_1VAR_2IN_1OUT(uchar1, /) -DECLOP_1VAR_2IN_1OUT(uchar1, %) -DECLOP_1VAR_2IN_1OUT(uchar1, &) -DECLOP_1VAR_2IN_1OUT(uchar1, |) -DECLOP_1VAR_2IN_1OUT(uchar1, ^) -DECLOP_1VAR_2IN_1OUT(uchar1, <<) -DECLOP_1VAR_2IN_1OUT(uchar1, >>) - - -DECLOP_1VAR_ASSIGN(uchar1, +=) -DECLOP_1VAR_ASSIGN(uchar1, -=) -DECLOP_1VAR_ASSIGN(uchar1, *=) -DECLOP_1VAR_ASSIGN(uchar1, /=) -DECLOP_1VAR_ASSIGN(uchar1, %=) -DECLOP_1VAR_ASSIGN(uchar1, &=) -DECLOP_1VAR_ASSIGN(uchar1, |=) -DECLOP_1VAR_ASSIGN(uchar1, ^=) -DECLOP_1VAR_ASSIGN(uchar1, <<=) -DECLOP_1VAR_ASSIGN(uchar1, >>=) - -DECLOP_1VAR_PREOP(uchar1, ++) -DECLOP_1VAR_PREOP(uchar1, --) - -DECLOP_1VAR_POSTOP(uchar1, ++) -DECLOP_1VAR_POSTOP(uchar1, --) - -DECLOP_1VAR_COMP(uchar1, ==) -DECLOP_1VAR_COMP(uchar1, !=) -DECLOP_1VAR_COMP(uchar1, <) -DECLOP_1VAR_COMP(uchar1, >) -DECLOP_1VAR_COMP(uchar1, <=) -DECLOP_1VAR_COMP(uchar1, >=) - -DECLOP_1VAR_COMP(uchar1, &&) -DECLOP_1VAR_COMP(uchar1, ||) - -DECLOP_1VAR_1IN_1OUT(uchar1, ~) -DECLOP_1VAR_1IN_BOOLOUT(uchar1, !) - -DECLOP_1VAR_SCALE_PRODUCT(uchar1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, float) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, double) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(uchar1, signed long long) - -// UNSIGNED CHAR2 - -DECLOP_2VAR_2IN_1OUT(uchar2, +) -DECLOP_2VAR_2IN_1OUT(uchar2, -) -DECLOP_2VAR_2IN_1OUT(uchar2, *) -DECLOP_2VAR_2IN_1OUT(uchar2, /) -DECLOP_2VAR_2IN_1OUT(uchar2, %) -DECLOP_2VAR_2IN_1OUT(uchar2, &) -DECLOP_2VAR_2IN_1OUT(uchar2, |) -DECLOP_2VAR_2IN_1OUT(uchar2, ^) -DECLOP_2VAR_2IN_1OUT(uchar2, <<) -DECLOP_2VAR_2IN_1OUT(uchar2, >>) - -DECLOP_2VAR_ASSIGN(uchar2, +=) -DECLOP_2VAR_ASSIGN(uchar2, -=) -DECLOP_2VAR_ASSIGN(uchar2, *=) -DECLOP_2VAR_ASSIGN(uchar2, /=) -DECLOP_2VAR_ASSIGN(uchar2, %=) -DECLOP_2VAR_ASSIGN(uchar2, &=) -DECLOP_2VAR_ASSIGN(uchar2, |=) -DECLOP_2VAR_ASSIGN(uchar2, ^=) -DECLOP_2VAR_ASSIGN(uchar2, <<=) -DECLOP_2VAR_ASSIGN(uchar2, >>=) - -DECLOP_2VAR_PREOP(uchar2, ++) -DECLOP_2VAR_PREOP(uchar2, --) - -DECLOP_2VAR_POSTOP(uchar2, ++) -DECLOP_2VAR_POSTOP(uchar2, --) - -DECLOP_2VAR_COMP(uchar2, ==) -DECLOP_2VAR_COMP(uchar2, !=) -DECLOP_2VAR_COMP(uchar2, <) -DECLOP_2VAR_COMP(uchar2, >) -DECLOP_2VAR_COMP(uchar2, <=) -DECLOP_2VAR_COMP(uchar2, >=) - -DECLOP_2VAR_COMP(uchar2, &&) -DECLOP_2VAR_COMP(uchar2, ||) - -DECLOP_2VAR_1IN_1OUT(uchar2, ~) -DECLOP_2VAR_1IN_BOOLOUT(uchar2, !) - -DECLOP_2VAR_SCALE_PRODUCT(uchar2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, float) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, double) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(uchar2, signed long long) - -// UNSIGNED CHAR3 - -DECLOP_3VAR_2IN_1OUT(uchar3, +) -DECLOP_3VAR_2IN_1OUT(uchar3, -) -DECLOP_3VAR_2IN_1OUT(uchar3, *) -DECLOP_3VAR_2IN_1OUT(uchar3, /) -DECLOP_3VAR_2IN_1OUT(uchar3, %) -DECLOP_3VAR_2IN_1OUT(uchar3, &) -DECLOP_3VAR_2IN_1OUT(uchar3, |) -DECLOP_3VAR_2IN_1OUT(uchar3, ^) -DECLOP_3VAR_2IN_1OUT(uchar3, <<) -DECLOP_3VAR_2IN_1OUT(uchar3, >>) - -DECLOP_3VAR_ASSIGN(uchar3, +=) -DECLOP_3VAR_ASSIGN(uchar3, -=) -DECLOP_3VAR_ASSIGN(uchar3, *=) -DECLOP_3VAR_ASSIGN(uchar3, /=) -DECLOP_3VAR_ASSIGN(uchar3, %=) -DECLOP_3VAR_ASSIGN(uchar3, &=) -DECLOP_3VAR_ASSIGN(uchar3, |=) -DECLOP_3VAR_ASSIGN(uchar3, ^=) -DECLOP_3VAR_ASSIGN(uchar3, <<=) -DECLOP_3VAR_ASSIGN(uchar3, >>=) - -DECLOP_3VAR_PREOP(uchar3, ++) -DECLOP_3VAR_PREOP(uchar3, --) - -DECLOP_3VAR_POSTOP(uchar3, ++) -DECLOP_3VAR_POSTOP(uchar3, --) - -DECLOP_3VAR_COMP(uchar3, ==) -DECLOP_3VAR_COMP(uchar3, !=) -DECLOP_3VAR_COMP(uchar3, <) -DECLOP_3VAR_COMP(uchar3, >) -DECLOP_3VAR_COMP(uchar3, <=) -DECLOP_3VAR_COMP(uchar3, >=) - -DECLOP_3VAR_COMP(uchar3, &&) -DECLOP_3VAR_COMP(uchar3, ||) - -DECLOP_3VAR_1IN_1OUT(uchar3, ~) -DECLOP_3VAR_1IN_BOOLOUT(uchar3, !) - -DECLOP_3VAR_SCALE_PRODUCT(uchar3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, float) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, double) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(uchar3, signed long long) - -// UNSIGNED CHAR4 - -DECLOP_4VAR_2IN_1OUT(uchar4, +) -DECLOP_4VAR_2IN_1OUT(uchar4, -) -DECLOP_4VAR_2IN_1OUT(uchar4, *) -DECLOP_4VAR_2IN_1OUT(uchar4, /) -DECLOP_4VAR_2IN_1OUT(uchar4, %) -DECLOP_4VAR_2IN_1OUT(uchar4, &) -DECLOP_4VAR_2IN_1OUT(uchar4, |) -DECLOP_4VAR_2IN_1OUT(uchar4, ^) -DECLOP_4VAR_2IN_1OUT(uchar4, <<) -DECLOP_4VAR_2IN_1OUT(uchar4, >>) - -DECLOP_4VAR_ASSIGN(uchar4, +=) -DECLOP_4VAR_ASSIGN(uchar4, -=) -DECLOP_4VAR_ASSIGN(uchar4, *=) -DECLOP_4VAR_ASSIGN(uchar4, /=) -DECLOP_4VAR_ASSIGN(uchar4, %=) -DECLOP_4VAR_ASSIGN(uchar4, &=) -DECLOP_4VAR_ASSIGN(uchar4, |=) -DECLOP_4VAR_ASSIGN(uchar4, ^=) -DECLOP_4VAR_ASSIGN(uchar4, <<=) -DECLOP_4VAR_ASSIGN(uchar4, >>=) - -DECLOP_4VAR_PREOP(uchar4, ++) -DECLOP_4VAR_PREOP(uchar4, --) - -DECLOP_4VAR_POSTOP(uchar4, ++) -DECLOP_4VAR_POSTOP(uchar4, --) - -DECLOP_4VAR_COMP(uchar4, ==) -DECLOP_4VAR_COMP(uchar4, !=) -DECLOP_4VAR_COMP(uchar4, <) -DECLOP_4VAR_COMP(uchar4, >) -DECLOP_4VAR_COMP(uchar4, <=) -DECLOP_4VAR_COMP(uchar4, >=) - -DECLOP_4VAR_COMP(uchar4, &&) -DECLOP_4VAR_COMP(uchar4, ||) - -DECLOP_4VAR_1IN_1OUT(uchar4, ~) -DECLOP_4VAR_1IN_BOOLOUT(uchar4, !) - -DECLOP_4VAR_SCALE_PRODUCT(uchar4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, float) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, double) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(uchar4, signed long long) - -// SIGNED CHAR1 - -DECLOP_1VAR_2IN_1OUT(char1, +) -DECLOP_1VAR_2IN_1OUT(char1, -) -DECLOP_1VAR_2IN_1OUT(char1, *) -DECLOP_1VAR_2IN_1OUT(char1, /) -DECLOP_1VAR_2IN_1OUT(char1, %) -DECLOP_1VAR_2IN_1OUT(char1, &) -DECLOP_1VAR_2IN_1OUT(char1, |) -DECLOP_1VAR_2IN_1OUT(char1, ^) -DECLOP_1VAR_2IN_1OUT(char1, <<) -DECLOP_1VAR_2IN_1OUT(char1, >>) - - -DECLOP_1VAR_ASSIGN(char1, +=) -DECLOP_1VAR_ASSIGN(char1, -=) -DECLOP_1VAR_ASSIGN(char1, *=) -DECLOP_1VAR_ASSIGN(char1, /=) -DECLOP_1VAR_ASSIGN(char1, %=) -DECLOP_1VAR_ASSIGN(char1, &=) -DECLOP_1VAR_ASSIGN(char1, |=) -DECLOP_1VAR_ASSIGN(char1, ^=) -DECLOP_1VAR_ASSIGN(char1, <<=) -DECLOP_1VAR_ASSIGN(char1, >>=) - -DECLOP_1VAR_PREOP(char1, ++) -DECLOP_1VAR_PREOP(char1, --) - -DECLOP_1VAR_POSTOP(char1, ++) -DECLOP_1VAR_POSTOP(char1, --) - -DECLOP_1VAR_COMP(char1, ==) -DECLOP_1VAR_COMP(char1, !=) -DECLOP_1VAR_COMP(char1, <) -DECLOP_1VAR_COMP(char1, >) -DECLOP_1VAR_COMP(char1, <=) -DECLOP_1VAR_COMP(char1, >=) - -DECLOP_1VAR_COMP(char1, &&) -DECLOP_1VAR_COMP(char1, ||) - -DECLOP_1VAR_1IN_1OUT(char1, ~) -DECLOP_1VAR_1IN_BOOLOUT(char1, !) - -DECLOP_1VAR_SCALE_PRODUCT(char1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(char1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(char1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(char1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(char1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(char1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(char1, float) -DECLOP_1VAR_SCALE_PRODUCT(char1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(char1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(char1, double) -DECLOP_1VAR_SCALE_PRODUCT(char1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(char1, signed long long) - -// SIGNED CHAR2 - -DECLOP_2VAR_2IN_1OUT(char2, +) -DECLOP_2VAR_2IN_1OUT(char2, -) -DECLOP_2VAR_2IN_1OUT(char2, *) -DECLOP_2VAR_2IN_1OUT(char2, /) -DECLOP_2VAR_2IN_1OUT(char2, %) -DECLOP_2VAR_2IN_1OUT(char2, &) -DECLOP_2VAR_2IN_1OUT(char2, |) -DECLOP_2VAR_2IN_1OUT(char2, ^) -DECLOP_2VAR_2IN_1OUT(char2, <<) -DECLOP_2VAR_2IN_1OUT(char2, >>) - -DECLOP_2VAR_ASSIGN(char2, +=) -DECLOP_2VAR_ASSIGN(char2, -=) -DECLOP_2VAR_ASSIGN(char2, *=) -DECLOP_2VAR_ASSIGN(char2, /=) -DECLOP_2VAR_ASSIGN(char2, %=) -DECLOP_2VAR_ASSIGN(char2, &=) -DECLOP_2VAR_ASSIGN(char2, |=) -DECLOP_2VAR_ASSIGN(char2, ^=) -DECLOP_2VAR_ASSIGN(char2, <<=) -DECLOP_2VAR_ASSIGN(char2, >>=) - -DECLOP_2VAR_PREOP(char2, ++) -DECLOP_2VAR_PREOP(char2, --) - -DECLOP_2VAR_POSTOP(char2, ++) -DECLOP_2VAR_POSTOP(char2, --) - -DECLOP_2VAR_COMP(char2, ==) -DECLOP_2VAR_COMP(char2, !=) -DECLOP_2VAR_COMP(char2, <) -DECLOP_2VAR_COMP(char2, >) -DECLOP_2VAR_COMP(char2, <=) -DECLOP_2VAR_COMP(char2, >=) - -DECLOP_2VAR_COMP(char2, &&) -DECLOP_2VAR_COMP(char2, ||) - -DECLOP_2VAR_1IN_1OUT(char2, ~) -DECLOP_2VAR_1IN_BOOLOUT(char2, !) - -DECLOP_2VAR_SCALE_PRODUCT(char2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(char2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(char2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(char2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(char2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(char2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(char2, float) -DECLOP_2VAR_SCALE_PRODUCT(char2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(char2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(char2, double) -DECLOP_2VAR_SCALE_PRODUCT(char2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(char2, signed long long) - -// SIGNED CHAR3 - -DECLOP_3VAR_2IN_1OUT(char3, +) -DECLOP_3VAR_2IN_1OUT(char3, -) -DECLOP_3VAR_2IN_1OUT(char3, *) -DECLOP_3VAR_2IN_1OUT(char3, /) -DECLOP_3VAR_2IN_1OUT(char3, %) -DECLOP_3VAR_2IN_1OUT(char3, &) -DECLOP_3VAR_2IN_1OUT(char3, |) -DECLOP_3VAR_2IN_1OUT(char3, ^) -DECLOP_3VAR_2IN_1OUT(char3, <<) -DECLOP_3VAR_2IN_1OUT(char3, >>) - -DECLOP_3VAR_ASSIGN(char3, +=) -DECLOP_3VAR_ASSIGN(char3, -=) -DECLOP_3VAR_ASSIGN(char3, *=) -DECLOP_3VAR_ASSIGN(char3, /=) -DECLOP_3VAR_ASSIGN(char3, %=) -DECLOP_3VAR_ASSIGN(char3, &=) -DECLOP_3VAR_ASSIGN(char3, |=) -DECLOP_3VAR_ASSIGN(char3, ^=) -DECLOP_3VAR_ASSIGN(char3, <<=) -DECLOP_3VAR_ASSIGN(char3, >>=) - -DECLOP_3VAR_PREOP(char3, ++) -DECLOP_3VAR_PREOP(char3, --) - -DECLOP_3VAR_POSTOP(char3, ++) -DECLOP_3VAR_POSTOP(char3, --) - -DECLOP_3VAR_COMP(char3, ==) -DECLOP_3VAR_COMP(char3, !=) -DECLOP_3VAR_COMP(char3, <) -DECLOP_3VAR_COMP(char3, >) -DECLOP_3VAR_COMP(char3, <=) -DECLOP_3VAR_COMP(char3, >=) - -DECLOP_3VAR_COMP(char3, &&) -DECLOP_3VAR_COMP(char3, ||) - -DECLOP_3VAR_1IN_1OUT(char3, ~) -DECLOP_3VAR_1IN_BOOLOUT(char3, !) - -DECLOP_3VAR_SCALE_PRODUCT(char3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(char3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(char3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(char3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(char3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(char3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(char3, float) -DECLOP_3VAR_SCALE_PRODUCT(char3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(char3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(char3, double) -DECLOP_3VAR_SCALE_PRODUCT(char3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(char3, signed long long) - -// SIGNED CHAR4 - -DECLOP_4VAR_2IN_1OUT(char4, +) -DECLOP_4VAR_2IN_1OUT(char4, -) -DECLOP_4VAR_2IN_1OUT(char4, *) -DECLOP_4VAR_2IN_1OUT(char4, /) -DECLOP_4VAR_2IN_1OUT(char4, %) -DECLOP_4VAR_2IN_1OUT(char4, &) -DECLOP_4VAR_2IN_1OUT(char4, |) -DECLOP_4VAR_2IN_1OUT(char4, ^) -DECLOP_4VAR_2IN_1OUT(char4, <<) -DECLOP_4VAR_2IN_1OUT(char4, >>) - -DECLOP_4VAR_ASSIGN(char4, +=) -DECLOP_4VAR_ASSIGN(char4, -=) -DECLOP_4VAR_ASSIGN(char4, *=) -DECLOP_4VAR_ASSIGN(char4, /=) -DECLOP_4VAR_ASSIGN(char4, %=) -DECLOP_4VAR_ASSIGN(char4, &=) -DECLOP_4VAR_ASSIGN(char4, |=) -DECLOP_4VAR_ASSIGN(char4, ^=) -DECLOP_4VAR_ASSIGN(char4, <<=) -DECLOP_4VAR_ASSIGN(char4, >>=) - -DECLOP_4VAR_PREOP(char4, ++) -DECLOP_4VAR_PREOP(char4, --) - -DECLOP_4VAR_POSTOP(char4, ++) -DECLOP_4VAR_POSTOP(char4, --) - -DECLOP_4VAR_COMP(char4, ==) -DECLOP_4VAR_COMP(char4, !=) -DECLOP_4VAR_COMP(char4, <) -DECLOP_4VAR_COMP(char4, >) -DECLOP_4VAR_COMP(char4, <=) -DECLOP_4VAR_COMP(char4, >=) - -DECLOP_4VAR_COMP(char4, &&) -DECLOP_4VAR_COMP(char4, ||) - -DECLOP_4VAR_1IN_1OUT(char4, ~) -DECLOP_4VAR_1IN_BOOLOUT(char4, !) - -DECLOP_4VAR_SCALE_PRODUCT(char4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(char4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(char4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(char4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(char4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(char4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(char4, float) -DECLOP_4VAR_SCALE_PRODUCT(char4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(char4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(char4, double) -DECLOP_4VAR_SCALE_PRODUCT(char4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(char4, signed long long) - -// UNSIGNED SHORT1 - -DECLOP_1VAR_2IN_1OUT(ushort1, +) -DECLOP_1VAR_2IN_1OUT(ushort1, -) -DECLOP_1VAR_2IN_1OUT(ushort1, *) -DECLOP_1VAR_2IN_1OUT(ushort1, /) -DECLOP_1VAR_2IN_1OUT(ushort1, %) -DECLOP_1VAR_2IN_1OUT(ushort1, &) -DECLOP_1VAR_2IN_1OUT(ushort1, |) -DECLOP_1VAR_2IN_1OUT(ushort1, ^) -DECLOP_1VAR_2IN_1OUT(ushort1, <<) -DECLOP_1VAR_2IN_1OUT(ushort1, >>) - - -DECLOP_1VAR_ASSIGN(ushort1, +=) -DECLOP_1VAR_ASSIGN(ushort1, -=) -DECLOP_1VAR_ASSIGN(ushort1, *=) -DECLOP_1VAR_ASSIGN(ushort1, /=) -DECLOP_1VAR_ASSIGN(ushort1, %=) -DECLOP_1VAR_ASSIGN(ushort1, &=) -DECLOP_1VAR_ASSIGN(ushort1, |=) -DECLOP_1VAR_ASSIGN(ushort1, ^=) -DECLOP_1VAR_ASSIGN(ushort1, <<=) -DECLOP_1VAR_ASSIGN(ushort1, >>=) - -DECLOP_1VAR_PREOP(ushort1, ++) -DECLOP_1VAR_PREOP(ushort1, --) - -DECLOP_1VAR_POSTOP(ushort1, ++) -DECLOP_1VAR_POSTOP(ushort1, --) - -DECLOP_1VAR_COMP(ushort1, ==) -DECLOP_1VAR_COMP(ushort1, !=) -DECLOP_1VAR_COMP(ushort1, <) -DECLOP_1VAR_COMP(ushort1, >) -DECLOP_1VAR_COMP(ushort1, <=) -DECLOP_1VAR_COMP(ushort1, >=) - -DECLOP_1VAR_COMP(ushort1, &&) -DECLOP_1VAR_COMP(ushort1, ||) - -DECLOP_1VAR_1IN_1OUT(ushort1, ~) -DECLOP_1VAR_1IN_BOOLOUT(ushort1, !) - -DECLOP_1VAR_SCALE_PRODUCT(ushort1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, float) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, double) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(ushort1, signed long long) - -// UNSIGNED SHORT2 - -DECLOP_2VAR_2IN_1OUT(ushort2, +) -DECLOP_2VAR_2IN_1OUT(ushort2, -) -DECLOP_2VAR_2IN_1OUT(ushort2, *) -DECLOP_2VAR_2IN_1OUT(ushort2, /) -DECLOP_2VAR_2IN_1OUT(ushort2, %) -DECLOP_2VAR_2IN_1OUT(ushort2, &) -DECLOP_2VAR_2IN_1OUT(ushort2, |) -DECLOP_2VAR_2IN_1OUT(ushort2, ^) -DECLOP_2VAR_2IN_1OUT(ushort2, <<) -DECLOP_2VAR_2IN_1OUT(ushort2, >>) - -DECLOP_2VAR_ASSIGN(ushort2, +=) -DECLOP_2VAR_ASSIGN(ushort2, -=) -DECLOP_2VAR_ASSIGN(ushort2, *=) -DECLOP_2VAR_ASSIGN(ushort2, /=) -DECLOP_2VAR_ASSIGN(ushort2, %=) -DECLOP_2VAR_ASSIGN(ushort2, &=) -DECLOP_2VAR_ASSIGN(ushort2, |=) -DECLOP_2VAR_ASSIGN(ushort2, ^=) -DECLOP_2VAR_ASSIGN(ushort2, <<=) -DECLOP_2VAR_ASSIGN(ushort2, >>=) - -DECLOP_2VAR_PREOP(ushort2, ++) -DECLOP_2VAR_PREOP(ushort2, --) - -DECLOP_2VAR_POSTOP(ushort2, ++) -DECLOP_2VAR_POSTOP(ushort2, --) - -DECLOP_2VAR_COMP(ushort2, ==) -DECLOP_2VAR_COMP(ushort2, !=) -DECLOP_2VAR_COMP(ushort2, <) -DECLOP_2VAR_COMP(ushort2, >) -DECLOP_2VAR_COMP(ushort2, <=) -DECLOP_2VAR_COMP(ushort2, >=) - -DECLOP_2VAR_COMP(ushort2, &&) -DECLOP_2VAR_COMP(ushort2, ||) - -DECLOP_2VAR_1IN_1OUT(ushort2, ~) -DECLOP_2VAR_1IN_BOOLOUT(ushort2, !) - -DECLOP_2VAR_SCALE_PRODUCT(ushort2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, float) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, double) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(ushort2, signed long long) - -// UNSIGNED SHORT3 - -DECLOP_3VAR_2IN_1OUT(ushort3, +) -DECLOP_3VAR_2IN_1OUT(ushort3, -) -DECLOP_3VAR_2IN_1OUT(ushort3, *) -DECLOP_3VAR_2IN_1OUT(ushort3, /) -DECLOP_3VAR_2IN_1OUT(ushort3, %) -DECLOP_3VAR_2IN_1OUT(ushort3, &) -DECLOP_3VAR_2IN_1OUT(ushort3, |) -DECLOP_3VAR_2IN_1OUT(ushort3, ^) -DECLOP_3VAR_2IN_1OUT(ushort3, <<) -DECLOP_3VAR_2IN_1OUT(ushort3, >>) - -DECLOP_3VAR_ASSIGN(ushort3, +=) -DECLOP_3VAR_ASSIGN(ushort3, -=) -DECLOP_3VAR_ASSIGN(ushort3, *=) -DECLOP_3VAR_ASSIGN(ushort3, /=) -DECLOP_3VAR_ASSIGN(ushort3, %=) -DECLOP_3VAR_ASSIGN(ushort3, &=) -DECLOP_3VAR_ASSIGN(ushort3, |=) -DECLOP_3VAR_ASSIGN(ushort3, ^=) -DECLOP_3VAR_ASSIGN(ushort3, <<=) -DECLOP_3VAR_ASSIGN(ushort3, >>=) - -DECLOP_3VAR_PREOP(ushort3, ++) -DECLOP_3VAR_PREOP(ushort3, --) - -DECLOP_3VAR_POSTOP(ushort3, ++) -DECLOP_3VAR_POSTOP(ushort3, --) - -DECLOP_3VAR_COMP(ushort3, ==) -DECLOP_3VAR_COMP(ushort3, !=) -DECLOP_3VAR_COMP(ushort3, <) -DECLOP_3VAR_COMP(ushort3, >) -DECLOP_3VAR_COMP(ushort3, <=) -DECLOP_3VAR_COMP(ushort3, >=) - -DECLOP_3VAR_COMP(ushort3, &&) -DECLOP_3VAR_COMP(ushort3, ||) - -DECLOP_3VAR_1IN_1OUT(ushort3, ~) -DECLOP_3VAR_1IN_BOOLOUT(ushort3, !) - -DECLOP_3VAR_SCALE_PRODUCT(ushort3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, float) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, double) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(ushort3, signed long long) - -// UNSIGNED SHORT4 - -DECLOP_4VAR_2IN_1OUT(ushort4, +) -DECLOP_4VAR_2IN_1OUT(ushort4, -) -DECLOP_4VAR_2IN_1OUT(ushort4, *) -DECLOP_4VAR_2IN_1OUT(ushort4, /) -DECLOP_4VAR_2IN_1OUT(ushort4, %) -DECLOP_4VAR_2IN_1OUT(ushort4, &) -DECLOP_4VAR_2IN_1OUT(ushort4, |) -DECLOP_4VAR_2IN_1OUT(ushort4, ^) -DECLOP_4VAR_2IN_1OUT(ushort4, <<) -DECLOP_4VAR_2IN_1OUT(ushort4, >>) - -DECLOP_4VAR_ASSIGN(ushort4, +=) -DECLOP_4VAR_ASSIGN(ushort4, -=) -DECLOP_4VAR_ASSIGN(ushort4, *=) -DECLOP_4VAR_ASSIGN(ushort4, /=) -DECLOP_4VAR_ASSIGN(ushort4, %=) -DECLOP_4VAR_ASSIGN(ushort4, &=) -DECLOP_4VAR_ASSIGN(ushort4, |=) -DECLOP_4VAR_ASSIGN(ushort4, ^=) -DECLOP_4VAR_ASSIGN(ushort4, <<=) -DECLOP_4VAR_ASSIGN(ushort4, >>=) - -DECLOP_4VAR_PREOP(ushort4, ++) -DECLOP_4VAR_PREOP(ushort4, --) - -DECLOP_4VAR_POSTOP(ushort4, ++) -DECLOP_4VAR_POSTOP(ushort4, --) - -DECLOP_4VAR_COMP(ushort4, ==) -DECLOP_4VAR_COMP(ushort4, !=) -DECLOP_4VAR_COMP(ushort4, <) -DECLOP_4VAR_COMP(ushort4, >) -DECLOP_4VAR_COMP(ushort4, <=) -DECLOP_4VAR_COMP(ushort4, >=) - -DECLOP_4VAR_COMP(ushort4, &&) -DECLOP_4VAR_COMP(ushort4, ||) - -DECLOP_4VAR_1IN_1OUT(ushort4, ~) -DECLOP_4VAR_1IN_BOOLOUT(ushort4, !) - -DECLOP_4VAR_SCALE_PRODUCT(ushort4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, float) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, double) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(ushort4, signed long long) - -// SIGNED SHORT1 - -DECLOP_1VAR_2IN_1OUT(short1, +) -DECLOP_1VAR_2IN_1OUT(short1, -) -DECLOP_1VAR_2IN_1OUT(short1, *) -DECLOP_1VAR_2IN_1OUT(short1, /) -DECLOP_1VAR_2IN_1OUT(short1, %) -DECLOP_1VAR_2IN_1OUT(short1, &) -DECLOP_1VAR_2IN_1OUT(short1, |) -DECLOP_1VAR_2IN_1OUT(short1, ^) -DECLOP_1VAR_2IN_1OUT(short1, <<) -DECLOP_1VAR_2IN_1OUT(short1, >>) - - -DECLOP_1VAR_ASSIGN(short1, +=) -DECLOP_1VAR_ASSIGN(short1, -=) -DECLOP_1VAR_ASSIGN(short1, *=) -DECLOP_1VAR_ASSIGN(short1, /=) -DECLOP_1VAR_ASSIGN(short1, %=) -DECLOP_1VAR_ASSIGN(short1, &=) -DECLOP_1VAR_ASSIGN(short1, |=) -DECLOP_1VAR_ASSIGN(short1, ^=) -DECLOP_1VAR_ASSIGN(short1, <<=) -DECLOP_1VAR_ASSIGN(short1, >>=) - -DECLOP_1VAR_PREOP(short1, ++) -DECLOP_1VAR_PREOP(short1, --) - -DECLOP_1VAR_POSTOP(short1, ++) -DECLOP_1VAR_POSTOP(short1, --) - -DECLOP_1VAR_COMP(short1, ==) -DECLOP_1VAR_COMP(short1, !=) -DECLOP_1VAR_COMP(short1, <) -DECLOP_1VAR_COMP(short1, >) -DECLOP_1VAR_COMP(short1, <=) -DECLOP_1VAR_COMP(short1, >=) - -DECLOP_1VAR_COMP(short1, &&) -DECLOP_1VAR_COMP(short1, ||) - -DECLOP_1VAR_1IN_1OUT(short1, ~) -DECLOP_1VAR_1IN_BOOLOUT(short1, !) - -DECLOP_1VAR_SCALE_PRODUCT(short1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(short1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(short1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(short1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(short1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(short1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(short1, float) -DECLOP_1VAR_SCALE_PRODUCT(short1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(short1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(short1, double) -DECLOP_1VAR_SCALE_PRODUCT(short1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(short1, signed long long) - -// SIGNED SHORT2 - -DECLOP_2VAR_2IN_1OUT(short2, +) -DECLOP_2VAR_2IN_1OUT(short2, -) -DECLOP_2VAR_2IN_1OUT(short2, *) -DECLOP_2VAR_2IN_1OUT(short2, /) -DECLOP_2VAR_2IN_1OUT(short2, %) -DECLOP_2VAR_2IN_1OUT(short2, &) -DECLOP_2VAR_2IN_1OUT(short2, |) -DECLOP_2VAR_2IN_1OUT(short2, ^) -DECLOP_2VAR_2IN_1OUT(short2, <<) -DECLOP_2VAR_2IN_1OUT(short2, >>) - -DECLOP_2VAR_ASSIGN(short2, +=) -DECLOP_2VAR_ASSIGN(short2, -=) -DECLOP_2VAR_ASSIGN(short2, *=) -DECLOP_2VAR_ASSIGN(short2, /=) -DECLOP_2VAR_ASSIGN(short2, %=) -DECLOP_2VAR_ASSIGN(short2, &=) -DECLOP_2VAR_ASSIGN(short2, |=) -DECLOP_2VAR_ASSIGN(short2, ^=) -DECLOP_2VAR_ASSIGN(short2, <<=) -DECLOP_2VAR_ASSIGN(short2, >>=) - -DECLOP_2VAR_PREOP(short2, ++) -DECLOP_2VAR_PREOP(short2, --) - -DECLOP_2VAR_POSTOP(short2, ++) -DECLOP_2VAR_POSTOP(short2, --) - -DECLOP_2VAR_COMP(short2, ==) -DECLOP_2VAR_COMP(short2, !=) -DECLOP_2VAR_COMP(short2, <) -DECLOP_2VAR_COMP(short2, >) -DECLOP_2VAR_COMP(short2, <=) -DECLOP_2VAR_COMP(short2, >=) - -DECLOP_2VAR_COMP(short2, &&) -DECLOP_2VAR_COMP(short2, ||) - -DECLOP_2VAR_1IN_1OUT(short2, ~) -DECLOP_2VAR_1IN_BOOLOUT(short2, !) - -DECLOP_2VAR_SCALE_PRODUCT(short2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(short2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(short2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(short2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(short2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(short2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(short2, float) -DECLOP_2VAR_SCALE_PRODUCT(short2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(short2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(short2, double) -DECLOP_2VAR_SCALE_PRODUCT(short2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(short2, signed long long) - -// SIGNED SHORT3 - -DECLOP_3VAR_2IN_1OUT(short3, +) -DECLOP_3VAR_2IN_1OUT(short3, -) -DECLOP_3VAR_2IN_1OUT(short3, *) -DECLOP_3VAR_2IN_1OUT(short3, /) -DECLOP_3VAR_2IN_1OUT(short3, %) -DECLOP_3VAR_2IN_1OUT(short3, &) -DECLOP_3VAR_2IN_1OUT(short3, |) -DECLOP_3VAR_2IN_1OUT(short3, ^) -DECLOP_3VAR_2IN_1OUT(short3, <<) -DECLOP_3VAR_2IN_1OUT(short3, >>) - -DECLOP_3VAR_ASSIGN(short3, +=) -DECLOP_3VAR_ASSIGN(short3, -=) -DECLOP_3VAR_ASSIGN(short3, *=) -DECLOP_3VAR_ASSIGN(short3, /=) -DECLOP_3VAR_ASSIGN(short3, %=) -DECLOP_3VAR_ASSIGN(short3, &=) -DECLOP_3VAR_ASSIGN(short3, |=) -DECLOP_3VAR_ASSIGN(short3, ^=) -DECLOP_3VAR_ASSIGN(short3, <<=) -DECLOP_3VAR_ASSIGN(short3, >>=) - -DECLOP_3VAR_PREOP(short3, ++) -DECLOP_3VAR_PREOP(short3, --) - -DECLOP_3VAR_POSTOP(short3, ++) -DECLOP_3VAR_POSTOP(short3, --) - -DECLOP_3VAR_COMP(short3, ==) -DECLOP_3VAR_COMP(short3, !=) -DECLOP_3VAR_COMP(short3, <) -DECLOP_3VAR_COMP(short3, >) -DECLOP_3VAR_COMP(short3, <=) -DECLOP_3VAR_COMP(short3, >=) - -DECLOP_3VAR_COMP(short3, &&) -DECLOP_3VAR_COMP(short3, ||) - -DECLOP_3VAR_1IN_1OUT(short3, ~) -DECLOP_3VAR_1IN_BOOLOUT(short3, !) - -DECLOP_3VAR_SCALE_PRODUCT(short3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(short3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(short3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(short3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(short3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(short3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(short3, float) -DECLOP_3VAR_SCALE_PRODUCT(short3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(short3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(short3, double) -DECLOP_3VAR_SCALE_PRODUCT(short3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(short3, signed long long) - -// SIGNED SHORT4 - -DECLOP_4VAR_2IN_1OUT(short4, +) -DECLOP_4VAR_2IN_1OUT(short4, -) -DECLOP_4VAR_2IN_1OUT(short4, *) -DECLOP_4VAR_2IN_1OUT(short4, /) -DECLOP_4VAR_2IN_1OUT(short4, %) -DECLOP_4VAR_2IN_1OUT(short4, &) -DECLOP_4VAR_2IN_1OUT(short4, |) -DECLOP_4VAR_2IN_1OUT(short4, ^) -DECLOP_4VAR_2IN_1OUT(short4, <<) -DECLOP_4VAR_2IN_1OUT(short4, >>) - -DECLOP_4VAR_ASSIGN(short4, +=) -DECLOP_4VAR_ASSIGN(short4, -=) -DECLOP_4VAR_ASSIGN(short4, *=) -DECLOP_4VAR_ASSIGN(short4, /=) -DECLOP_4VAR_ASSIGN(short4, %=) -DECLOP_4VAR_ASSIGN(short4, &=) -DECLOP_4VAR_ASSIGN(short4, |=) -DECLOP_4VAR_ASSIGN(short4, ^=) -DECLOP_4VAR_ASSIGN(short4, <<=) -DECLOP_4VAR_ASSIGN(short4, >>=) - -DECLOP_4VAR_PREOP(short4, ++) -DECLOP_4VAR_PREOP(short4, --) - -DECLOP_4VAR_POSTOP(short4, ++) -DECLOP_4VAR_POSTOP(short4, --) - -DECLOP_4VAR_COMP(short4, ==) -DECLOP_4VAR_COMP(short4, !=) -DECLOP_4VAR_COMP(short4, <) -DECLOP_4VAR_COMP(short4, >) -DECLOP_4VAR_COMP(short4, <=) -DECLOP_4VAR_COMP(short4, >=) - -DECLOP_4VAR_COMP(short4, &&) -DECLOP_4VAR_COMP(short4, ||) - -DECLOP_4VAR_1IN_1OUT(short4, ~) -DECLOP_4VAR_1IN_BOOLOUT(short4, !) - -DECLOP_4VAR_SCALE_PRODUCT(short4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(short4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(short4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(short4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(short4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(short4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(short4, float) -DECLOP_4VAR_SCALE_PRODUCT(short4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(short4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(short4, double) -DECLOP_4VAR_SCALE_PRODUCT(short4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(short4, signed long long) - -// UNSIGNED INT1 - -DECLOP_1VAR_2IN_1OUT(uint1, +) -DECLOP_1VAR_2IN_1OUT(uint1, -) -DECLOP_1VAR_2IN_1OUT(uint1, *) -DECLOP_1VAR_2IN_1OUT(uint1, /) -DECLOP_1VAR_2IN_1OUT(uint1, %) -DECLOP_1VAR_2IN_1OUT(uint1, &) -DECLOP_1VAR_2IN_1OUT(uint1, |) -DECLOP_1VAR_2IN_1OUT(uint1, ^) -DECLOP_1VAR_2IN_1OUT(uint1, <<) -DECLOP_1VAR_2IN_1OUT(uint1, >>) - - -DECLOP_1VAR_ASSIGN(uint1, +=) -DECLOP_1VAR_ASSIGN(uint1, -=) -DECLOP_1VAR_ASSIGN(uint1, *=) -DECLOP_1VAR_ASSIGN(uint1, /=) -DECLOP_1VAR_ASSIGN(uint1, %=) -DECLOP_1VAR_ASSIGN(uint1, &=) -DECLOP_1VAR_ASSIGN(uint1, |=) -DECLOP_1VAR_ASSIGN(uint1, ^=) -DECLOP_1VAR_ASSIGN(uint1, <<=) -DECLOP_1VAR_ASSIGN(uint1, >>=) - -DECLOP_1VAR_PREOP(uint1, ++) -DECLOP_1VAR_PREOP(uint1, --) - -DECLOP_1VAR_POSTOP(uint1, ++) -DECLOP_1VAR_POSTOP(uint1, --) - -DECLOP_1VAR_COMP(uint1, ==) -DECLOP_1VAR_COMP(uint1, !=) -DECLOP_1VAR_COMP(uint1, <) -DECLOP_1VAR_COMP(uint1, >) -DECLOP_1VAR_COMP(uint1, <=) -DECLOP_1VAR_COMP(uint1, >=) - -DECLOP_1VAR_COMP(uint1, &&) -DECLOP_1VAR_COMP(uint1, ||) - -DECLOP_1VAR_1IN_1OUT(uint1, ~) -DECLOP_1VAR_1IN_BOOLOUT(uint1, !) - -DECLOP_1VAR_SCALE_PRODUCT(uint1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(uint1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(uint1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(uint1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(uint1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(uint1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(uint1, float) -DECLOP_1VAR_SCALE_PRODUCT(uint1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(uint1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(uint1, double) -DECLOP_1VAR_SCALE_PRODUCT(uint1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(uint1, signed long long) - -// UNSIGNED INT2 - -DECLOP_2VAR_2IN_1OUT(uint2, +) -DECLOP_2VAR_2IN_1OUT(uint2, -) -DECLOP_2VAR_2IN_1OUT(uint2, *) -DECLOP_2VAR_2IN_1OUT(uint2, /) -DECLOP_2VAR_2IN_1OUT(uint2, %) -DECLOP_2VAR_2IN_1OUT(uint2, &) -DECLOP_2VAR_2IN_1OUT(uint2, |) -DECLOP_2VAR_2IN_1OUT(uint2, ^) -DECLOP_2VAR_2IN_1OUT(uint2, <<) -DECLOP_2VAR_2IN_1OUT(uint2, >>) - -DECLOP_2VAR_ASSIGN(uint2, +=) -DECLOP_2VAR_ASSIGN(uint2, -=) -DECLOP_2VAR_ASSIGN(uint2, *=) -DECLOP_2VAR_ASSIGN(uint2, /=) -DECLOP_2VAR_ASSIGN(uint2, %=) -DECLOP_2VAR_ASSIGN(uint2, &=) -DECLOP_2VAR_ASSIGN(uint2, |=) -DECLOP_2VAR_ASSIGN(uint2, ^=) -DECLOP_2VAR_ASSIGN(uint2, <<=) -DECLOP_2VAR_ASSIGN(uint2, >>=) - -DECLOP_2VAR_PREOP(uint2, ++) -DECLOP_2VAR_PREOP(uint2, --) - -DECLOP_2VAR_POSTOP(uint2, ++) -DECLOP_2VAR_POSTOP(uint2, --) - -DECLOP_2VAR_COMP(uint2, ==) -DECLOP_2VAR_COMP(uint2, !=) -DECLOP_2VAR_COMP(uint2, <) -DECLOP_2VAR_COMP(uint2, >) -DECLOP_2VAR_COMP(uint2, <=) -DECLOP_2VAR_COMP(uint2, >=) - -DECLOP_2VAR_COMP(uint2, &&) -DECLOP_2VAR_COMP(uint2, ||) - -DECLOP_2VAR_1IN_1OUT(uint2, ~) -DECLOP_2VAR_1IN_BOOLOUT(uint2, !) - -DECLOP_2VAR_SCALE_PRODUCT(uint2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(uint2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(uint2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(uint2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(uint2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(uint2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(uint2, float) -DECLOP_2VAR_SCALE_PRODUCT(uint2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(uint2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(uint2, double) -DECLOP_2VAR_SCALE_PRODUCT(uint2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(uint2, signed long long) - -// UNSIGNED INT3 - -DECLOP_3VAR_2IN_1OUT(uint3, +) -DECLOP_3VAR_2IN_1OUT(uint3, -) -DECLOP_3VAR_2IN_1OUT(uint3, *) -DECLOP_3VAR_2IN_1OUT(uint3, /) -DECLOP_3VAR_2IN_1OUT(uint3, %) -DECLOP_3VAR_2IN_1OUT(uint3, &) -DECLOP_3VAR_2IN_1OUT(uint3, |) -DECLOP_3VAR_2IN_1OUT(uint3, ^) -DECLOP_3VAR_2IN_1OUT(uint3, <<) -DECLOP_3VAR_2IN_1OUT(uint3, >>) - -DECLOP_3VAR_ASSIGN(uint3, +=) -DECLOP_3VAR_ASSIGN(uint3, -=) -DECLOP_3VAR_ASSIGN(uint3, *=) -DECLOP_3VAR_ASSIGN(uint3, /=) -DECLOP_3VAR_ASSIGN(uint3, %=) -DECLOP_3VAR_ASSIGN(uint3, &=) -DECLOP_3VAR_ASSIGN(uint3, |=) -DECLOP_3VAR_ASSIGN(uint3, ^=) -DECLOP_3VAR_ASSIGN(uint3, <<=) -DECLOP_3VAR_ASSIGN(uint3, >>=) - -DECLOP_3VAR_PREOP(uint3, ++) -DECLOP_3VAR_PREOP(uint3, --) - -DECLOP_3VAR_POSTOP(uint3, ++) -DECLOP_3VAR_POSTOP(uint3, --) - -DECLOP_3VAR_COMP(uint3, ==) -DECLOP_3VAR_COMP(uint3, !=) -DECLOP_3VAR_COMP(uint3, <) -DECLOP_3VAR_COMP(uint3, >) -DECLOP_3VAR_COMP(uint3, <=) -DECLOP_3VAR_COMP(uint3, >=) - -DECLOP_3VAR_COMP(uint3, &&) -DECLOP_3VAR_COMP(uint3, ||) - -DECLOP_3VAR_1IN_1OUT(uint3, ~) -DECLOP_3VAR_1IN_BOOLOUT(uint3, !) - -DECLOP_3VAR_SCALE_PRODUCT(uint3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(uint3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(uint3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(uint3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(uint3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(uint3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(uint3, float) -DECLOP_3VAR_SCALE_PRODUCT(uint3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(uint3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(uint3, double) -DECLOP_3VAR_SCALE_PRODUCT(uint3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(uint3, signed long long) - -// UNSIGNED INT4 - -DECLOP_4VAR_2IN_1OUT(uint4, +) -DECLOP_4VAR_2IN_1OUT(uint4, -) -DECLOP_4VAR_2IN_1OUT(uint4, *) -DECLOP_4VAR_2IN_1OUT(uint4, /) -DECLOP_4VAR_2IN_1OUT(uint4, %) -DECLOP_4VAR_2IN_1OUT(uint4, &) -DECLOP_4VAR_2IN_1OUT(uint4, |) -DECLOP_4VAR_2IN_1OUT(uint4, ^) -DECLOP_4VAR_2IN_1OUT(uint4, <<) -DECLOP_4VAR_2IN_1OUT(uint4, >>) - -DECLOP_4VAR_ASSIGN(uint4, +=) -DECLOP_4VAR_ASSIGN(uint4, -=) -DECLOP_4VAR_ASSIGN(uint4, *=) -DECLOP_4VAR_ASSIGN(uint4, /=) -DECLOP_4VAR_ASSIGN(uint4, %=) -DECLOP_4VAR_ASSIGN(uint4, &=) -DECLOP_4VAR_ASSIGN(uint4, |=) -DECLOP_4VAR_ASSIGN(uint4, ^=) -DECLOP_4VAR_ASSIGN(uint4, <<=) -DECLOP_4VAR_ASSIGN(uint4, >>=) - -DECLOP_4VAR_PREOP(uint4, ++) -DECLOP_4VAR_PREOP(uint4, --) - -DECLOP_4VAR_POSTOP(uint4, ++) -DECLOP_4VAR_POSTOP(uint4, --) - -DECLOP_4VAR_COMP(uint4, ==) -DECLOP_4VAR_COMP(uint4, !=) -DECLOP_4VAR_COMP(uint4, <) -DECLOP_4VAR_COMP(uint4, >) -DECLOP_4VAR_COMP(uint4, <=) -DECLOP_4VAR_COMP(uint4, >=) - -DECLOP_4VAR_COMP(uint4, &&) -DECLOP_4VAR_COMP(uint4, ||) - -DECLOP_4VAR_1IN_1OUT(uint4, ~) -DECLOP_4VAR_1IN_BOOLOUT(uint4, !) - -DECLOP_4VAR_SCALE_PRODUCT(uint4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(uint4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(uint4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(uint4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(uint4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(uint4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(uint4, float) -DECLOP_4VAR_SCALE_PRODUCT(uint4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(uint4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(uint4, double) -DECLOP_4VAR_SCALE_PRODUCT(uint4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(uint4, signed long long) - -// SIGNED INT1 - -DECLOP_1VAR_2IN_1OUT(int1, +) -DECLOP_1VAR_2IN_1OUT(int1, -) -DECLOP_1VAR_2IN_1OUT(int1, *) -DECLOP_1VAR_2IN_1OUT(int1, /) -DECLOP_1VAR_2IN_1OUT(int1, %) -DECLOP_1VAR_2IN_1OUT(int1, &) -DECLOP_1VAR_2IN_1OUT(int1, |) -DECLOP_1VAR_2IN_1OUT(int1, ^) -DECLOP_1VAR_2IN_1OUT(int1, <<) -DECLOP_1VAR_2IN_1OUT(int1, >>) - - -DECLOP_1VAR_ASSIGN(int1, +=) -DECLOP_1VAR_ASSIGN(int1, -=) -DECLOP_1VAR_ASSIGN(int1, *=) -DECLOP_1VAR_ASSIGN(int1, /=) -DECLOP_1VAR_ASSIGN(int1, %=) -DECLOP_1VAR_ASSIGN(int1, &=) -DECLOP_1VAR_ASSIGN(int1, |=) -DECLOP_1VAR_ASSIGN(int1, ^=) -DECLOP_1VAR_ASSIGN(int1, <<=) -DECLOP_1VAR_ASSIGN(int1, >>=) - -DECLOP_1VAR_PREOP(int1, ++) -DECLOP_1VAR_PREOP(int1, --) - -DECLOP_1VAR_POSTOP(int1, ++) -DECLOP_1VAR_POSTOP(int1, --) - -DECLOP_1VAR_COMP(int1, ==) -DECLOP_1VAR_COMP(int1, !=) -DECLOP_1VAR_COMP(int1, <) -DECLOP_1VAR_COMP(int1, >) -DECLOP_1VAR_COMP(int1, <=) -DECLOP_1VAR_COMP(int1, >=) - -DECLOP_1VAR_COMP(int1, &&) -DECLOP_1VAR_COMP(int1, ||) - -DECLOP_1VAR_1IN_1OUT(int1, ~) -DECLOP_1VAR_1IN_BOOLOUT(int1, !) - -DECLOP_1VAR_SCALE_PRODUCT(int1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(int1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(int1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(int1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(int1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(int1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(int1, float) -DECLOP_1VAR_SCALE_PRODUCT(int1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(int1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(int1, double) -DECLOP_1VAR_SCALE_PRODUCT(int1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(int1, signed long long) - -// SIGNED INT2 - -DECLOP_2VAR_2IN_1OUT(int2, +) -DECLOP_2VAR_2IN_1OUT(int2, -) -DECLOP_2VAR_2IN_1OUT(int2, *) -DECLOP_2VAR_2IN_1OUT(int2, /) -DECLOP_2VAR_2IN_1OUT(int2, %) -DECLOP_2VAR_2IN_1OUT(int2, &) -DECLOP_2VAR_2IN_1OUT(int2, |) -DECLOP_2VAR_2IN_1OUT(int2, ^) -DECLOP_2VAR_2IN_1OUT(int2, <<) -DECLOP_2VAR_2IN_1OUT(int2, >>) - -DECLOP_2VAR_ASSIGN(int2, +=) -DECLOP_2VAR_ASSIGN(int2, -=) -DECLOP_2VAR_ASSIGN(int2, *=) -DECLOP_2VAR_ASSIGN(int2, /=) -DECLOP_2VAR_ASSIGN(int2, %=) -DECLOP_2VAR_ASSIGN(int2, &=) -DECLOP_2VAR_ASSIGN(int2, |=) -DECLOP_2VAR_ASSIGN(int2, ^=) -DECLOP_2VAR_ASSIGN(int2, <<=) -DECLOP_2VAR_ASSIGN(int2, >>=) - -DECLOP_2VAR_PREOP(int2, ++) -DECLOP_2VAR_PREOP(int2, --) - -DECLOP_2VAR_POSTOP(int2, ++) -DECLOP_2VAR_POSTOP(int2, --) - -DECLOP_2VAR_COMP(int2, ==) -DECLOP_2VAR_COMP(int2, !=) -DECLOP_2VAR_COMP(int2, <) -DECLOP_2VAR_COMP(int2, >) -DECLOP_2VAR_COMP(int2, <=) -DECLOP_2VAR_COMP(int2, >=) - -DECLOP_2VAR_COMP(int2, &&) -DECLOP_2VAR_COMP(int2, ||) - -DECLOP_2VAR_1IN_1OUT(int2, ~) -DECLOP_2VAR_1IN_BOOLOUT(int2, !) - -DECLOP_2VAR_SCALE_PRODUCT(int2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(int2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(int2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(int2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(int2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(int2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(int2, float) -DECLOP_2VAR_SCALE_PRODUCT(int2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(int2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(int2, double) -DECLOP_2VAR_SCALE_PRODUCT(int2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(int2, signed long long) - -// SIGNED INT3 - -DECLOP_3VAR_2IN_1OUT(int3, +) -DECLOP_3VAR_2IN_1OUT(int3, -) -DECLOP_3VAR_2IN_1OUT(int3, *) -DECLOP_3VAR_2IN_1OUT(int3, /) -DECLOP_3VAR_2IN_1OUT(int3, %) -DECLOP_3VAR_2IN_1OUT(int3, &) -DECLOP_3VAR_2IN_1OUT(int3, |) -DECLOP_3VAR_2IN_1OUT(int3, ^) -DECLOP_3VAR_2IN_1OUT(int3, <<) -DECLOP_3VAR_2IN_1OUT(int3, >>) - -DECLOP_3VAR_ASSIGN(int3, +=) -DECLOP_3VAR_ASSIGN(int3, -=) -DECLOP_3VAR_ASSIGN(int3, *=) -DECLOP_3VAR_ASSIGN(int3, /=) -DECLOP_3VAR_ASSIGN(int3, %=) -DECLOP_3VAR_ASSIGN(int3, &=) -DECLOP_3VAR_ASSIGN(int3, |=) -DECLOP_3VAR_ASSIGN(int3, ^=) -DECLOP_3VAR_ASSIGN(int3, <<=) -DECLOP_3VAR_ASSIGN(int3, >>=) - -DECLOP_3VAR_PREOP(int3, ++) -DECLOP_3VAR_PREOP(int3, --) - -DECLOP_3VAR_POSTOP(int3, ++) -DECLOP_3VAR_POSTOP(int3, --) - -DECLOP_3VAR_COMP(int3, ==) -DECLOP_3VAR_COMP(int3, !=) -DECLOP_3VAR_COMP(int3, <) -DECLOP_3VAR_COMP(int3, >) -DECLOP_3VAR_COMP(int3, <=) -DECLOP_3VAR_COMP(int3, >=) - -DECLOP_3VAR_COMP(int3, &&) -DECLOP_3VAR_COMP(int3, ||) - -DECLOP_3VAR_1IN_1OUT(int3, ~) -DECLOP_3VAR_1IN_BOOLOUT(int3, !) - -DECLOP_3VAR_SCALE_PRODUCT(int3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(int3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(int3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(int3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(int3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(int3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(int3, float) -DECLOP_3VAR_SCALE_PRODUCT(int3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(int3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(int3, double) -DECLOP_3VAR_SCALE_PRODUCT(int3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(int3, signed long long) - -// SIGNED INT4 - -DECLOP_4VAR_2IN_1OUT(int4, +) -DECLOP_4VAR_2IN_1OUT(int4, -) -DECLOP_4VAR_2IN_1OUT(int4, *) -DECLOP_4VAR_2IN_1OUT(int4, /) -DECLOP_4VAR_2IN_1OUT(int4, %) -DECLOP_4VAR_2IN_1OUT(int4, &) -DECLOP_4VAR_2IN_1OUT(int4, |) -DECLOP_4VAR_2IN_1OUT(int4, ^) -DECLOP_4VAR_2IN_1OUT(int4, <<) -DECLOP_4VAR_2IN_1OUT(int4, >>) - -DECLOP_4VAR_ASSIGN(int4, +=) -DECLOP_4VAR_ASSIGN(int4, -=) -DECLOP_4VAR_ASSIGN(int4, *=) -DECLOP_4VAR_ASSIGN(int4, /=) -DECLOP_4VAR_ASSIGN(int4, %=) -DECLOP_4VAR_ASSIGN(int4, &=) -DECLOP_4VAR_ASSIGN(int4, |=) -DECLOP_4VAR_ASSIGN(int4, ^=) -DECLOP_4VAR_ASSIGN(int4, <<=) -DECLOP_4VAR_ASSIGN(int4, >>=) - -DECLOP_4VAR_PREOP(int4, ++) -DECLOP_4VAR_PREOP(int4, --) - -DECLOP_4VAR_POSTOP(int4, ++) -DECLOP_4VAR_POSTOP(int4, --) - -DECLOP_4VAR_COMP(int4, ==) -DECLOP_4VAR_COMP(int4, !=) -DECLOP_4VAR_COMP(int4, <) -DECLOP_4VAR_COMP(int4, >) -DECLOP_4VAR_COMP(int4, <=) -DECLOP_4VAR_COMP(int4, >=) - -DECLOP_4VAR_COMP(int4, &&) -DECLOP_4VAR_COMP(int4, ||) - -DECLOP_4VAR_1IN_1OUT(int4, ~) -DECLOP_4VAR_1IN_BOOLOUT(int4, !) - -DECLOP_4VAR_SCALE_PRODUCT(int4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(int4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(int4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(int4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(int4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(int4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(int4, float) -DECLOP_4VAR_SCALE_PRODUCT(int4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(int4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(int4, double) -DECLOP_4VAR_SCALE_PRODUCT(int4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(int4, signed long long) - -// FLOAT1 - -DECLOP_1VAR_2IN_1OUT(float1, +) -DECLOP_1VAR_2IN_1OUT(float1, -) -DECLOP_1VAR_2IN_1OUT(float1, *) -DECLOP_1VAR_2IN_1OUT(float1, /) - -DECLOP_1VAR_ASSIGN(float1, +=) -DECLOP_1VAR_ASSIGN(float1, -=) -DECLOP_1VAR_ASSIGN(float1, *=) -DECLOP_1VAR_ASSIGN(float1, /=) - -DECLOP_1VAR_PREOP(float1, ++) -DECLOP_1VAR_PREOP(float1, --) - -DECLOP_1VAR_POSTOP(float1, ++) -DECLOP_1VAR_POSTOP(float1, --) - -DECLOP_1VAR_COMP(float1, ==) -DECLOP_1VAR_COMP(float1, !=) -DECLOP_1VAR_COMP(float1, <) -DECLOP_1VAR_COMP(float1, >) -DECLOP_1VAR_COMP(float1, <=) -DECLOP_1VAR_COMP(float1, >=) - -DECLOP_1VAR_SCALE_PRODUCT(float1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(float1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(float1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(float1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(float1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(float1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(float1, float) -DECLOP_1VAR_SCALE_PRODUCT(float1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(float1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(float1, double) -DECLOP_1VAR_SCALE_PRODUCT(float1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(float1, signed long long) - -// FLOAT2 - -DECLOP_2VAR_2IN_1OUT(float2, +) -DECLOP_2VAR_2IN_1OUT(float2, -) -DECLOP_2VAR_2IN_1OUT(float2, *) -DECLOP_2VAR_2IN_1OUT(float2, /) - -DECLOP_2VAR_ASSIGN(float2, +=) -DECLOP_2VAR_ASSIGN(float2, -=) -DECLOP_2VAR_ASSIGN(float2, *=) -DECLOP_2VAR_ASSIGN(float2, /=) - -DECLOP_2VAR_PREOP(float2, ++) -DECLOP_2VAR_PREOP(float2, --) - -DECLOP_2VAR_POSTOP(float2, ++) -DECLOP_2VAR_POSTOP(float2, --) - -DECLOP_2VAR_COMP(float2, ==) -DECLOP_2VAR_COMP(float2, !=) -DECLOP_2VAR_COMP(float2, <) -DECLOP_2VAR_COMP(float2, >) -DECLOP_2VAR_COMP(float2, <=) -DECLOP_2VAR_COMP(float2, >=) - -DECLOP_2VAR_SCALE_PRODUCT(float2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(float2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(float2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(float2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(float2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(float2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(float2, float) -DECLOP_2VAR_SCALE_PRODUCT(float2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(float2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(float2, double) -DECLOP_2VAR_SCALE_PRODUCT(float2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(float2, signed long long) - -// FLOAT3 - -DECLOP_3VAR_2IN_1OUT(float3, +) -DECLOP_3VAR_2IN_1OUT(float3, -) -DECLOP_3VAR_2IN_1OUT(float3, *) -DECLOP_3VAR_2IN_1OUT(float3, /) - -DECLOP_3VAR_ASSIGN(float3, +=) -DECLOP_3VAR_ASSIGN(float3, -=) -DECLOP_3VAR_ASSIGN(float3, *=) -DECLOP_3VAR_ASSIGN(float3, /=) - -DECLOP_3VAR_PREOP(float3, ++) -DECLOP_3VAR_PREOP(float3, --) - -DECLOP_3VAR_POSTOP(float3, ++) -DECLOP_3VAR_POSTOP(float3, --) - -DECLOP_3VAR_COMP(float3, ==) -DECLOP_3VAR_COMP(float3, !=) -DECLOP_3VAR_COMP(float3, <) -DECLOP_3VAR_COMP(float3, >) -DECLOP_3VAR_COMP(float3, <=) -DECLOP_3VAR_COMP(float3, >=) - -DECLOP_3VAR_SCALE_PRODUCT(float3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(float3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(float3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(float3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(float3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(float3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(float3, float) -DECLOP_3VAR_SCALE_PRODUCT(float3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(float3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(float3, double) -DECLOP_3VAR_SCALE_PRODUCT(float3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(float3, signed long long) - -// FLOAT4 - -DECLOP_4VAR_2IN_1OUT(float4, +) -DECLOP_4VAR_2IN_1OUT(float4, -) -DECLOP_4VAR_2IN_1OUT(float4, *) -DECLOP_4VAR_2IN_1OUT(float4, /) - -DECLOP_4VAR_ASSIGN(float4, +=) -DECLOP_4VAR_ASSIGN(float4, -=) -DECLOP_4VAR_ASSIGN(float4, *=) -DECLOP_4VAR_ASSIGN(float4, /=) - -DECLOP_4VAR_PREOP(float4, ++) -DECLOP_4VAR_PREOP(float4, --) - -DECLOP_4VAR_POSTOP(float4, ++) -DECLOP_4VAR_POSTOP(float4, --) - -DECLOP_4VAR_COMP(float4, ==) -DECLOP_4VAR_COMP(float4, !=) -DECLOP_4VAR_COMP(float4, <) -DECLOP_4VAR_COMP(float4, >) -DECLOP_4VAR_COMP(float4, <=) -DECLOP_4VAR_COMP(float4, >=) - -DECLOP_4VAR_SCALE_PRODUCT(float4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(float4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(float4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(float4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(float4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(float4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(float4, float) -DECLOP_4VAR_SCALE_PRODUCT(float4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(float4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(float4, double) -DECLOP_4VAR_SCALE_PRODUCT(float4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(float4, signed long long) - -// DOUBLE1 - -DECLOP_1VAR_2IN_1OUT(double1, +) -DECLOP_1VAR_2IN_1OUT(double1, -) -DECLOP_1VAR_2IN_1OUT(double1, *) -DECLOP_1VAR_2IN_1OUT(double1, /) - -DECLOP_1VAR_ASSIGN(double1, +=) -DECLOP_1VAR_ASSIGN(double1, -=) -DECLOP_1VAR_ASSIGN(double1, *=) -DECLOP_1VAR_ASSIGN(double1, /=) - -DECLOP_1VAR_PREOP(double1, ++) -DECLOP_1VAR_PREOP(double1, --) - -DECLOP_1VAR_POSTOP(double1, ++) -DECLOP_1VAR_POSTOP(double1, --) - -DECLOP_1VAR_COMP(double1, ==) -DECLOP_1VAR_COMP(double1, !=) -DECLOP_1VAR_COMP(double1, <) -DECLOP_1VAR_COMP(double1, >) -DECLOP_1VAR_COMP(double1, <=) -DECLOP_1VAR_COMP(double1, >=) - -DECLOP_1VAR_SCALE_PRODUCT(double1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(double1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(double1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(double1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(double1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(double1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(double1, float) -DECLOP_1VAR_SCALE_PRODUCT(double1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(double1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(double1, double) -DECLOP_1VAR_SCALE_PRODUCT(double1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(double1, signed long long) - -// DOUBLE2 - -DECLOP_2VAR_2IN_1OUT(double2, +) -DECLOP_2VAR_2IN_1OUT(double2, -) -DECLOP_2VAR_2IN_1OUT(double2, *) -DECLOP_2VAR_2IN_1OUT(double2, /) - -DECLOP_2VAR_ASSIGN(double2, +=) -DECLOP_2VAR_ASSIGN(double2, -=) -DECLOP_2VAR_ASSIGN(double2, *=) -DECLOP_2VAR_ASSIGN(double2, /=) - -DECLOP_2VAR_PREOP(double2, ++) -DECLOP_2VAR_PREOP(double2, --) - -DECLOP_2VAR_POSTOP(double2, ++) -DECLOP_2VAR_POSTOP(double2, --) - -DECLOP_2VAR_COMP(double2, ==) -DECLOP_2VAR_COMP(double2, !=) -DECLOP_2VAR_COMP(double2, <) -DECLOP_2VAR_COMP(double2, >) -DECLOP_2VAR_COMP(double2, <=) -DECLOP_2VAR_COMP(double2, >=) - -DECLOP_2VAR_SCALE_PRODUCT(double2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(double2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(double2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(double2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(double2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(double2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(double2, float) -DECLOP_2VAR_SCALE_PRODUCT(double2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(double2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(double2, double) -DECLOP_2VAR_SCALE_PRODUCT(double2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(double2, signed long long) - -// DOUBLE3 - -DECLOP_3VAR_2IN_1OUT(double3, +) -DECLOP_3VAR_2IN_1OUT(double3, -) -DECLOP_3VAR_2IN_1OUT(double3, *) -DECLOP_3VAR_2IN_1OUT(double3, /) - -DECLOP_3VAR_ASSIGN(double3, +=) -DECLOP_3VAR_ASSIGN(double3, -=) -DECLOP_3VAR_ASSIGN(double3, *=) -DECLOP_3VAR_ASSIGN(double3, /=) - -DECLOP_3VAR_PREOP(double3, ++) -DECLOP_3VAR_PREOP(double3, --) - -DECLOP_3VAR_POSTOP(double3, ++) -DECLOP_3VAR_POSTOP(double3, --) - -DECLOP_3VAR_COMP(double3, ==) -DECLOP_3VAR_COMP(double3, !=) -DECLOP_3VAR_COMP(double3, <) -DECLOP_3VAR_COMP(double3, >) -DECLOP_3VAR_COMP(double3, <=) -DECLOP_3VAR_COMP(double3, >=) - -DECLOP_3VAR_SCALE_PRODUCT(double3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(double3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(double3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(double3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(double3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(double3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(double3, float) -DECLOP_3VAR_SCALE_PRODUCT(double3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(double3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(double3, double) -DECLOP_3VAR_SCALE_PRODUCT(double3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(double3, signed long long) - -// DOUBLE4 - -DECLOP_4VAR_2IN_1OUT(double4, +) -DECLOP_4VAR_2IN_1OUT(double4, -) -DECLOP_4VAR_2IN_1OUT(double4, *) -DECLOP_4VAR_2IN_1OUT(double4, /) - -DECLOP_4VAR_ASSIGN(double4, +=) -DECLOP_4VAR_ASSIGN(double4, -=) -DECLOP_4VAR_ASSIGN(double4, *=) -DECLOP_4VAR_ASSIGN(double4, /=) - -DECLOP_4VAR_PREOP(double4, ++) -DECLOP_4VAR_PREOP(double4, --) - -DECLOP_4VAR_POSTOP(double4, ++) -DECLOP_4VAR_POSTOP(double4, --) - -DECLOP_4VAR_COMP(double4, ==) -DECLOP_4VAR_COMP(double4, !=) -DECLOP_4VAR_COMP(double4, <) -DECLOP_4VAR_COMP(double4, >) -DECLOP_4VAR_COMP(double4, <=) -DECLOP_4VAR_COMP(double4, >=) - -DECLOP_4VAR_SCALE_PRODUCT(double4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(double4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(double4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(double4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(double4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(double4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(double4, float) -DECLOP_4VAR_SCALE_PRODUCT(double4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(double4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(double4, double) -DECLOP_4VAR_SCALE_PRODUCT(double4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(double4, signed long long) - -// UNSIGNED LONG1 - -DECLOP_1VAR_2IN_1OUT(ulong1, +) -DECLOP_1VAR_2IN_1OUT(ulong1, -) -DECLOP_1VAR_2IN_1OUT(ulong1, *) -DECLOP_1VAR_2IN_1OUT(ulong1, /) -DECLOP_1VAR_2IN_1OUT(ulong1, %) -DECLOP_1VAR_2IN_1OUT(ulong1, &) -DECLOP_1VAR_2IN_1OUT(ulong1, |) -DECLOP_1VAR_2IN_1OUT(ulong1, ^) -DECLOP_1VAR_2IN_1OUT(ulong1, <<) -DECLOP_1VAR_2IN_1OUT(ulong1, >>) - - -DECLOP_1VAR_ASSIGN(ulong1, +=) -DECLOP_1VAR_ASSIGN(ulong1, -=) -DECLOP_1VAR_ASSIGN(ulong1, *=) -DECLOP_1VAR_ASSIGN(ulong1, /=) -DECLOP_1VAR_ASSIGN(ulong1, %=) -DECLOP_1VAR_ASSIGN(ulong1, &=) -DECLOP_1VAR_ASSIGN(ulong1, |=) -DECLOP_1VAR_ASSIGN(ulong1, ^=) -DECLOP_1VAR_ASSIGN(ulong1, <<=) -DECLOP_1VAR_ASSIGN(ulong1, >>=) - -DECLOP_1VAR_PREOP(ulong1, ++) -DECLOP_1VAR_PREOP(ulong1, --) - -DECLOP_1VAR_POSTOP(ulong1, ++) -DECLOP_1VAR_POSTOP(ulong1, --) - -DECLOP_1VAR_COMP(ulong1, ==) -DECLOP_1VAR_COMP(ulong1, !=) -DECLOP_1VAR_COMP(ulong1, <) -DECLOP_1VAR_COMP(ulong1, >) -DECLOP_1VAR_COMP(ulong1, <=) -DECLOP_1VAR_COMP(ulong1, >=) - -DECLOP_1VAR_COMP(ulong1, &&) -DECLOP_1VAR_COMP(ulong1, ||) - -DECLOP_1VAR_1IN_1OUT(ulong1, ~) -DECLOP_1VAR_1IN_BOOLOUT(ulong1, !) - -DECLOP_1VAR_SCALE_PRODUCT(ulong1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, float) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, double) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(ulong1, signed long long) - -// UNSIGNED LONG2 - -DECLOP_2VAR_2IN_1OUT(ulong2, +) -DECLOP_2VAR_2IN_1OUT(ulong2, -) -DECLOP_2VAR_2IN_1OUT(ulong2, *) -DECLOP_2VAR_2IN_1OUT(ulong2, /) -DECLOP_2VAR_2IN_1OUT(ulong2, %) -DECLOP_2VAR_2IN_1OUT(ulong2, &) -DECLOP_2VAR_2IN_1OUT(ulong2, |) -DECLOP_2VAR_2IN_1OUT(ulong2, ^) -DECLOP_2VAR_2IN_1OUT(ulong2, <<) -DECLOP_2VAR_2IN_1OUT(ulong2, >>) - -DECLOP_2VAR_ASSIGN(ulong2, +=) -DECLOP_2VAR_ASSIGN(ulong2, -=) -DECLOP_2VAR_ASSIGN(ulong2, *=) -DECLOP_2VAR_ASSIGN(ulong2, /=) -DECLOP_2VAR_ASSIGN(ulong2, %=) -DECLOP_2VAR_ASSIGN(ulong2, &=) -DECLOP_2VAR_ASSIGN(ulong2, |=) -DECLOP_2VAR_ASSIGN(ulong2, ^=) -DECLOP_2VAR_ASSIGN(ulong2, <<=) -DECLOP_2VAR_ASSIGN(ulong2, >>=) - -DECLOP_2VAR_PREOP(ulong2, ++) -DECLOP_2VAR_PREOP(ulong2, --) - -DECLOP_2VAR_POSTOP(ulong2, ++) -DECLOP_2VAR_POSTOP(ulong2, --) - -DECLOP_2VAR_COMP(ulong2, ==) -DECLOP_2VAR_COMP(ulong2, !=) -DECLOP_2VAR_COMP(ulong2, <) -DECLOP_2VAR_COMP(ulong2, >) -DECLOP_2VAR_COMP(ulong2, <=) -DECLOP_2VAR_COMP(ulong2, >=) - -DECLOP_2VAR_COMP(ulong2, &&) -DECLOP_2VAR_COMP(ulong2, ||) - -DECLOP_2VAR_1IN_1OUT(ulong2, ~) -DECLOP_2VAR_1IN_BOOLOUT(ulong2, !) - -DECLOP_2VAR_SCALE_PRODUCT(ulong2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, float) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, double) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(ulong2, signed long long) - -// UNSIGNED LONG3 - -DECLOP_3VAR_2IN_1OUT(ulong3, +) -DECLOP_3VAR_2IN_1OUT(ulong3, -) -DECLOP_3VAR_2IN_1OUT(ulong3, *) -DECLOP_3VAR_2IN_1OUT(ulong3, /) -DECLOP_3VAR_2IN_1OUT(ulong3, %) -DECLOP_3VAR_2IN_1OUT(ulong3, &) -DECLOP_3VAR_2IN_1OUT(ulong3, |) -DECLOP_3VAR_2IN_1OUT(ulong3, ^) -DECLOP_3VAR_2IN_1OUT(ulong3, <<) -DECLOP_3VAR_2IN_1OUT(ulong3, >>) - -DECLOP_3VAR_ASSIGN(ulong3, +=) -DECLOP_3VAR_ASSIGN(ulong3, -=) -DECLOP_3VAR_ASSIGN(ulong3, *=) -DECLOP_3VAR_ASSIGN(ulong3, /=) -DECLOP_3VAR_ASSIGN(ulong3, %=) -DECLOP_3VAR_ASSIGN(ulong3, &=) -DECLOP_3VAR_ASSIGN(ulong3, |=) -DECLOP_3VAR_ASSIGN(ulong3, ^=) -DECLOP_3VAR_ASSIGN(ulong3, <<=) -DECLOP_3VAR_ASSIGN(ulong3, >>=) - -DECLOP_3VAR_PREOP(ulong3, ++) -DECLOP_3VAR_PREOP(ulong3, --) - -DECLOP_3VAR_POSTOP(ulong3, ++) -DECLOP_3VAR_POSTOP(ulong3, --) - -DECLOP_3VAR_COMP(ulong3, ==) -DECLOP_3VAR_COMP(ulong3, !=) -DECLOP_3VAR_COMP(ulong3, <) -DECLOP_3VAR_COMP(ulong3, >) -DECLOP_3VAR_COMP(ulong3, <=) -DECLOP_3VAR_COMP(ulong3, >=) - -DECLOP_3VAR_COMP(ulong3, &&) -DECLOP_3VAR_COMP(ulong3, ||) - -DECLOP_3VAR_1IN_1OUT(ulong3, ~) -DECLOP_3VAR_1IN_BOOLOUT(ulong3, !) - -DECLOP_3VAR_SCALE_PRODUCT(ulong3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, float) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, double) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(ulong3, signed long long) - -// UNSIGNED LONG4 - -DECLOP_4VAR_2IN_1OUT(ulong4, +) -DECLOP_4VAR_2IN_1OUT(ulong4, -) -DECLOP_4VAR_2IN_1OUT(ulong4, *) -DECLOP_4VAR_2IN_1OUT(ulong4, /) -DECLOP_4VAR_2IN_1OUT(ulong4, %) -DECLOP_4VAR_2IN_1OUT(ulong4, &) -DECLOP_4VAR_2IN_1OUT(ulong4, |) -DECLOP_4VAR_2IN_1OUT(ulong4, ^) -DECLOP_4VAR_2IN_1OUT(ulong4, <<) -DECLOP_4VAR_2IN_1OUT(ulong4, >>) - -DECLOP_4VAR_ASSIGN(ulong4, +=) -DECLOP_4VAR_ASSIGN(ulong4, -=) -DECLOP_4VAR_ASSIGN(ulong4, *=) -DECLOP_4VAR_ASSIGN(ulong4, /=) -DECLOP_4VAR_ASSIGN(ulong4, %=) -DECLOP_4VAR_ASSIGN(ulong4, &=) -DECLOP_4VAR_ASSIGN(ulong4, |=) -DECLOP_4VAR_ASSIGN(ulong4, ^=) -DECLOP_4VAR_ASSIGN(ulong4, <<=) -DECLOP_4VAR_ASSIGN(ulong4, >>=) - -DECLOP_4VAR_PREOP(ulong4, ++) -DECLOP_4VAR_PREOP(ulong4, --) - -DECLOP_4VAR_POSTOP(ulong4, ++) -DECLOP_4VAR_POSTOP(ulong4, --) - -DECLOP_4VAR_COMP(ulong4, ==) -DECLOP_4VAR_COMP(ulong4, !=) -DECLOP_4VAR_COMP(ulong4, <) -DECLOP_4VAR_COMP(ulong4, >) -DECLOP_4VAR_COMP(ulong4, <=) -DECLOP_4VAR_COMP(ulong4, >=) - -DECLOP_4VAR_COMP(ulong4, &&) -DECLOP_4VAR_COMP(ulong4, ||) - -DECLOP_4VAR_1IN_1OUT(ulong4, ~) -DECLOP_4VAR_1IN_BOOLOUT(ulong4, !) - -DECLOP_4VAR_SCALE_PRODUCT(ulong4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, float) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, double) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(ulong4, signed long long) - -// SIGNED LONG1 - -DECLOP_1VAR_2IN_1OUT(long1, +) -DECLOP_1VAR_2IN_1OUT(long1, -) -DECLOP_1VAR_2IN_1OUT(long1, *) -DECLOP_1VAR_2IN_1OUT(long1, /) -DECLOP_1VAR_2IN_1OUT(long1, %) -DECLOP_1VAR_2IN_1OUT(long1, &) -DECLOP_1VAR_2IN_1OUT(long1, |) -DECLOP_1VAR_2IN_1OUT(long1, ^) -DECLOP_1VAR_2IN_1OUT(long1, <<) -DECLOP_1VAR_2IN_1OUT(long1, >>) - - -DECLOP_1VAR_ASSIGN(long1, +=) -DECLOP_1VAR_ASSIGN(long1, -=) -DECLOP_1VAR_ASSIGN(long1, *=) -DECLOP_1VAR_ASSIGN(long1, /=) -DECLOP_1VAR_ASSIGN(long1, %=) -DECLOP_1VAR_ASSIGN(long1, &=) -DECLOP_1VAR_ASSIGN(long1, |=) -DECLOP_1VAR_ASSIGN(long1, ^=) -DECLOP_1VAR_ASSIGN(long1, <<=) -DECLOP_1VAR_ASSIGN(long1, >>=) - -DECLOP_1VAR_PREOP(long1, ++) -DECLOP_1VAR_PREOP(long1, --) - -DECLOP_1VAR_POSTOP(long1, ++) -DECLOP_1VAR_POSTOP(long1, --) - -DECLOP_1VAR_COMP(long1, ==) -DECLOP_1VAR_COMP(long1, !=) -DECLOP_1VAR_COMP(long1, <) -DECLOP_1VAR_COMP(long1, >) -DECLOP_1VAR_COMP(long1, <=) -DECLOP_1VAR_COMP(long1, >=) - -DECLOP_1VAR_COMP(long1, &&) -DECLOP_1VAR_COMP(long1, ||) - -DECLOP_1VAR_1IN_1OUT(long1, ~) -DECLOP_1VAR_1IN_BOOLOUT(long1, !) - -DECLOP_1VAR_SCALE_PRODUCT(long1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(long1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(long1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(long1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(long1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(long1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(long1, float) -DECLOP_1VAR_SCALE_PRODUCT(long1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(long1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(long1, double) -DECLOP_1VAR_SCALE_PRODUCT(long1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(long1, signed long long) - -// SIGNED LONG2 - -DECLOP_2VAR_2IN_1OUT(long2, +) -DECLOP_2VAR_2IN_1OUT(long2, -) -DECLOP_2VAR_2IN_1OUT(long2, *) -DECLOP_2VAR_2IN_1OUT(long2, /) -DECLOP_2VAR_2IN_1OUT(long2, %) -DECLOP_2VAR_2IN_1OUT(long2, &) -DECLOP_2VAR_2IN_1OUT(long2, |) -DECLOP_2VAR_2IN_1OUT(long2, ^) -DECLOP_2VAR_2IN_1OUT(long2, <<) -DECLOP_2VAR_2IN_1OUT(long2, >>) - -DECLOP_2VAR_ASSIGN(long2, +=) -DECLOP_2VAR_ASSIGN(long2, -=) -DECLOP_2VAR_ASSIGN(long2, *=) -DECLOP_2VAR_ASSIGN(long2, /=) -DECLOP_2VAR_ASSIGN(long2, %=) -DECLOP_2VAR_ASSIGN(long2, &=) -DECLOP_2VAR_ASSIGN(long2, |=) -DECLOP_2VAR_ASSIGN(long2, ^=) -DECLOP_2VAR_ASSIGN(long2, <<=) -DECLOP_2VAR_ASSIGN(long2, >>=) - -DECLOP_2VAR_PREOP(long2, ++) -DECLOP_2VAR_PREOP(long2, --) - -DECLOP_2VAR_POSTOP(long2, ++) -DECLOP_2VAR_POSTOP(long2, --) - -DECLOP_2VAR_COMP(long2, ==) -DECLOP_2VAR_COMP(long2, !=) -DECLOP_2VAR_COMP(long2, <) -DECLOP_2VAR_COMP(long2, >) -DECLOP_2VAR_COMP(long2, <=) -DECLOP_2VAR_COMP(long2, >=) - -DECLOP_2VAR_COMP(long2, &&) -DECLOP_2VAR_COMP(long2, ||) - -DECLOP_2VAR_1IN_1OUT(long2, ~) -DECLOP_2VAR_1IN_BOOLOUT(long2, !) - -DECLOP_2VAR_SCALE_PRODUCT(long2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(long2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(long2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(long2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(long2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(long2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(long2, float) -DECLOP_2VAR_SCALE_PRODUCT(long2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(long2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(long2, double) -DECLOP_2VAR_SCALE_PRODUCT(long2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(long2, signed long long) - -// SIGNED LONG3 - -DECLOP_3VAR_2IN_1OUT(long3, +) -DECLOP_3VAR_2IN_1OUT(long3, -) -DECLOP_3VAR_2IN_1OUT(long3, *) -DECLOP_3VAR_2IN_1OUT(long3, /) -DECLOP_3VAR_2IN_1OUT(long3, %) -DECLOP_3VAR_2IN_1OUT(long3, &) -DECLOP_3VAR_2IN_1OUT(long3, |) -DECLOP_3VAR_2IN_1OUT(long3, ^) -DECLOP_3VAR_2IN_1OUT(long3, <<) -DECLOP_3VAR_2IN_1OUT(long3, >>) - -DECLOP_3VAR_ASSIGN(long3, +=) -DECLOP_3VAR_ASSIGN(long3, -=) -DECLOP_3VAR_ASSIGN(long3, *=) -DECLOP_3VAR_ASSIGN(long3, /=) -DECLOP_3VAR_ASSIGN(long3, %=) -DECLOP_3VAR_ASSIGN(long3, &=) -DECLOP_3VAR_ASSIGN(long3, |=) -DECLOP_3VAR_ASSIGN(long3, ^=) -DECLOP_3VAR_ASSIGN(long3, <<=) -DECLOP_3VAR_ASSIGN(long3, >>=) - -DECLOP_3VAR_PREOP(long3, ++) -DECLOP_3VAR_PREOP(long3, --) - -DECLOP_3VAR_POSTOP(long3, ++) -DECLOP_3VAR_POSTOP(long3, --) - -DECLOP_3VAR_COMP(long3, ==) -DECLOP_3VAR_COMP(long3, !=) -DECLOP_3VAR_COMP(long3, <) -DECLOP_3VAR_COMP(long3, >) -DECLOP_3VAR_COMP(long3, <=) -DECLOP_3VAR_COMP(long3, >=) - -DECLOP_3VAR_COMP(long3, &&) -DECLOP_3VAR_COMP(long3, ||) - -DECLOP_3VAR_1IN_1OUT(long3, ~) -DECLOP_3VAR_1IN_BOOLOUT(long3, !) - -DECLOP_3VAR_SCALE_PRODUCT(long3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(long3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(long3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(long3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(long3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(long3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(long3, float) -DECLOP_3VAR_SCALE_PRODUCT(long3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(long3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(long3, double) -DECLOP_3VAR_SCALE_PRODUCT(long3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(long3, signed long long) - -// SIGNED LONG4 - -DECLOP_4VAR_2IN_1OUT(long4, +) -DECLOP_4VAR_2IN_1OUT(long4, -) -DECLOP_4VAR_2IN_1OUT(long4, *) -DECLOP_4VAR_2IN_1OUT(long4, /) -DECLOP_4VAR_2IN_1OUT(long4, %) -DECLOP_4VAR_2IN_1OUT(long4, &) -DECLOP_4VAR_2IN_1OUT(long4, |) -DECLOP_4VAR_2IN_1OUT(long4, ^) -DECLOP_4VAR_2IN_1OUT(long4, <<) -DECLOP_4VAR_2IN_1OUT(long4, >>) - -DECLOP_4VAR_ASSIGN(long4, +=) -DECLOP_4VAR_ASSIGN(long4, -=) -DECLOP_4VAR_ASSIGN(long4, *=) -DECLOP_4VAR_ASSIGN(long4, /=) -DECLOP_4VAR_ASSIGN(long4, %=) -DECLOP_4VAR_ASSIGN(long4, &=) -DECLOP_4VAR_ASSIGN(long4, |=) -DECLOP_4VAR_ASSIGN(long4, ^=) -DECLOP_4VAR_ASSIGN(long4, <<=) -DECLOP_4VAR_ASSIGN(long4, >>=) - -DECLOP_4VAR_PREOP(long4, ++) -DECLOP_4VAR_PREOP(long4, --) - -DECLOP_4VAR_POSTOP(long4, ++) -DECLOP_4VAR_POSTOP(long4, --) - -DECLOP_4VAR_COMP(long4, ==) -DECLOP_4VAR_COMP(long4, !=) -DECLOP_4VAR_COMP(long4, <) -DECLOP_4VAR_COMP(long4, >) -DECLOP_4VAR_COMP(long4, <=) -DECLOP_4VAR_COMP(long4, >=) - -DECLOP_4VAR_COMP(long4, &&) -DECLOP_4VAR_COMP(long4, ||) - -DECLOP_4VAR_1IN_1OUT(long4, ~) -DECLOP_4VAR_1IN_BOOLOUT(long4, !) - -DECLOP_4VAR_SCALE_PRODUCT(long4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(long4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(long4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(long4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(long4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(long4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(long4, float) -DECLOP_4VAR_SCALE_PRODUCT(long4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(long4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(long4, double) -DECLOP_4VAR_SCALE_PRODUCT(long4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(long4, signed long long) - -// UNSIGNED LONGLONG1 - -DECLOP_1VAR_2IN_1OUT(ulonglong1, +) -DECLOP_1VAR_2IN_1OUT(ulonglong1, -) -DECLOP_1VAR_2IN_1OUT(ulonglong1, *) -DECLOP_1VAR_2IN_1OUT(ulonglong1, /) -DECLOP_1VAR_2IN_1OUT(ulonglong1, %) -DECLOP_1VAR_2IN_1OUT(ulonglong1, &) -DECLOP_1VAR_2IN_1OUT(ulonglong1, |) -DECLOP_1VAR_2IN_1OUT(ulonglong1, ^) -DECLOP_1VAR_2IN_1OUT(ulonglong1, <<) -DECLOP_1VAR_2IN_1OUT(ulonglong1, >>) - - -DECLOP_1VAR_ASSIGN(ulonglong1, +=) -DECLOP_1VAR_ASSIGN(ulonglong1, -=) -DECLOP_1VAR_ASSIGN(ulonglong1, *=) -DECLOP_1VAR_ASSIGN(ulonglong1, /=) -DECLOP_1VAR_ASSIGN(ulonglong1, %=) -DECLOP_1VAR_ASSIGN(ulonglong1, &=) -DECLOP_1VAR_ASSIGN(ulonglong1, |=) -DECLOP_1VAR_ASSIGN(ulonglong1, ^=) -DECLOP_1VAR_ASSIGN(ulonglong1, <<=) -DECLOP_1VAR_ASSIGN(ulonglong1, >>=) - -DECLOP_1VAR_PREOP(ulonglong1, ++) -DECLOP_1VAR_PREOP(ulonglong1, --) - -DECLOP_1VAR_POSTOP(ulonglong1, ++) -DECLOP_1VAR_POSTOP(ulonglong1, --) - -DECLOP_1VAR_COMP(ulonglong1, ==) -DECLOP_1VAR_COMP(ulonglong1, !=) -DECLOP_1VAR_COMP(ulonglong1, <) -DECLOP_1VAR_COMP(ulonglong1, >) -DECLOP_1VAR_COMP(ulonglong1, <=) -DECLOP_1VAR_COMP(ulonglong1, >=) - -DECLOP_1VAR_COMP(ulonglong1, &&) -DECLOP_1VAR_COMP(ulonglong1, ||) - -DECLOP_1VAR_1IN_1OUT(ulonglong1, ~) -DECLOP_1VAR_1IN_BOOLOUT(ulonglong1, !) - -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, float) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, double) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(ulonglong1, signed long long) - -// UNSIGNED LONGLONG2 - -DECLOP_2VAR_2IN_1OUT(ulonglong2, +) -DECLOP_2VAR_2IN_1OUT(ulonglong2, -) -DECLOP_2VAR_2IN_1OUT(ulonglong2, *) -DECLOP_2VAR_2IN_1OUT(ulonglong2, /) -DECLOP_2VAR_2IN_1OUT(ulonglong2, %) -DECLOP_2VAR_2IN_1OUT(ulonglong2, &) -DECLOP_2VAR_2IN_1OUT(ulonglong2, |) -DECLOP_2VAR_2IN_1OUT(ulonglong2, ^) -DECLOP_2VAR_2IN_1OUT(ulonglong2, <<) -DECLOP_2VAR_2IN_1OUT(ulonglong2, >>) - -DECLOP_2VAR_ASSIGN(ulonglong2, +=) -DECLOP_2VAR_ASSIGN(ulonglong2, -=) -DECLOP_2VAR_ASSIGN(ulonglong2, *=) -DECLOP_2VAR_ASSIGN(ulonglong2, /=) -DECLOP_2VAR_ASSIGN(ulonglong2, %=) -DECLOP_2VAR_ASSIGN(ulonglong2, &=) -DECLOP_2VAR_ASSIGN(ulonglong2, |=) -DECLOP_2VAR_ASSIGN(ulonglong2, ^=) -DECLOP_2VAR_ASSIGN(ulonglong2, <<=) -DECLOP_2VAR_ASSIGN(ulonglong2, >>=) - -DECLOP_2VAR_PREOP(ulonglong2, ++) -DECLOP_2VAR_PREOP(ulonglong2, --) - -DECLOP_2VAR_POSTOP(ulonglong2, ++) -DECLOP_2VAR_POSTOP(ulonglong2, --) - -DECLOP_2VAR_COMP(ulonglong2, ==) -DECLOP_2VAR_COMP(ulonglong2, !=) -DECLOP_2VAR_COMP(ulonglong2, <) -DECLOP_2VAR_COMP(ulonglong2, >) -DECLOP_2VAR_COMP(ulonglong2, <=) -DECLOP_2VAR_COMP(ulonglong2, >=) - -DECLOP_2VAR_COMP(ulonglong2, &&) -DECLOP_2VAR_COMP(ulonglong2, ||) - -DECLOP_2VAR_1IN_1OUT(ulonglong2, ~) -DECLOP_2VAR_1IN_BOOLOUT(ulonglong2, !) - -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, float) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, double) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(ulonglong2, signed long long) - -// UNSIGNED LONGLONG3 - -DECLOP_3VAR_2IN_1OUT(ulonglong3, +) -DECLOP_3VAR_2IN_1OUT(ulonglong3, -) -DECLOP_3VAR_2IN_1OUT(ulonglong3, *) -DECLOP_3VAR_2IN_1OUT(ulonglong3, /) -DECLOP_3VAR_2IN_1OUT(ulonglong3, %) -DECLOP_3VAR_2IN_1OUT(ulonglong3, &) -DECLOP_3VAR_2IN_1OUT(ulonglong3, |) -DECLOP_3VAR_2IN_1OUT(ulonglong3, ^) -DECLOP_3VAR_2IN_1OUT(ulonglong3, <<) -DECLOP_3VAR_2IN_1OUT(ulonglong3, >>) - -DECLOP_3VAR_ASSIGN(ulonglong3, +=) -DECLOP_3VAR_ASSIGN(ulonglong3, -=) -DECLOP_3VAR_ASSIGN(ulonglong3, *=) -DECLOP_3VAR_ASSIGN(ulonglong3, /=) -DECLOP_3VAR_ASSIGN(ulonglong3, %=) -DECLOP_3VAR_ASSIGN(ulonglong3, &=) -DECLOP_3VAR_ASSIGN(ulonglong3, |=) -DECLOP_3VAR_ASSIGN(ulonglong3, ^=) -DECLOP_3VAR_ASSIGN(ulonglong3, <<=) -DECLOP_3VAR_ASSIGN(ulonglong3, >>=) - -DECLOP_3VAR_PREOP(ulonglong3, ++) -DECLOP_3VAR_PREOP(ulonglong3, --) - -DECLOP_3VAR_POSTOP(ulonglong3, ++) -DECLOP_3VAR_POSTOP(ulonglong3, --) - -DECLOP_3VAR_COMP(ulonglong3, ==) -DECLOP_3VAR_COMP(ulonglong3, !=) -DECLOP_3VAR_COMP(ulonglong3, <) -DECLOP_3VAR_COMP(ulonglong3, >) -DECLOP_3VAR_COMP(ulonglong3, <=) -DECLOP_3VAR_COMP(ulonglong3, >=) - -DECLOP_3VAR_COMP(ulonglong3, &&) -DECLOP_3VAR_COMP(ulonglong3, ||) - -DECLOP_3VAR_1IN_1OUT(ulonglong3, ~) -DECLOP_3VAR_1IN_BOOLOUT(ulonglong3, !) - -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, float) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, double) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(ulonglong3, signed long long) - -// UNSIGNED LONGLONG4 - -DECLOP_4VAR_2IN_1OUT(ulonglong4, +) -DECLOP_4VAR_2IN_1OUT(ulonglong4, -) -DECLOP_4VAR_2IN_1OUT(ulonglong4, *) -DECLOP_4VAR_2IN_1OUT(ulonglong4, /) -DECLOP_4VAR_2IN_1OUT(ulonglong4, %) -DECLOP_4VAR_2IN_1OUT(ulonglong4, &) -DECLOP_4VAR_2IN_1OUT(ulonglong4, |) -DECLOP_4VAR_2IN_1OUT(ulonglong4, ^) -DECLOP_4VAR_2IN_1OUT(ulonglong4, <<) -DECLOP_4VAR_2IN_1OUT(ulonglong4, >>) - -DECLOP_4VAR_ASSIGN(ulonglong4, +=) -DECLOP_4VAR_ASSIGN(ulonglong4, -=) -DECLOP_4VAR_ASSIGN(ulonglong4, *=) -DECLOP_4VAR_ASSIGN(ulonglong4, /=) -DECLOP_4VAR_ASSIGN(ulonglong4, %=) -DECLOP_4VAR_ASSIGN(ulonglong4, &=) -DECLOP_4VAR_ASSIGN(ulonglong4, |=) -DECLOP_4VAR_ASSIGN(ulonglong4, ^=) -DECLOP_4VAR_ASSIGN(ulonglong4, <<=) -DECLOP_4VAR_ASSIGN(ulonglong4, >>=) - -DECLOP_4VAR_PREOP(ulonglong4, ++) -DECLOP_4VAR_PREOP(ulonglong4, --) - -DECLOP_4VAR_POSTOP(ulonglong4, ++) -DECLOP_4VAR_POSTOP(ulonglong4, --) - -DECLOP_4VAR_COMP(ulonglong4, ==) -DECLOP_4VAR_COMP(ulonglong4, !=) -DECLOP_4VAR_COMP(ulonglong4, <) -DECLOP_4VAR_COMP(ulonglong4, >) -DECLOP_4VAR_COMP(ulonglong4, <=) -DECLOP_4VAR_COMP(ulonglong4, >=) - -DECLOP_4VAR_COMP(ulonglong4, &&) -DECLOP_4VAR_COMP(ulonglong4, ||) - -DECLOP_4VAR_1IN_1OUT(ulonglong4, ~) -DECLOP_4VAR_1IN_BOOLOUT(ulonglong4, !) - -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, float) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, double) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(ulonglong4, signed long long) - -// SIGNED LONGLONG1 - -DECLOP_1VAR_2IN_1OUT(longlong1, +) -DECLOP_1VAR_2IN_1OUT(longlong1, -) -DECLOP_1VAR_2IN_1OUT(longlong1, *) -DECLOP_1VAR_2IN_1OUT(longlong1, /) -DECLOP_1VAR_2IN_1OUT(longlong1, %) -DECLOP_1VAR_2IN_1OUT(longlong1, &) -DECLOP_1VAR_2IN_1OUT(longlong1, |) -DECLOP_1VAR_2IN_1OUT(longlong1, ^) -DECLOP_1VAR_2IN_1OUT(longlong1, <<) -DECLOP_1VAR_2IN_1OUT(longlong1, >>) - - -DECLOP_1VAR_ASSIGN(longlong1, +=) -DECLOP_1VAR_ASSIGN(longlong1, -=) -DECLOP_1VAR_ASSIGN(longlong1, *=) -DECLOP_1VAR_ASSIGN(longlong1, /=) -DECLOP_1VAR_ASSIGN(longlong1, %=) -DECLOP_1VAR_ASSIGN(longlong1, &=) -DECLOP_1VAR_ASSIGN(longlong1, |=) -DECLOP_1VAR_ASSIGN(longlong1, ^=) -DECLOP_1VAR_ASSIGN(longlong1, <<=) -DECLOP_1VAR_ASSIGN(longlong1, >>=) - -DECLOP_1VAR_PREOP(longlong1, ++) -DECLOP_1VAR_PREOP(longlong1, --) - -DECLOP_1VAR_POSTOP(longlong1, ++) -DECLOP_1VAR_POSTOP(longlong1, --) - -DECLOP_1VAR_COMP(longlong1, ==) -DECLOP_1VAR_COMP(longlong1, !=) -DECLOP_1VAR_COMP(longlong1, <) -DECLOP_1VAR_COMP(longlong1, >) -DECLOP_1VAR_COMP(longlong1, <=) -DECLOP_1VAR_COMP(longlong1, >=) - -DECLOP_1VAR_COMP(longlong1, &&) -DECLOP_1VAR_COMP(longlong1, ||) - -DECLOP_1VAR_1IN_1OUT(longlong1, ~) -DECLOP_1VAR_1IN_BOOLOUT(longlong1, !) - -DECLOP_1VAR_SCALE_PRODUCT(longlong1, unsigned char) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, signed char) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, unsigned short) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, signed short) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, unsigned int) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, signed int) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, float) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, unsigned long) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, signed long) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, double) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, unsigned long long) -DECLOP_1VAR_SCALE_PRODUCT(longlong1, signed long long) - -// SIGNED LONGLONG2 - -DECLOP_2VAR_2IN_1OUT(longlong2, +) -DECLOP_2VAR_2IN_1OUT(longlong2, -) -DECLOP_2VAR_2IN_1OUT(longlong2, *) -DECLOP_2VAR_2IN_1OUT(longlong2, /) -DECLOP_2VAR_2IN_1OUT(longlong2, %) -DECLOP_2VAR_2IN_1OUT(longlong2, &) -DECLOP_2VAR_2IN_1OUT(longlong2, |) -DECLOP_2VAR_2IN_1OUT(longlong2, ^) -DECLOP_2VAR_2IN_1OUT(longlong2, <<) -DECLOP_2VAR_2IN_1OUT(longlong2, >>) - -DECLOP_2VAR_ASSIGN(longlong2, +=) -DECLOP_2VAR_ASSIGN(longlong2, -=) -DECLOP_2VAR_ASSIGN(longlong2, *=) -DECLOP_2VAR_ASSIGN(longlong2, /=) -DECLOP_2VAR_ASSIGN(longlong2, %=) -DECLOP_2VAR_ASSIGN(longlong2, &=) -DECLOP_2VAR_ASSIGN(longlong2, |=) -DECLOP_2VAR_ASSIGN(longlong2, ^=) -DECLOP_2VAR_ASSIGN(longlong2, <<=) -DECLOP_2VAR_ASSIGN(longlong2, >>=) - -DECLOP_2VAR_PREOP(longlong2, ++) -DECLOP_2VAR_PREOP(longlong2, --) - -DECLOP_2VAR_POSTOP(longlong2, ++) -DECLOP_2VAR_POSTOP(longlong2, --) - -DECLOP_2VAR_COMP(longlong2, ==) -DECLOP_2VAR_COMP(longlong2, !=) -DECLOP_2VAR_COMP(longlong2, <) -DECLOP_2VAR_COMP(longlong2, >) -DECLOP_2VAR_COMP(longlong2, <=) -DECLOP_2VAR_COMP(longlong2, >=) - -DECLOP_2VAR_COMP(longlong2, &&) -DECLOP_2VAR_COMP(longlong2, ||) - -DECLOP_2VAR_1IN_1OUT(longlong2, ~) -DECLOP_2VAR_1IN_BOOLOUT(longlong2, !) - -DECLOP_2VAR_SCALE_PRODUCT(longlong2, unsigned char) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, signed char) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, unsigned short) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, signed short) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, unsigned int) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, signed int) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, float) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, unsigned long) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, signed long) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, double) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, unsigned long long) -DECLOP_2VAR_SCALE_PRODUCT(longlong2, signed long long) - -// SIGNED LONGLONG3 - -DECLOP_3VAR_2IN_1OUT(longlong3, +) -DECLOP_3VAR_2IN_1OUT(longlong3, -) -DECLOP_3VAR_2IN_1OUT(longlong3, *) -DECLOP_3VAR_2IN_1OUT(longlong3, /) -DECLOP_3VAR_2IN_1OUT(longlong3, %) -DECLOP_3VAR_2IN_1OUT(longlong3, &) -DECLOP_3VAR_2IN_1OUT(longlong3, |) -DECLOP_3VAR_2IN_1OUT(longlong3, ^) -DECLOP_3VAR_2IN_1OUT(longlong3, <<) -DECLOP_3VAR_2IN_1OUT(longlong3, >>) - -DECLOP_3VAR_ASSIGN(longlong3, +=) -DECLOP_3VAR_ASSIGN(longlong3, -=) -DECLOP_3VAR_ASSIGN(longlong3, *=) -DECLOP_3VAR_ASSIGN(longlong3, /=) -DECLOP_3VAR_ASSIGN(longlong3, %=) -DECLOP_3VAR_ASSIGN(longlong3, &=) -DECLOP_3VAR_ASSIGN(longlong3, |=) -DECLOP_3VAR_ASSIGN(longlong3, ^=) -DECLOP_3VAR_ASSIGN(longlong3, <<=) -DECLOP_3VAR_ASSIGN(longlong3, >>=) - -DECLOP_3VAR_PREOP(longlong3, ++) -DECLOP_3VAR_PREOP(longlong3, --) - -DECLOP_3VAR_POSTOP(longlong3, ++) -DECLOP_3VAR_POSTOP(longlong3, --) - -DECLOP_3VAR_COMP(longlong3, ==) -DECLOP_3VAR_COMP(longlong3, !=) -DECLOP_3VAR_COMP(longlong3, <) -DECLOP_3VAR_COMP(longlong3, >) -DECLOP_3VAR_COMP(longlong3, <=) -DECLOP_3VAR_COMP(longlong3, >=) - -DECLOP_3VAR_COMP(longlong3, &&) -DECLOP_3VAR_COMP(longlong3, ||) - -DECLOP_3VAR_1IN_1OUT(longlong3, ~) -DECLOP_3VAR_1IN_BOOLOUT(longlong3, !) - -DECLOP_3VAR_SCALE_PRODUCT(longlong3, unsigned char) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, signed char) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, unsigned short) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, signed short) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, unsigned int) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, signed int) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, float) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, unsigned long) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, signed long) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, double) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, unsigned long long) -DECLOP_3VAR_SCALE_PRODUCT(longlong3, signed long long) - -// SIGNED LONGLONG4 - -DECLOP_4VAR_2IN_1OUT(longlong4, +) -DECLOP_4VAR_2IN_1OUT(longlong4, -) -DECLOP_4VAR_2IN_1OUT(longlong4, *) -DECLOP_4VAR_2IN_1OUT(longlong4, /) -DECLOP_4VAR_2IN_1OUT(longlong4, %) -DECLOP_4VAR_2IN_1OUT(longlong4, &) -DECLOP_4VAR_2IN_1OUT(longlong4, |) -DECLOP_4VAR_2IN_1OUT(longlong4, ^) -DECLOP_4VAR_2IN_1OUT(longlong4, <<) -DECLOP_4VAR_2IN_1OUT(longlong4, >>) - -DECLOP_4VAR_ASSIGN(longlong4, +=) -DECLOP_4VAR_ASSIGN(longlong4, -=) -DECLOP_4VAR_ASSIGN(longlong4, *=) -DECLOP_4VAR_ASSIGN(longlong4, /=) -DECLOP_4VAR_ASSIGN(longlong4, %=) -DECLOP_4VAR_ASSIGN(longlong4, &=) -DECLOP_4VAR_ASSIGN(longlong4, |=) -DECLOP_4VAR_ASSIGN(longlong4, ^=) -DECLOP_4VAR_ASSIGN(longlong4, <<=) -DECLOP_4VAR_ASSIGN(longlong4, >>=) - -DECLOP_4VAR_PREOP(longlong4, ++) -DECLOP_4VAR_PREOP(longlong4, --) - -DECLOP_4VAR_POSTOP(longlong4, ++) -DECLOP_4VAR_POSTOP(longlong4, --) - -DECLOP_4VAR_COMP(longlong4, ==) -DECLOP_4VAR_COMP(longlong4, !=) -DECLOP_4VAR_COMP(longlong4, <) -DECLOP_4VAR_COMP(longlong4, >) -DECLOP_4VAR_COMP(longlong4, <=) -DECLOP_4VAR_COMP(longlong4, >=) - -DECLOP_4VAR_COMP(longlong4, &&) -DECLOP_4VAR_COMP(longlong4, ||) - -DECLOP_4VAR_1IN_1OUT(longlong4, ~) -DECLOP_4VAR_1IN_BOOLOUT(longlong4, !) - -DECLOP_4VAR_SCALE_PRODUCT(longlong4, unsigned char) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, signed char) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, unsigned short) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, signed short) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, unsigned int) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, signed int) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, float) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, unsigned long) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, signed long) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, double) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, unsigned long long) -DECLOP_4VAR_SCALE_PRODUCT(longlong4, signed long long) #endif + #endif diff --git a/hipamd/src/hip_fp16.cpp b/hipamd/src/hip_fp16.cpp index c2b7b47597..8e8f003f56 100644 --- a/hipamd/src/hip_fp16.cpp +++ b/hipamd/src/hip_fp16.cpp @@ -90,11 +90,11 @@ __device__ bool __hgt(__half a, __half b) { } __device__ bool __hisinf(__half a) { - return a == __hInfValue.h ? true : false; + return a == HINF ? true : false; } __device__ bool __hisnan(__half a) { - return a > __hInfValue.h ? true : false; + return a > HINF ? true : false; } __device__ bool __hle(__half a, __half b) { @@ -114,75 +114,75 @@ Half2 Comparision Functions */ __device__ bool __hbeq2(__half2 a, __half2 b) { - return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false); + return (a.x == b.x ? true : false) && (a.y == b.y ? true : false); } __device__ bool __hbge2(__half2 a, __half2 b) { - return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false); + return (a.x >= b.x ? true : false) && (a.y >= b.y ? true : false); } __device__ bool __hbgt2(__half2 a, __half2 b) { - return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false); + return (a.x > b.x ? true : false) && (a.y > b.y ? true : false); } __device__ bool __hble2(__half2 a, __half2 b) { - return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false); + return (a.x <= b.x ? true : false) && (a.y <= b.y ? true : false); } __device__ bool __hblt2(__half2 a, __half2 b) { - return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false); + return (a.x < b.x ? true : false) && (a.y < b.y ? true : false); } __device__ bool __hbne2(__half2 a, __half2 b) { - return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false); + return (a.x != b.x ? true : false) && (a.y != b.y ? true : false); } __device__ __half2 __heq2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] == b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] == b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x == b.x) ? (__half)1 : (__half)0; + c.y = (a.y == b.y) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hge2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] >= b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] >= b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x >= b.x) ? (__half)1 : (__half)0; + c.y = (a.y >= b.y) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hgt2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] > b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] > b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x > b.x) ? (__half)1 : (__half)0; + c.y = (a.y > b.y) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hisnan2(__half2 a) { __half2 c; - c.p[0] = (a.p[0] > __hInfValue.h) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] > __hInfValue.h) ? (__half)1 : (__half)0; + c.x = (a.x > HINF) ? (__half)1 : (__half)0; + c.y = (a.y > HINF) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hle2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] <= b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] <= b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x <= b.x) ? (__half)1 : (__half)0; + c.y = (a.y <= b.y) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hlt2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] < b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] < b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x < b.x) ? (__half)1 : (__half)0; + c.y = (a.y < b.y) ? (__half)1 : (__half)0; return c; } __device__ __half2 __hne2(__half2 a, __half2 b) { __half2 c; - c.p[0] = (a.p[0] != b.p[0]) ? (__half)1 : (__half)0; - c.p[1] = (a.p[1] != b.p[1]) ? (__half)1 : (__half)0; + c.x = (a.x != b.x) ? (__half)1 : (__half)0; + c.y = (a.y != b.y) ? (__half)1 : (__half)0; return c; } @@ -191,8 +191,8 @@ Conversion instructions */ __device__ __half2 __float22half2_rn(const float2 a) { __half2 b; - b.p[0] = (__half)a.x; - b.p[1] = (__half)a.y; + b.x = (__half)a.x; + b.y = (__half)a.y; return b; } @@ -202,8 +202,8 @@ __device__ __half __float2half(const float a) { __device__ __half2 __float2half2_rn(const float a) { __half2 b; - b.p[0] = (__half)a; - b.p[1] = (__half)a; + b.x = (__half)a; + b.y = (__half)a; return b; } @@ -225,15 +225,15 @@ __device__ __half __float2half_rz(const float a) { __device__ __half2 __floats2half2_rn(const float a, const float b) { __half2 c; - c.p[0] = (__half)a; - c.p[1] = (__half)b; + c.x = (__half)a; + c.y = (__half)b; return c; } __device__ float2 __half22float2(const __half2 a) { float2 b; - b.x = (float)a.p[0]; - b.y = (float)a.p[1]; + b.x = (float)a.x; + b.y = (float)a.y; return b; } @@ -243,8 +243,8 @@ __device__ float __half2float(const __half a) { __device__ __half2 half2half2(const __half a) { __half2 b; - b.p[0] = a; - b.p[1] = a; + b.x = a; + b.y = a; return b; } @@ -358,30 +358,30 @@ __device__ unsigned short int __half_as_ushort(const __half h) { __device__ __half2 __halves2half2(const __half a, const __half b) { __half2 c; - c.p[0] = a; - c.p[1] = b; + c.x = a; + c.y = b; return c; } __device__ float __high2float(const __half2 a) { - return (float)a.p[1]; + return (float)a.y; } __device__ __half __high2half(const __half2 a) { - return a.p[1]; + return a.y; } __device__ __half2 __high2half2(const __half2 a) { __half2 b; - b.p[0] = a.p[1]; - b.p[1] = a.p[1]; + b.x = a.y; + b.y = a.y; return b; } __device__ __half2 __highs2half2(const __half2 a, const __half2 b) { __half2 c; - c.p[0] = a.p[1]; - c.p[1] = b.p[1]; + c.x = a.y; + c.y = b.y; return c; } @@ -418,38 +418,38 @@ __device__ __half __ll2half_rz(long long int i){ } __device__ float __low2float(const __half2 a) { - return (float)a.p[0]; + return (float)a.x; } __device__ __half __low2half(const __half2 a) { - return a.p[0]; + return a.x; } __device__ __half2 __low2half2(const __half2 a, const __half2 b) { __half2 c; - c.p[0] = a.p[0]; - c.p[1] = b.p[0]; + c.x = a.x; + c.y = b.x; return c; } __device__ __half2 __low2half2(const __half2 a) { __half2 b; - b.p[0] = a.p[0]; - b.p[1] = a.p[0]; + b.x = a.x; + b.y = a.x; return b; } __device__ __half2 __lowhigh2highlow(const __half2 a) { __half2 b; - b.p[0] = a.p[1]; - b.p[1] = a.p[0]; + b.x = a.y; + b.y = a.x; return b; } __device__ __half2 __lows2half2(const __half2 a, const __half2 b) { __half2 c; - c.p[0] = a.p[0]; - c.p[1] = b.p[0]; + c.y = a.x; + c.y = b.x; return c; } @@ -542,346 +542,4 @@ typedef struct{ }; } struct_float; -#if __clang_major__ == 3 -static __device__ float cvt_half_to_float(__half a){ - struct_float ret = {0}; - if(a.x == 0){ - return 0.0f; - } - if(a.x == 0x8000){ - return -0.0f; - } - ret.u = ((a.x&0x8000)<<16) | (((a.x&0x7c00)+0x1C000)<<13) | ((a.x&0x03FF)<<13); - return ret.f; -} - -static __device__ __half cvt_float_to_half(float b){ - struct_float f = {0}; - __half ret = {0}; - f.f = b; - if(f.f == 0.0f){ - ret.x = 0; - return ret; - } - if(f.f == -0.0f){ - ret.x = 0x8000; - return ret; - } - ret.x = ((f.u>>16)&0x8000)|((((f.u&0x7f800000)-0x38000000)>>13)&0x7c00)|((f.u>>13)&0x03ff); - return ret; -} - - -__device__ __half __soft_hadd(const __half a, const __half b){ - return cvt_float_to_half(cvt_half_to_float(a)+cvt_half_to_float(b)); -} - -__device__ __half __soft_hadd_sat(const __half a, const __half b){ - float f = cvt_half_to_float(a) + cvt_half_to_float(b); - return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f))); -} - -__device__ __half __soft_hfma(const __half a, const __half b, const __half c){ - return cvt_float_to_half(fmaf(cvt_half_to_float(a), cvt_half_to_float(b), cvt_half_to_float(c))); -} - -__device__ __half __soft_hfma_sat(const __half a, const __half b, const __half c){ - float f = fmaf(cvt_half_to_float(a), cvt_half_to_float(b), cvt_half_to_float(c)); - return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f))); -} - -__device__ __half __soft_hmul(const __half a, const __half b){ - return cvt_float_to_half(cvt_half_to_float(a)*cvt_half_to_float(b)); -} - -__device__ __half __soft_hmul_sat(const __half a, const __half b){ - float f = cvt_half_to_float(a) * cvt_half_to_float(b); - return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f))); -} - -__device__ __half __soft_hneq(const __half a){ - __half ret = {a.x}; - ret.x ^= 1 << 15; - return ret; -} - -__device__ __half __soft_hsub(const __half a, const __half b){ - return cvt_float_to_half(cvt_half_to_float(a)-cvt_half_to_float(b)); -} - -__device__ __half __soft_hsub_sat(const __half a, const __half b){ - float f = cvt_half_to_float(a) - cvt_half_to_float(b); - return (f < 0.0f ? __half_value_zero_float : (f > 1.0f ? __half_value_one_float: cvt_float_to_half(f))); -} - - -/* -Half2 Arithmetic Instructions -*/ - -__device__ __half2 __soft_hadd2(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hadd(a.p[1], b.p[1]); - ret.p[0] = __soft_hadd(a.p[0], b.p[0]); - return ret; -} - -__device__ __half2 __soft_hadd2_sat(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hadd_sat(a.p[1], b.p[1]); - ret.p[0] = __soft_hadd_sat(a.p[0], b.p[0]); - return ret; -} - -__device__ __half2 __soft_hfma2(const __half2 a, const __half2 b, const __half2 c){ - __half2 ret; - ret.p[1] = __soft_hfma(a.p[1], b.p[1], c.p[1]); - ret.p[0] = __soft_hfma(a.p[0], b.p[0], c.p[0]); - return ret; -} - -__device__ __half2 __soft_hfma2_sat(const __half2 a, const __half2 b, const __half2 c){ - __half2 ret; - ret.p[1] = __soft_hfma_sat(a.p[1], b.p[1], c.p[1]); - ret.p[0] = __soft_hfma_sat(a.p[0], b.p[0], c.p[0]); - return ret; -} - -__device__ __half2 __soft_hmul2(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hmul(a.p[1], b.p[1]); - ret.p[0] = __soft_hmul(a.p[0], b.p[0]); - return ret; -} - -__device__ __half2 __soft_hmul2_sat(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hmul_sat(a.p[1], b.p[1]); - ret.p[0] = __soft_hmul_sat(a.p[0], b.p[0]); - return ret; -} - -__device__ __half2 __soft_hneq2(const __half2 a){ - __half2 ret; - ret.p[1] = __soft_hneq(a.p[1]); - ret.p[0] = __soft_hneq(a.p[0]); - return ret; -} - -__device__ __half2 __soft_hsub2(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hsub(a.p[1], b.p[1]); - ret.p[0] = __soft_hsub(a.p[0], b.p[0]); - return ret; -} - -__device__ __half2 __soft_hsub2_sat(const __half2 a, const __half2 b){ - __half2 ret; - ret.p[1] = __soft_hsub_sat(a.p[1], b.p[1]); - ret.p[0] = __soft_hsub_sat(a.p[0], b.p[0]); - return ret; -} - -/* -Half Cmps -*/ - -__device__ bool __soft_heq(const __half a, const __half b){ - return (a.x == b.x ? true:false); -} - -__device__ bool __soft_hge(const __half a, const __half b){ - return (cvt_half_to_float(a) >= cvt_half_to_float(b)); -} - -__device__ bool __soft_hgt(const __half a, const __half b){ - return (cvt_half_to_float(a) > cvt_half_to_float(b)); -} - -__device__ bool __soft_hisinf(const __half a){ - return ((a.x == __half_neg_inf) ? -1 : (a.x == __half_pos_inf) ? 1 : 0); -} - -__device__ bool __soft_hisnan(const __half a){ - if(((a.x & __half_pos_inf) == a.x) || ((a.x & __half_neg_inf) == a.x)){ - return true; - }else{ - return false; - } -} - -__device__ bool __soft_hle(const __half a, const __half b){ - return (cvt_half_to_float(a) <= cvt_half_to_float(b)); -} - -__device__ bool __soft_hlt(const __half a, const __half b){ - return (cvt_half_to_float(a) < cvt_half_to_float(b)); -} - -__device__ bool __soft_hne(const __half a, const __half b){ - return a.x == b.x ? false : true; -} - -/* -Half2 Cmps -*/ - -__device__ bool __soft_hbeq2(const __half2 a, const __half2 b){ - return __soft_heq(a.p[1], b.p[1]) && __soft_heq(a.p[0], b.p[0]); -} - -__device__ bool __soft_hbge2(const __half2 a, const __half2 b){ - return __soft_hge(a.p[1], b.p[1]) && __soft_hge(a.p[0], b.p[0]); -} - -__device__ bool __soft_hbgt2(const __half2 a, const __half2 b){ - return __soft_hgt(a.p[1], b.p[1]) && __soft_hgt(a.p[0], b.p[0]); -} - -__device__ bool __soft_hble2(const __half2 a, const __half2 b){ - return __soft_hle(a.p[1], b.p[1]) && __soft_hle(a.p[0], b.p[0]); -} - -__device__ bool __soft_hblt2(const __half2 a, const __half2 b){ - return __soft_hlt(a.p[1], b.p[1]) && __soft_hlt(a.p[0], b.p[0]); -} - -__device__ bool __soft_hbne2(const __half2 a, const __half2 b){ - return __soft_hne(a.p[1], b.p[1]) && __soft_hne(a.p[0], b.p[0]); -} - - - -__device__ __half2 __soft_heq2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_heq(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_heq(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hge2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_hge(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_hge(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hgt2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_hgt(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_hgt(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hisnan2(const __half2 a){ - __half2 ret = {0}; - ret.p[1] = __soft_hisnan(a.p[1]) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = __soft_hisnan(a.p[0]) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hle2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_hle(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_hle(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hlt2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_hlt(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_hlt(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -__device__ __half2 __soft_hne2(const __half2 a, const __half2 b){ - __half2 ret = {0}; - ret.p[1] = (__soft_hne(a.p[1], b.p[1])) ? __half_value_one_float : __half_value_zero_float; - ret.p[0] = (__soft_hne(a.p[0], b.p[0])) ? __half_value_one_float : __half_value_zero_float; - return ret; -} - -/* -Half Cnvs and Data Mvmnt -*/ - -__device__ __half2 __soft_float22half2_rn(const float2 a){ - __half2 ret = {0}; - ret.p[1] = cvt_float_to_half(a.x); - ret.p[0] = cvt_float_to_half(a.y); - return ret; -} - -__device__ __half __soft_float2half(const float a){ - return cvt_float_to_half(a); -} - -__device__ __half2 __soft_float2half2_rn(const float a){ - __half ret = cvt_float_to_half(a); - return {ret, ret}; -} - -__device__ __half2 __soft_floats2half2_rn(const float a, const float b){ - return {cvt_float_to_half(a), cvt_float_to_half(b)}; -} - -__device__ float2 __soft_half22float2(const __half2 a){ - return {cvt_half_to_float(a.p[1]), cvt_half_to_float(a.p[0])}; -} - -__device__ float __soft_half2float(const __half a){ - return cvt_half_to_float(a); -} - -__device__ __half2 __soft_half2half2(const __half a){ - return {a,a}; -} - -__device__ __half2 __soft_halves2half2(const __half a, const __half b){ - return {a,b}; -} - -__device__ float __soft_high2float(const __half2 a){ - return cvt_half_to_float(a.p[1]); -} - -__device__ __half __soft_high2half(const __half2 a){ - return a.p[1]; -} - -__device__ __half2 __soft_high2half2(const __half2 a){ - return {a.p[1], a.p[1]}; -} - -__device__ __half2 __soft_highs2half2(const __half2 a, const __half2 b){ - return {a.p[1], b.p[1]}; -} - -__device__ float __soft_low2float(const __half2 a){ - return cvt_half_to_float(a.p[0]); -} - -__device__ __half __soft_low2half(const __half2 a){ - return a.p[0]; -} - -__device__ __half2 __soft_low2half2(const __half2 a){ - return {a.p[0], a.p[0]}; -} - -__device__ __half2 __soft_lows2half2(const __half2 a, const __half2 b){ - return {a.p[0], b.p[0]}; -} - -__device__ __half2 __soft_lowhigh2highlow(const __half2 a){ - return {a.p[0], a.p[1]}; -} - -__device__ __half2 __soft_low2half2(const __half2 a, const __half2 b){ - return {a.p[0], b.p[0]}; -} - - - -#endif diff --git a/hipamd/src/hip_hc_gfx803.ll b/hipamd/src/hip_hc_gfx803.ll index 0080fc7d81..7e3d0e37dd 100644 --- a/hipamd/src/hip_hc_gfx803.ll +++ b/hipamd/src/hip_hc_gfx803.ll @@ -2,89 +2,122 @@ target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64: target triple = "amdgcn--amdhsa" -define i32 @__hip_hc_ir_hadd2_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_add_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b) - ret i32 %1 +define <2 x half> @__hip_hc_ir_hadd2_int(<2 x half> %a, <2 x half> %b) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = bitcast <2 x half> %b to i32 + %3 = tail call i32 asm sideeffect "v_add_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2) + tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2) + %4 = bitcast i32 %3 to <2 x half> + ret <2 x half> %4 } -define i32 @__hip_hc_ir_hfma2_int(i32 %a, i32 %b, i32 %c) #1 { - %1 = tail call i32 asm sideeffect "v_mad_f16 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c) - tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b) - tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %1, i32 %c) - ret i32 %1 +define <2 x half> @__hip_hc_ir_hfma2_int(<2 x half> %a, <2 x half> %b, <2 x half> %c) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = bitcast <2 x half> %b to i32 + %3 = bitcast <2 x half> %c to i32 + %4 = tail call i32 asm sideeffect "v_mad_f16 $0, $1, $2, $3","=v,v,v,v"(i32 %1, i32 %2, i32 %3) + tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %4, i32 %1, i32 %2) + tail call void asm sideeffect "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %4, i32 %4, i32 %3) + %5 = bitcast i32 %4 to <2 x half> + ret <2 x half> %5 } -define i32 @__hip_hc_ir_hmul2_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_mul_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b) - ret i32 %1 +define <2 x half> @__hip_hc_ir_hmul2_int(<2 x half> %a, <2 x half> %b) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = bitcast <2 x half> %b to i32 + %3 = tail call i32 asm sideeffect "v_mul_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2) + tail call void asm sideeffect "v_mul_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2) + %4 = bitcast i32 %3 to <2 x half> + ret <2 x half> %4 } -define i32 @__hip_hc_ir_hsub2_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_sub_f16 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - tail call void asm sideeffect "v_sub_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %1, i32 %a, i32 %b) - ret i32 %1 +define <2 x half> @__hip_hc_ir_hsub2_int(<2 x half> %a, <2 x half> %b) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = bitcast <2 x half> %b to i32 + %3 = tail call i32 asm sideeffect "v_sub_f16 $0, $1, $2","=v,v,v"(i32 %1, i32 %2) + tail call void asm sideeffect "v_sub_f16_sdwa $0, $1, $2 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1 src1_sel:WORD_1","v,v,v"(i32 %3, i32 %1, i32 %2) + %4 = bitcast i32 %3 to <2 x half> + ret <2 x half> %4 } -define i32 @__hip_hc_ir_h2ceil_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_ceil_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_ceil_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2ceil_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_ceil_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_ceil_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2cos_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_cos_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_cos_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2cos_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_cos_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_cos_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2exp2_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_exp_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_exp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2exp2_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_exp_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_exp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2floor_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_floor_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_floor_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2floor_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_floor_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_floor_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2log2_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_log_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_log_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2log2_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_log_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_log_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2rcp_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_rcp_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_rcp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2rcp_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_rcp_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_rcp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2rsqrt_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_rsq_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_rsq_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2rsqrt_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_rsq_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_rsq_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2sin_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_sin_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_sin_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2sin_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_sin_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_sin_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2sqrt_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_sqrt_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_sqrt_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2sqrt_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_sqrt_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_sqrt_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } -define i32 @__hip_hc_ir_h2trunc_int(i32 %a) #1 { - %1 = tail call i32 asm sideeffect "v_trunc_f16 $0, $1","=v,v"(i32 %a) - tail call void asm sideeffect "v_trunc_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a) - ret i32 %1 +define <2 x half> @__hip_hc_ir_h2trunc_int(<2 x half> %a) #1 { + %1 = bitcast <2 x half> %a to i32 + %2 = tail call i32 asm sideeffect "v_trunc_f16 $0, $1","=v,v"(i32 %1) + tail call void asm sideeffect "v_trunc_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %2, i32 %1) + %3 = bitcast i32 %2 to <2 x half> + ret <2 x half> %3 } attributes #1 = { alwaysinline nounwind }