From 145adc186af007aa34546ba95ff582f2e281dbcf Mon Sep 17 00:00:00 2001 From: JAR Date: Sat, 13 Apr 2013 17:12:56 -0400 Subject: [PATCH 1/2] Added support for more OpenCL datatypes and builtin functions --- src/libcoprthr/opencl_lift.h | 2038 ++++++++++++---------- src/libcoprthr/ser_engine.h | 28 + src/libcoprthr/sl_engine.h | 28 + src/libstdcl/clforka.h | 122 +- src/libstdcl/stdcl_kernel_openclcompat.h | 51 +- tools/xclnm/xclnm_gram.y | 64 + tools/xclnm/xclnm_scan.l | 34 +- 7 files changed, 1444 insertions(+), 921 deletions(-) diff --git a/src/libcoprthr/opencl_lift.h b/src/libcoprthr/opencl_lift.h index a2938fe8..00ec93ca 100644 --- a/src/libcoprthr/opencl_lift.h +++ b/src/libcoprthr/opencl_lift.h @@ -62,18 +62,60 @@ typedef unsigned int uint; #endif -typedef unsigned int uint; typedef unsigned char uchar; +typedef unsigned short ushort; +typedef unsigned int uint; -extern "C" { - -} #define CLK_LOCAL_MEM_FENCE 1 #define CLK_GLOBAL_MEM_FENCE 1 #define prefetch( x, y ) +/*** math constants [6.11.2] ***/ + +#undef M_E +#undef M_LOG2E +#undef M_LOG10E +#undef M_LN2 +#undef M_LN10 +#undef M_PI +#undef M_PI_2 +#undef M_PI_4 +#undef M_1_PI +#undef M_2_PI +#undef M_2_SQRTPI +#undef M_SQRT2 +#undef M_SQRT1_2 + +#define M_E 2.718281828459045090796 +#define M_LOG2E 1.442695040888963387005 +#define M_LOG10E 0.434294481903251816668 +#define M_LN2 0.693147180559945286227 +#define M_LN10 2.302585092994045901094 +#define M_PI 3.141592653589793115998 +#define M_PI_2 1.570796326794896557999 +#define M_PI_4 0.785398163397448278999 +#define M_1_PI 0.318309886183790691216 +#define M_2_PI 0.636619772367581382433 +#define M_2_SQRTPI 1.128379167095512558561 +#define M_SQRT2 1.414213562373095145475 +#define M_SQRT1_2 0.707106781186547572737 + +#define M_E_F 2.71828174591064f +#define M_LOG2E_F 1.44269502162933f +#define M_LOG10E_F 0.43429449200630f +#define M_LN2_F 0.69314718246460f +#define M_LN10_F 2.30258512496948f +#define M_PI_F 3.14159274101257f +#define M_PI_2_F 1.57079637050629f +#define M_PI_4_F 0.78539818525314f +#define M_1_PI_F 0.31830987334251f +#define M_2_PI_F 0.63661974668503f +#define M_2_SQRTPI_F 1.12837922573090f +#define M_SQRT2_F 1.41421353816986f +#define M_SQRT1_2_F 0.70710676908493f + /*** adress space qualifiers [6.5] ***/ #define __global @@ -112,6 +154,7 @@ typedef union { \ __##cltype##2 vec; \ struct { ctype x,y; }; \ struct { ctype s0,s1; }; \ + struct { ctype lo,hi; }; \ ctype s[2]; \ } _##cltype##2; \ typedef _##cltype##2 cltype##2; @@ -125,6 +168,7 @@ typedef union { \ struct { ctype s0,s1,s2,s3; }; \ struct { cltype##2 xy,zw; }; \ struct { cltype##2 s01,s23; }; \ + struct { cltype##2 lo,hi; }; \ ctype s[4]; \ } _##cltype##4; \ typedef _##cltype##4 cltype##4; @@ -134,13 +178,34 @@ typedef ctype __##cltype##8 __attribute__((vector_size(8*sizeof(ctype)))); \ typedef union { \ typedef ctype type_t; \ __##cltype##8 vec; \ + struct { ctype x,y,z,w,__spacer4,__spacer5,__spacer6,__spacer7; }; \ struct { ctype s0,s1,s2,s3,s4,s5,s6,s7; }; \ + struct { cltype##2 xy,zw,__spacer45,__spacer67; }; \ struct { cltype##2 s01,s23,s45,s67; }; \ struct { cltype##4 s0123,s4567; }; \ + struct { cltype##4 lo,hi; }; \ ctype s[8]; \ } _##cltype##8; \ typedef _##cltype##8 cltype##8; +#define VECTOR_TYPE_16( ctype, cltype ) \ +typedef ctype __##cltype##16 __attribute__((vector_size(16*sizeof(ctype)))); \ +typedef union { \ + typedef ctype type_t; \ + __##cltype##16 vec; \ + struct { ctype x,y,z,w,__spacer4,__spacer5,__spacer6,__spacer7,__spacer8,__spacer9,sa,sb,sc,sd,se,sf; }; \ + struct { ctype s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sA,sB,sC,sD,sE,sF; }; \ + struct { cltype##2 xy,zw,__spacer45,__spacer67,__spacer89,sab,scd,sef; }; \ + struct { cltype##2 s01,s23,s45,s67,s89,sAB,sCD,sEF; }; \ + struct { cltype##4 xyzw,__spacer4567,s89ab,scdef; }; \ + struct { cltype##4 s0123,s4567,s89AB,sCDEF; }; \ + struct { cltype##8 s01234567,s89abcdef; }; \ + struct { cltype##8 __spacer01234567,s89ABCDEF; }; \ + struct { cltype##8 lo,hi; }; \ + ctype s[16]; \ +} _##cltype##16; \ +typedef _##cltype##16 cltype##16; + VECTOR_TYPE_2(char,char) #define __vector_char2(a,b) (__char2){a,b} #define _vector_char2(a,b) (_char2){(__char2){a,b}} @@ -151,6 +216,16 @@ VECTOR_TYPE_2( unsigned char, uchar ) #define _vector_uchar2(a,b) (_uchar2){(__uchar2){a,b}} #define vector_uchar2(a,b) _vector_uchar2(a,b) +VECTOR_TYPE_2(short,short) +#define __vector_short2(a,b) (__short2){a,b} +#define _vector_short2(a,b) (_short2){(__short2){a,b}} +#define vector_short2(a,b) _vector_short2(a,b) + +VECTOR_TYPE_2( unsigned short, ushort ) +#define __vector_ushort2(a,b) (__ushort2){a,b} +#define _vector_ushort2(a,b) (_ushort2){(__ushort2){a,b}} +#define vector_ushort2(a,b) _vector_ushort2(a,b) + VECTOR_TYPE_2( int, int ) #define __vector_int2(a,b) (__int2){a,b} #define _vector_int2(a,b) (_int2){(__int2){a,b}} @@ -191,6 +266,16 @@ VECTOR_TYPE_4( unsigned char, uchar ) #define _vector_uchar4(a,b,c,d) (_uchar4){(__uchar4){a,b,c,d}} #define vector_uchar4(a,b,c,d) _vector_uchar4(a,b,c,d) +VECTOR_TYPE_4( short, short ) +#define __vector_short4(a,b,c,d) (__short4){a,b,c,d} +#define _vector_short4(a,b,c,d) (_short4){(__short4){a,b,c,d}} +#define vector_short4(a,b,c,d) _vector_short4(a,b,c,d) + +VECTOR_TYPE_4( unsigned short, ushort ) +#define __vector_ushort4(a,b,c,d) (__ushort4){a,b,c,d} +#define _vector_ushort4(a,b,c,d) (_ushort4){(__ushort4){a,b,c,d}} +#define vector_ushort4(a,b,c,d) _vector_ushort4(a,b,c,d) + VECTOR_TYPE_4( int, int ) #define __vector_int4(a,b,c,d) (__int4){a,b,c,d} #define _vector_int4(a,b,c,d) (_int4){(__int4){a,b,c,d}} @@ -221,7 +306,6 @@ VECTOR_TYPE_4( double, double ) #define _vector_double4(a,b,c,d) (_double4){(__double4){a,b,c,d}} #define vector_double4(a,b,c,d) _vector_double4(a,b,c,d) - VECTOR_TYPE_8( char, char ) #define __vector_char8(s0,s1,s2,s3,s4,s5,s6,s7) \ (__char8){s0,s1,s2,s3,s4,s5,s6,s7} @@ -238,6 +322,22 @@ VECTOR_TYPE_8( unsigned char, uchar ) #define vector_uchar8(s0,s1,s2,s3,s4,s5,s6,s7) \ _vector_uchar8(s0,s1,s2,s3,s4,s5,s6,s7) +VECTOR_TYPE_8( short, short ) +#define __vector_short8(s0,s1,s2,s3,s4,s5,s6,s7) \ + (__short8){s0,s1,s2,s3,s4,s5,s6,s7} +#define _vector_short8(s0,s1,s2,s3,s4,s5,s6,s7) \ + (_short8){(__short8){s0,s1,s2,s3,s4,s5,s6,s7}} +#define vector_short8(s0,s1,s2,s3,s4,s5,s6,s7) \ + _vector_short8(s0,s1,s2,s3,s4,s5,s6,s7) + +VECTOR_TYPE_8( unsigned short, ushort ) +#define __vector_ushort8(s0,s1,s2,s3,s4,s5,s6,s7) \ + (__ushort8){s0,s1,s2,s3,s4,s5,s6,s7} +#define _vector_ushort8(s0,s1,s2,s3,s4,s5,s6,s7) \ + (_ushort8){(__ushort8){s0,s1,s2,s3,s4,s5,s6,s7}} +#define vector_ushort8(s0,s1,s2,s3,s4,s5,s6,s7) \ + _vector_ushort8(s0,s1,s2,s3,s4,s5,s6,s7) + VECTOR_TYPE_8( int, int ) #define __vector_int8(s0,s1,s2,s3,s4,s5,s6,s7) \ (__int8){s0,s1,s2,s3,s4,s5,s6,s7} @@ -286,478 +386,231 @@ VECTOR_TYPE_8( double, double ) #define vector_double8(s0,s1,s2,s3,s4,s5,s6,s7) \ _vector_double8(s0,s1,s2,s3,s4,s5,s6,s7) - - -#define GENERIC_UNARY_PLUS(type) \ -static __always_inline _##type operator + ( _##type& a ); \ -static __always_inline _##type operator + ( _##type& a ) { return a; } - -#define GENERIC_UNARY_OP(type,op) \ -static __always_inline _##type operator - ( _##type& a ); \ -static __always_inline _##type operator - ( _##type& a ) \ - { return (_##type){ -a.vec }; } - -#define EXPLICIT_UNARY_OP_VEC2(type,op) \ -static __always_inline _##type operator op ( _##type& a ); \ -static __always_inline _##type operator op ( _##type& a ) \ +VECTOR_TYPE_16( char, char ) +#define __vector_char16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__char16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_char16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_char16){(__char16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_char16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_char16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( unsigned char, uchar ) +#define __vector_uchar16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__uchar16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_uchar16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_uchar16){(__uchar16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_uchar16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_uchar16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( short, short ) +#define __vector_short16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__short16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_short16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_short16){(__short16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_short16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_short16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( unsigned short, ushort ) +#define __vector_ushort16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__ushort16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_ushort16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_ushort16){(__ushort16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_ushort16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_ushort16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( int, int ) +#define __vector_int16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__int16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_int16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_int16){(__int16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_int16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_int16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( unsigned int, uint ) +#define __vector_uint16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__uint16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_uint16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_uint16){(__uint16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_uint16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_uint16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( long long, long ) +#define __vector_long16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__long16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_long16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_long16){(__long16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_long16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_long16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( unsigned long long, ulong ) +#define __vector_ulong16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__ulong16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_ulong16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_ulong16){(__ulong16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_ulong16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_ulong16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( float, float ) +#define __vector_float16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__float16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_float16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_float16){(__float16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_float16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_float16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + +VECTOR_TYPE_16( double, double ) +#define __vector_double16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (__double16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf} +#define _vector_double16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + (_double16){(__double16){s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf}} +#define vector_double16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) \ + _vector_double16(s0,s1,s2,s3,s4,s5,s6,s7,s8,s9,sa,sb,sc,sd,se,sf) + + +#define GENERIC_UNARY_PLUS_SIZE(vectype,size) \ +static __always_inline _##vectype##size operator + ( _##vectype##size& a ); \ +static __always_inline _##vectype##size operator + ( _##vectype##size& a ) { return a; } + +#define GENERIC_UNARY_PLUS(vectype,noop) \ +GENERIC_UNARY_PLUS_SIZE(vectype,2) \ +GENERIC_UNARY_PLUS_SIZE(vectype,4) \ +GENERIC_UNARY_PLUS_SIZE(vectype,8) \ +GENERIC_UNARY_PLUS_SIZE(vectype,16) + +#define GENERIC_UNARY_OP_SIZE(vectype,op,size) \ +static __always_inline _##vectype##size operator op ( _##vectype##size& a ); \ +static __always_inline _##vectype##size operator op ( _##vectype##size& a ) { return (_##vectype##size){ op a.vec }; } + +#define GENERIC_UNARY_OP(vectype,op) \ +GENERIC_UNARY_OP_SIZE(vectype,op,2) \ +GENERIC_UNARY_OP_SIZE(vectype,op,4) \ +GENERIC_UNARY_OP_SIZE(vectype,op,8) \ +GENERIC_UNARY_OP_SIZE(vectype,op,16) + +#define EXPLICIT_UNARY_OP_ALL_TYPE_SIZE(vectype,op) \ +static __always_inline _##vectype##2 operator op ( _##vectype##2& a ); \ +static __always_inline _##vectype##2 operator op ( _##vectype##2& a ) \ + { \ + _##vectype##2 tmp \ + = _vector_##vectype##2( (vectype##2::type_t)op a.s0,(vectype##2::type_t)op a.s1 ); \ + return tmp; \ + } \ +static __always_inline _##vectype##4 operator op ( _##vectype##4& a ); \ +static __always_inline _##vectype##4 operator op ( _##vectype##4& a ) \ + { \ + _##vectype##4 tmp = _vector_##vectype##4( (vectype##4::type_t)op a.s0, \ + (vectype##4::type_t)op a.s1,(vectype##4::type_t)op a.s2,(vectype##4::type_t)op a.s3); \ + return tmp; \ + } \ +static __always_inline _##vectype##8 operator op ( _##vectype##8& a ); \ +static __always_inline _##vectype##8 operator op ( _##vectype##8& a ) \ + { \ + _##vectype##8 tmp = _vector_##vectype##8( (vectype##8::type_t)op a.s0,(vectype##8::type_t)op a.s1, \ + (vectype##8::type_t)op a.s2,(vectype##8::type_t)op a.s3,(vectype##8::type_t)op a.s4, \ + (vectype##8::type_t)op a.s5,(vectype##8::type_t)op a.s6,(vectype##8::type_t)op a.s7); \ + return tmp; \ + } \ +static __always_inline _##vectype##16 operator op ( _##vectype##16& a ); \ +static __always_inline _##vectype##16 operator op ( _##vectype##16& a ) \ + { \ + _##vectype##16 tmp = _vector_##vectype##16( (vectype##16::type_t)op a.s0,(vectype##16::type_t)op a.s1, \ + (vectype##16::type_t)op a.s2,(vectype##16::type_t)op a.s3,(vectype##16::type_t)op a.s4, \ + (vectype##16::type_t)op a.s5,(vectype##16::type_t)op a.s6,(vectype##16::type_t)op a.s7, \ + (vectype##16::type_t)op a.s8,(vectype##16::type_t)op a.s9,(vectype##16::type_t)op a.sA, \ + (vectype##16::type_t)op a.sB,(vectype##16::type_t)op a.sC,(vectype##16::type_t)op a.sD, \ + (vectype##16::type_t)op a.sE,(vectype##16::type_t)op a.sF); \ + return tmp; \ + } + +#define GENERIC_OP_ASSIGN_SIZE(vectype,op,size) \ +static __always_inline _##vectype##size operator op##= ( _##vectype##size& lhs, const _##vectype##size rhs ); \ +static __always_inline _##vectype##size operator op##= ( _##vectype##size& lhs, const _##vectype##size rhs ) { lhs.vec op##= rhs.vec; return lhs; } + +#define GENERIC_OP_ASSIGN(vectype,op) \ +GENERIC_OP_ASSIGN_SIZE(vectype,op,2) \ +GENERIC_OP_ASSIGN_SIZE(vectype,op,4) \ +GENERIC_OP_ASSIGN_SIZE(vectype,op,8) \ +GENERIC_OP_ASSIGN_SIZE(vectype,op,16) + +#define EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE(vectype,op) \ +static __always_inline _##vectype##2 \ + operator op##= ( _##vectype##2& lhs, const _##vectype##2 rhs ); \ +static __always_inline _##vectype##2 \ + operator op##= ( _##vectype##2& lhs, const _##vectype##2 rhs ) \ { \ - _##type tmp \ - = _vector_##type( (type::type_t)op a.s0,(type::type_t)op a.s1 ); \ - return tmp; \ - } - -#define EXPLICIT_UNARY_OP_VEC4(type,op) \ -static __always_inline _##type operator op ( _##type& a ); \ -static __always_inline _##type operator op ( _##type& a ) \ + lhs.vec = __vector_##vectype##2( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1 ); \ + return lhs; \ + } \ +static __always_inline _##vectype##4 \ + operator op##= ( _##vectype##4& lhs, const _##vectype##4 rhs ); \ +static __always_inline _##vectype##4 \ + operator op##= ( _##vectype##4& lhs, const _##vectype##4 rhs ) \ { \ - _##type tmp = _vector_##type( (type::type_t)op a.s0, \ - (type::type_t)op a.s1,(type::type_t)op a.s2,(type::type_t)op a.s3); \ - return tmp; \ - } - -#define EXPLICIT_UNARY_OP_VEC8(type,op) \ -static __always_inline _##type operator op ( _##type& a ); \ -static __always_inline _##type operator op ( _##type& a ) \ + lhs.vec = __vector_##vectype##4( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1, \ + lhs.s2 op rhs.s2, lhs.s3 op rhs.s3 ); \ + return lhs; \ + } \ +static __always_inline _##vectype##8 \ + operator op##= ( _##vectype##8& lhs, const _##vectype##8 rhs ); \ +static __always_inline _##vectype##8 \ + operator op##= ( _##vectype##8& lhs, const _##vectype##8 rhs ) \ { \ - _##type tmp = _vector_##type( \ - (type::type_t)op a.s0, \ - (type::type_t)op a.s1, \ - (type::type_t)op a.s2, \ - (type::type_t)op a.s3, \ - (type::type_t)op a.s4, \ - (type::type_t)op a.s5, \ - (type::type_t)op a.s6, \ - (type::type_t)op a.s7); \ - return tmp; \ - } - -#define GENERIC_OP_ASSIGN(type,op) \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ); \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ) \ - { lhs.vec op##= rhs.vec; return lhs; } - -#define EXPLICIT_OP_ASSIGN_VEC2(type,op) \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ); \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ) \ - { \ - lhs.vec = __vector_##type( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1 ); \ - return lhs; \ - } - -#define EXPLICIT_OP_ASSIGN_VEC4(type,op) \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ); \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ) \ - { \ - lhs.vec = __vector_##type( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1, \ - lhs.s2 op rhs.s2, lhs.s3 op rhs.s3 ); \ - return lhs; \ - } - -#define EXPLICIT_OP_ASSIGN_VEC8(type,op) \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ); \ -static __always_inline _##type \ - operator op##= ( _##type& lhs, const _##type rhs ) \ - { \ - lhs.vec = __vector_##type( \ - lhs.s0 op rhs.s0, lhs.s1 op rhs.s1, \ - lhs.s2 op rhs.s2, lhs.s3 op rhs.s3, \ - lhs.s4 op rhs.s4, lhs.s5 op rhs.s5, \ - lhs.s6 op rhs.s6, lhs.s7 op rhs.s7 ); \ - return lhs; \ + lhs.vec = __vector_##vectype##8( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1, \ + lhs.s2 op rhs.s2, lhs.s3 op rhs.s3, lhs.s4 op rhs.s4, \ + lhs.s5 op rhs.s5, lhs.s6 op rhs.s6, lhs.s7 op rhs.s7 ); \ + return lhs; \ + } \ +static __always_inline _##vectype##16 \ + operator op##= ( _##vectype##16& lhs, const _##vectype##16 rhs ); \ +static __always_inline _##vectype##16 \ + operator op##= ( _##vectype##16& lhs, const _##vectype##16 rhs ) \ + { \ + lhs.vec = __vector_##vectype##16( lhs.s0 op rhs.s0, lhs.s1 op rhs.s1, \ + lhs.s2 op rhs.s2, lhs.s3 op rhs.s3, lhs.s4 op rhs.s4, \ + lhs.s5 op rhs.s5, lhs.s6 op rhs.s6, lhs.s7 op rhs.s7, \ + lhs.s8 op rhs.s8, lhs.s9 op rhs.s9, lhs.sA op rhs.sA, \ + lhs.sB op rhs.sB, lhs.sC op rhs.sC, lhs.sD op rhs.sD, \ + lhs.sE op rhs.sE, lhs.sF op rhs.sF); \ + return lhs; \ } +#define INT_TYPE(m,...) \ +m(char,__VA_ARGS__) \ +m(uchar,__VA_ARGS__) \ +m(short,__VA_ARGS__) \ +m(ushort,__VA_ARGS__) \ +m(int,__VA_ARGS__) \ +m(uint,__VA_ARGS__) \ +m(long,__VA_ARGS__) \ +m(ulong,__VA_ARGS__) + +#define ALL_TYPE(m,...) \ +INT_TYPE(m,__VA_ARGS__) \ +m(float,__VA_ARGS__) \ +m(double,__VA_ARGS__) /* unary plus */ -GENERIC_UNARY_PLUS(char2) -GENERIC_UNARY_PLUS(uchar2) -GENERIC_UNARY_PLUS(int2) -GENERIC_UNARY_PLUS(uint2) -GENERIC_UNARY_PLUS(long2) -GENERIC_UNARY_PLUS(ulong2) -GENERIC_UNARY_PLUS(float2) -GENERIC_UNARY_PLUS(double2) -GENERIC_UNARY_PLUS(char4) -GENERIC_UNARY_PLUS(uchar4) -GENERIC_UNARY_PLUS(int4) -GENERIC_UNARY_PLUS(uint4) -GENERIC_UNARY_PLUS(long4) -GENERIC_UNARY_PLUS(ulong4) -GENERIC_UNARY_PLUS(float4) -GENERIC_UNARY_PLUS(double4) -GENERIC_UNARY_PLUS(char8) -GENERIC_UNARY_PLUS(uchar8) -GENERIC_UNARY_PLUS(int8) -GENERIC_UNARY_PLUS(uint8) -GENERIC_UNARY_PLUS(long8) -GENERIC_UNARY_PLUS(ulong8) -GENERIC_UNARY_PLUS(float8) -GENERIC_UNARY_PLUS(double8) - - -/* unary minus */ -#if GCC_VERSION < 40600 -EXPLICIT_UNARY_OP_VEC2(char2,-) -EXPLICIT_UNARY_OP_VEC2(uchar2,-) -EXPLICIT_UNARY_OP_VEC2(int2,-) -EXPLICIT_UNARY_OP_VEC2(uint2,-) -EXPLICIT_UNARY_OP_VEC2(long2,-) -EXPLICIT_UNARY_OP_VEC2(ulong2,-) -EXPLICIT_UNARY_OP_VEC2(float2,-) -EXPLICIT_UNARY_OP_VEC2(double2,-) -EXPLICIT_UNARY_OP_VEC4(char4,-) -EXPLICIT_UNARY_OP_VEC4(uchar4,-) -EXPLICIT_UNARY_OP_VEC4(int4,-) -EXPLICIT_UNARY_OP_VEC4(uint4,-) -EXPLICIT_UNARY_OP_VEC4(long4,-) -EXPLICIT_UNARY_OP_VEC4(ulong4,-) -EXPLICIT_UNARY_OP_VEC4(float4,-) -EXPLICIT_UNARY_OP_VEC4(double4,-) -EXPLICIT_UNARY_OP_VEC8(char8,-) -EXPLICIT_UNARY_OP_VEC8(uchar8,-) -EXPLICIT_UNARY_OP_VEC8(int8,-) -EXPLICIT_UNARY_OP_VEC8(uint8,-) -EXPLICIT_UNARY_OP_VEC8(long8,-) -EXPLICIT_UNARY_OP_VEC8(ulong8,-) -EXPLICIT_UNARY_OP_VEC8(float8,-) -EXPLICIT_UNARY_OP_VEC8(double8,-) -#else -GENERIC_UNARY_OP(char2,-) -GENERIC_UNARY_OP(uchar2,-) -GENERIC_UNARY_OP(int2,-) -GENERIC_UNARY_OP(uint2,-) -GENERIC_UNARY_OP(long2,-) -GENERIC_UNARY_OP(ulong2,-) -GENERIC_UNARY_OP(float2,-) -GENERIC_UNARY_OP(double2,-) -GENERIC_UNARY_OP(char4,-) -GENERIC_UNARY_OP(uchar4,-) -GENERIC_UNARY_OP(int4,-) -GENERIC_UNARY_OP(uint4,-) -GENERIC_UNARY_OP(long4,-) -GENERIC_UNARY_OP(ulong4,-) -GENERIC_UNARY_OP(float4,-) -GENERIC_UNARY_OP(double4,-) -#endif - -/* unary bitwise not */ -EXPLICIT_UNARY_OP_VEC2(char2,~) -EXPLICIT_UNARY_OP_VEC2(uchar2,~) -EXPLICIT_UNARY_OP_VEC2(int2,~) -EXPLICIT_UNARY_OP_VEC2(uint2,~) -EXPLICIT_UNARY_OP_VEC2(long2,~) -EXPLICIT_UNARY_OP_VEC2(ulong2,~) -EXPLICIT_UNARY_OP_VEC4(char4,~) -EXPLICIT_UNARY_OP_VEC4(uchar4,~) -EXPLICIT_UNARY_OP_VEC4(int4,~) -EXPLICIT_UNARY_OP_VEC4(uint4,~) -EXPLICIT_UNARY_OP_VEC4(long4,~) -EXPLICIT_UNARY_OP_VEC4(ulong4,~) -EXPLICIT_UNARY_OP_VEC8(char8,~) -EXPLICIT_UNARY_OP_VEC8(uchar8,~) -EXPLICIT_UNARY_OP_VEC8(int8,~) -EXPLICIT_UNARY_OP_VEC8(uint8,~) -EXPLICIT_UNARY_OP_VEC8(long8,~) -EXPLICIT_UNARY_OP_VEC8(ulong8,~) - - -/* add assign */ -#if GCC_VERSION < 40600 -EXPLICIT_OP_ASSIGN_VEC2(char2,+) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,+) -EXPLICIT_OP_ASSIGN_VEC2(int2,+) -EXPLICIT_OP_ASSIGN_VEC2(uint2,+) -EXPLICIT_OP_ASSIGN_VEC2(long2,+) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,+) -EXPLICIT_OP_ASSIGN_VEC2(float2,+) -EXPLICIT_OP_ASSIGN_VEC2(double2,+) -EXPLICIT_OP_ASSIGN_VEC4(char4,+) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,+) -EXPLICIT_OP_ASSIGN_VEC4(int4,+) -EXPLICIT_OP_ASSIGN_VEC4(uint4,+) -EXPLICIT_OP_ASSIGN_VEC4(long4,+) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,+) -EXPLICIT_OP_ASSIGN_VEC4(float4,+) -EXPLICIT_OP_ASSIGN_VEC4(double4,+) -EXPLICIT_OP_ASSIGN_VEC8(char8,+) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,+) -EXPLICIT_OP_ASSIGN_VEC8(int8,+) -EXPLICIT_OP_ASSIGN_VEC8(uint8,+) -EXPLICIT_OP_ASSIGN_VEC8(long8,+) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,+) -EXPLICIT_OP_ASSIGN_VEC8(float8,+) -EXPLICIT_OP_ASSIGN_VEC8(double8,+) -#else -GENERIC_OP_ASSIGN(char2,+) -GENERIC_OP_ASSIGN(uchar2,+) -GENERIC_OP_ASSIGN(int2,+) -GENERIC_OP_ASSIGN(uint2,+) -GENERIC_OP_ASSIGN(long2,+) -GENERIC_OP_ASSIGN(ulong2,+) -GENERIC_OP_ASSIGN(float2,+) -GENERIC_OP_ASSIGN(double2,+) -GENERIC_OP_ASSIGN(char4,+) -GENERIC_OP_ASSIGN(uchar4,+) -GENERIC_OP_ASSIGN(int4,+) -GENERIC_OP_ASSIGN(uint4,+) -GENERIC_OP_ASSIGN(long4,+) -GENERIC_OP_ASSIGN(ulong4,+) -GENERIC_OP_ASSIGN(float4,+) -GENERIC_OP_ASSIGN(double4,+) -GENERIC_OP_ASSIGN(char8,+) -GENERIC_OP_ASSIGN(uchar8,+) -GENERIC_OP_ASSIGN(int8,+) -GENERIC_OP_ASSIGN(uint8,+) -GENERIC_OP_ASSIGN(long8,+) -GENERIC_OP_ASSIGN(ulong8,+) -GENERIC_OP_ASSIGN(float8,+) -GENERIC_OP_ASSIGN(double8,+) -#endif +ALL_TYPE(GENERIC_UNARY_PLUS) -/* subtract assign */ #if GCC_VERSION < 40600 -EXPLICIT_OP_ASSIGN_VEC2(char2,-) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,-) -EXPLICIT_OP_ASSIGN_VEC2(int2,-) -EXPLICIT_OP_ASSIGN_VEC2(uint2,-) -EXPLICIT_OP_ASSIGN_VEC2(long2,-) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,-) -EXPLICIT_OP_ASSIGN_VEC2(float2,-) -EXPLICIT_OP_ASSIGN_VEC2(double2,-) -EXPLICIT_OP_ASSIGN_VEC4(char4,-) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,-) -EXPLICIT_OP_ASSIGN_VEC4(int4,-) -EXPLICIT_OP_ASSIGN_VEC4(uint4,-) -EXPLICIT_OP_ASSIGN_VEC4(long4,-) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,-) -EXPLICIT_OP_ASSIGN_VEC4(float4,-) -EXPLICIT_OP_ASSIGN_VEC4(double4,-) -EXPLICIT_OP_ASSIGN_VEC8(char8,-) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,-) -EXPLICIT_OP_ASSIGN_VEC8(int8,-) -EXPLICIT_OP_ASSIGN_VEC8(uint8,-) -EXPLICIT_OP_ASSIGN_VEC8(long8,-) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,-) -EXPLICIT_OP_ASSIGN_VEC8(float8,-) -EXPLICIT_OP_ASSIGN_VEC8(double8,-) +ALL_TYPE(EXPLICIT_UNARY_OP_ALL_TYPE_SIZE,-) /* unary minus */ +ALL_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,+) /* add assign */ +ALL_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,-) /* subtract assign */ +ALL_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,*) /* multiply assign */ +ALL_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,/) /* divide assign */ #else -GENERIC_OP_ASSIGN(char2,-) -GENERIC_OP_ASSIGN(uchar2,-) -GENERIC_OP_ASSIGN(int2,-) -GENERIC_OP_ASSIGN(uint2,-) -GENERIC_OP_ASSIGN(long2,-) -GENERIC_OP_ASSIGN(ulong2,-) -GENERIC_OP_ASSIGN(float2,-) -GENERIC_OP_ASSIGN(double2,-) -GENERIC_OP_ASSIGN(char4,-) -GENERIC_OP_ASSIGN(uchar4,-) -GENERIC_OP_ASSIGN(int4,-) -GENERIC_OP_ASSIGN(uint4,-) -GENERIC_OP_ASSIGN(long4,-) -GENERIC_OP_ASSIGN(ulong4,-) -GENERIC_OP_ASSIGN(float4,-) -GENERIC_OP_ASSIGN(double4,-) -GENERIC_OP_ASSIGN(char8,-) -GENERIC_OP_ASSIGN(uchar8,-) -GENERIC_OP_ASSIGN(int8,-) -GENERIC_OP_ASSIGN(uint8,-) -GENERIC_OP_ASSIGN(long8,-) -GENERIC_OP_ASSIGN(ulong8,-) -GENERIC_OP_ASSIGN(float8,-) -GENERIC_OP_ASSIGN(double8,-) +ALL_TYPE(GENERIC_UNARY_OP,-) /* unary minus */ +ALL_TYPE(GENERIC_OP_ASSIGN,+) /* add assign */ +ALL_TYPE(GENERIC_OP_ASSIGN,-) /* subtract assign */ +ALL_TYPE(GENERIC_OP_ASSIGN,*) /* multiply assign */ +ALL_TYPE(GENERIC_OP_ASSIGN,/) /* divide assign */ #endif -/* multiply assign */ -#if GCC_VERSION < 40600 -EXPLICIT_OP_ASSIGN_VEC2(char2,*) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,*) -EXPLICIT_OP_ASSIGN_VEC2(int2,*) -EXPLICIT_OP_ASSIGN_VEC2(uint2,*) -EXPLICIT_OP_ASSIGN_VEC2(long2,*) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,*) -EXPLICIT_OP_ASSIGN_VEC2(float2,*) -EXPLICIT_OP_ASSIGN_VEC2(double2,*) -EXPLICIT_OP_ASSIGN_VEC4(char4,*) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,*) -EXPLICIT_OP_ASSIGN_VEC4(int4,*) -EXPLICIT_OP_ASSIGN_VEC4(uint4,*) -EXPLICIT_OP_ASSIGN_VEC4(long4,*) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,*) -EXPLICIT_OP_ASSIGN_VEC4(float4,*) -EXPLICIT_OP_ASSIGN_VEC4(double4,*) -EXPLICIT_OP_ASSIGN_VEC8(char8,*) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,*) -EXPLICIT_OP_ASSIGN_VEC8(int8,*) -EXPLICIT_OP_ASSIGN_VEC8(uint8,*) -EXPLICIT_OP_ASSIGN_VEC8(long8,*) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,*) -EXPLICIT_OP_ASSIGN_VEC8(float8,*) -EXPLICIT_OP_ASSIGN_VEC8(double8,*) -#else -GENERIC_OP_ASSIGN(char2,*) -GENERIC_OP_ASSIGN(uchar2,*) -GENERIC_OP_ASSIGN(int2,*) -GENERIC_OP_ASSIGN(uint2,*) -GENERIC_OP_ASSIGN(long2,*) -GENERIC_OP_ASSIGN(ulong2,*) -GENERIC_OP_ASSIGN(float2,*) -GENERIC_OP_ASSIGN(double2,*) -GENERIC_OP_ASSIGN(char4,*) -GENERIC_OP_ASSIGN(uchar4,*) -GENERIC_OP_ASSIGN(int4,*) -GENERIC_OP_ASSIGN(uint4,*) -GENERIC_OP_ASSIGN(long4,*) -GENERIC_OP_ASSIGN(ulong4,*) -GENERIC_OP_ASSIGN(float4,*) -GENERIC_OP_ASSIGN(double4,*) -GENERIC_OP_ASSIGN(char8,*) -GENERIC_OP_ASSIGN(uchar8,*) -GENERIC_OP_ASSIGN(int8,*) -GENERIC_OP_ASSIGN(uint8,*) -GENERIC_OP_ASSIGN(long8,*) -GENERIC_OP_ASSIGN(ulong8,*) -GENERIC_OP_ASSIGN(float8,*) -GENERIC_OP_ASSIGN(double8,*) -#endif - -/* modulo assign */ -EXPLICIT_OP_ASSIGN_VEC2(char2,%) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,%) -EXPLICIT_OP_ASSIGN_VEC2(int2,%) -EXPLICIT_OP_ASSIGN_VEC2(uint2,%) -EXPLICIT_OP_ASSIGN_VEC2(long2,%) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,%) -EXPLICIT_OP_ASSIGN_VEC4(char4,%) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,%) -EXPLICIT_OP_ASSIGN_VEC4(int4,%) -EXPLICIT_OP_ASSIGN_VEC4(uint4,%) -EXPLICIT_OP_ASSIGN_VEC4(long4,%) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,%) -EXPLICIT_OP_ASSIGN_VEC8(char8,%) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,%) -EXPLICIT_OP_ASSIGN_VEC8(int8,%) -EXPLICIT_OP_ASSIGN_VEC8(uint8,%) -EXPLICIT_OP_ASSIGN_VEC8(long8,%) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,%) - -/* divide assign */ -#if GCC_VERSION < 40600 -EXPLICIT_OP_ASSIGN_VEC2(char2,/) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,/) -EXPLICIT_OP_ASSIGN_VEC2(int2,/) -EXPLICIT_OP_ASSIGN_VEC2(uint2,/) -EXPLICIT_OP_ASSIGN_VEC2(long2,/) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,/) -EXPLICIT_OP_ASSIGN_VEC2(float2,/) -EXPLICIT_OP_ASSIGN_VEC2(double2,/) -EXPLICIT_OP_ASSIGN_VEC4(char4,/) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,/) -EXPLICIT_OP_ASSIGN_VEC4(int4,/) -EXPLICIT_OP_ASSIGN_VEC4(uint4,/) -EXPLICIT_OP_ASSIGN_VEC4(long4,/) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,/) -EXPLICIT_OP_ASSIGN_VEC4(float4,/) -EXPLICIT_OP_ASSIGN_VEC4(double4,/) -EXPLICIT_OP_ASSIGN_VEC8(char8,/) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,/) -EXPLICIT_OP_ASSIGN_VEC8(int8,/) -EXPLICIT_OP_ASSIGN_VEC8(uint8,/) -EXPLICIT_OP_ASSIGN_VEC8(long8,/) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,/) -EXPLICIT_OP_ASSIGN_VEC8(float8,/) -EXPLICIT_OP_ASSIGN_VEC8(double8,/) -#else -GENERIC_OP_ASSIGN(char2,/) -GENERIC_OP_ASSIGN(uchar2,/) -GENERIC_OP_ASSIGN(int2,/) -GENERIC_OP_ASSIGN(uint2,/) -GENERIC_OP_ASSIGN(long2,/) -GENERIC_OP_ASSIGN(ulong2,/) -GENERIC_OP_ASSIGN(float2,/) -GENERIC_OP_ASSIGN(double2,/) -GENERIC_OP_ASSIGN(char4,/) -GENERIC_OP_ASSIGN(uchar4,/) -GENERIC_OP_ASSIGN(int4,/) -GENERIC_OP_ASSIGN(uint4,/) -GENERIC_OP_ASSIGN(long4,/) -GENERIC_OP_ASSIGN(ulong4,/) -GENERIC_OP_ASSIGN(float4,/) -GENERIC_OP_ASSIGN(double4,/) -GENERIC_OP_ASSIGN(char8,/) -GENERIC_OP_ASSIGN(uchar8,/) -GENERIC_OP_ASSIGN(int8,/) -GENERIC_OP_ASSIGN(uint8,/) -GENERIC_OP_ASSIGN(long8,/) -GENERIC_OP_ASSIGN(ulong8,/) -GENERIC_OP_ASSIGN(float8,/) -GENERIC_OP_ASSIGN(double8,/) -#endif - -/* bitwise and assign */ -EXPLICIT_OP_ASSIGN_VEC2(char2,&) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,&) -EXPLICIT_OP_ASSIGN_VEC2(int2,&) -EXPLICIT_OP_ASSIGN_VEC2(uint2,&) -EXPLICIT_OP_ASSIGN_VEC2(long2,&) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,&) -EXPLICIT_OP_ASSIGN_VEC4(char4,&) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,&) -EXPLICIT_OP_ASSIGN_VEC4(int4,&) -EXPLICIT_OP_ASSIGN_VEC4(uint4,&) -EXPLICIT_OP_ASSIGN_VEC4(long4,&) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,&) -EXPLICIT_OP_ASSIGN_VEC8(char8,&) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,&) -EXPLICIT_OP_ASSIGN_VEC8(int8,&) -EXPLICIT_OP_ASSIGN_VEC8(uint8,&) -EXPLICIT_OP_ASSIGN_VEC8(long8,&) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,&) - -/* bitwise or assign */ -EXPLICIT_OP_ASSIGN_VEC2(char2,|) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,|) -EXPLICIT_OP_ASSIGN_VEC2(int2,|) -EXPLICIT_OP_ASSIGN_VEC2(uint2,|) -EXPLICIT_OP_ASSIGN_VEC2(long2,|) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,|) -EXPLICIT_OP_ASSIGN_VEC4(char4,|) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,|) -EXPLICIT_OP_ASSIGN_VEC4(int4,|) -EXPLICIT_OP_ASSIGN_VEC4(uint4,|) -EXPLICIT_OP_ASSIGN_VEC4(long4,|) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,|) -EXPLICIT_OP_ASSIGN_VEC8(char8,|) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,|) -EXPLICIT_OP_ASSIGN_VEC8(int8,|) -EXPLICIT_OP_ASSIGN_VEC8(uint8,|) -EXPLICIT_OP_ASSIGN_VEC8(long8,|) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,|) - -/* bitwise xor assign */ -EXPLICIT_OP_ASSIGN_VEC2(char2,^) -EXPLICIT_OP_ASSIGN_VEC2(uchar2,^) -EXPLICIT_OP_ASSIGN_VEC2(int2,^) -EXPLICIT_OP_ASSIGN_VEC2(uint2,^) -EXPLICIT_OP_ASSIGN_VEC2(long2,^) -EXPLICIT_OP_ASSIGN_VEC2(ulong2,^) -EXPLICIT_OP_ASSIGN_VEC4(char4,^) -EXPLICIT_OP_ASSIGN_VEC4(uchar4,^) -EXPLICIT_OP_ASSIGN_VEC4(int4,^) -EXPLICIT_OP_ASSIGN_VEC4(uint4,^) -EXPLICIT_OP_ASSIGN_VEC4(long4,^) -EXPLICIT_OP_ASSIGN_VEC4(ulong4,^) -EXPLICIT_OP_ASSIGN_VEC8(char8,^) -EXPLICIT_OP_ASSIGN_VEC8(uchar8,^) -EXPLICIT_OP_ASSIGN_VEC8(int8,^) -EXPLICIT_OP_ASSIGN_VEC8(uint8,^) -EXPLICIT_OP_ASSIGN_VEC8(long8,^) -EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) +INT_TYPE(EXPLICIT_UNARY_OP_ALL_TYPE_SIZE,~) /* unary bitwise not */ +INT_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,%) /* modulo assign */ +INT_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,&) /* bitwise and assign */ +INT_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,|) /* bitwise or assign */ +INT_TYPE(EXPLICIT_OP_ASSIGN_ALL_TYPE_SIZE,^) /* bitwise xor assign */ /*** other builtin data types [6.1.3] ***/ @@ -769,6 +622,8 @@ EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) #define convert_char(x) static_cast(x) #define convert_uchar(x) static_cast(x) +#define convert_short(x) static_cast(x) +#define convert_ushort(x) static_cast(x) #define convert_int(x) static_cast(x) #define convert_uint(x) static_cast(x) #define convert_long(x) static_cast(x) @@ -778,6 +633,8 @@ EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) #define convert_char2(x) static_cast(x) #define convert_uchar2(x) static_cast(x) +#define convert_short2(x) static_cast(x) +#define convert_ushort2(x) static_cast(x) #define convert_int2(x) static_cast(x) #define convert_uint2(x) static_cast(x) #define convert_long2(x) static_cast(x) @@ -787,6 +644,8 @@ EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) #define convert_char4(x) static_cast(x) #define convert_uchar4(x) static_cast(x) +#define convert_short4(x) static_cast(x) +#define convert_ushort4(x) static_cast(x) #define convert_int4(x) static_cast(x) #define convert_uint4(x) static_cast(x) #define convert_long4(x) static_cast(x) @@ -796,6 +655,8 @@ EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) #define convert_char8(x) static_cast(x) #define convert_uchar8(x) static_cast(x) +#define convert_short8(x) static_cast(x) +#define convert_ushort8(x) static_cast(x) #define convert_int8(x) static_cast(x) #define convert_uint8(x) static_cast(x) #define convert_long8(x) static_cast(x) @@ -803,21 +664,17 @@ EXPLICIT_OP_ASSIGN_VEC8(ulong8,^) #define convert_float8(x) static_cast(x) #define convert_double8(x) static_cast(x) +#define convert_char16(x) static_cast(x) +#define convert_uchar16(x) static_cast(x) +#define convert_short16(x) static_cast(x) +#define convert_ushort16(x) static_cast(x) +#define convert_int16(x) static_cast(x) +#define convert_uint16(x) static_cast(x) +#define convert_long16(x) static_cast(x) +#define convert_ulong16(x) static_cast(x) +#define convert_float16(x) static_cast(x) +#define convert_double16(x) static_cast(x) -/* -static __always_inline double as_double( float2 f2 ); -static __always_inline double as_double( float2 f2 ) - { return *(double*)(&f2); } - -template < typename T > -static __always_inline float2 as_float2( T x ) { return *(float2*)(&x); } - -template < typename T > -static __always_inline uint2 as_uint2( T x ) { return *(uint2*)(&x); } - -template < typename T > -static __always_inline int2 as_int2( T x ) { return *(int2*)(&x); } -*/ #define AS_TYPE(type) \ template < typename T > \ @@ -825,6 +682,8 @@ static __always_inline type as_##type( T x ) { return *(type*)(&x); } AS_TYPE(char2) AS_TYPE(uchar2) +AS_TYPE(short2) +AS_TYPE(ushort2) AS_TYPE(int2) AS_TYPE(uint2) AS_TYPE(long2) @@ -833,6 +692,8 @@ AS_TYPE(float2) AS_TYPE(double2) AS_TYPE(char4) AS_TYPE(uchar4) +AS_TYPE(short4) +AS_TYPE(ushort4) AS_TYPE(int4) AS_TYPE(uint4) AS_TYPE(long4) @@ -841,323 +702,229 @@ AS_TYPE(float4) AS_TYPE(double4) AS_TYPE(char8) AS_TYPE(uchar8) +AS_TYPE(short8) +AS_TYPE(ushort8) AS_TYPE(int8) AS_TYPE(uint8) AS_TYPE(long8) AS_TYPE(ulong8) AS_TYPE(float8) AS_TYPE(double8) +AS_TYPE(char16) +AS_TYPE(uchar16) +AS_TYPE(short16) +AS_TYPE(ushort16) +AS_TYPE(int16) +AS_TYPE(uint16) +AS_TYPE(long16) +AS_TYPE(ulong16) +AS_TYPE(float16) +AS_TYPE(double16) /*** *** operators for vector data types [6.3] ***/ -#define GENERIC_BINOP_VEC2(type,vectype,op) \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ) \ - { return (_##vectype){ a.vec op b.vec }; } \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ) \ - { return (_##vectype){ __vector_##vectype(a,a) op b.vec }; } \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ) \ - { return (_##vectype){ a.vec op __vector_##vectype(b,b) }; } - -#define GENERIC_BINOP_VEC4(type,vectype,op) \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ) \ - { return (_##vectype){ a.vec op b.vec }; } \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ) \ - { return (_##vectype){ __vector_##vectype(a,a,a,a) op b.vec }; } \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ) \ - { return (_##vectype){ a.vec op __vector_##vectype(b,b,b,b) }; } - -#define EXPLICIT_BINOP_VEC2(type,vectype,op) \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ) \ +#define GENERIC_BINOP_ALL_TYPE_SIZE(vectype,type,op) \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, const _##vectype##2 b ); \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, const _##vectype##2 b ) \ + { return (_##vectype##2){ a.vec op b.vec }; } \ +static __always_inline _##vectype##2 \ + operator op ( type a, const _##vectype##2 b ); \ +static __always_inline _##vectype##2 \ + operator op ( type a, const _##vectype##2 b ) \ + { return (_##vectype##2){ __vector_##vectype##2(a,a) op b.vec }; } \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, type b ); \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, type b ) \ + { return (_##vectype##2){ a.vec op __vector_##vectype##2(b,b) }; } \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, const _##vectype##4 b ); \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, const _##vectype##4 b ) \ + { return (_##vectype##4){ a.vec op b.vec }; } \ +static __always_inline _##vectype##4 \ + operator op ( type a, const _##vectype##4 b ); \ +static __always_inline _##vectype##4 \ + operator op ( type a, const _##vectype##4 b ) \ + { return (_##vectype##4){ __vector_##vectype##4(a,a,a,a) op b.vec }; } \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, type b ); \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, type b ) \ + { return (_##vectype##4){ a.vec op __vector_##vectype##4(b,b,b,b) }; } \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, const _##vectype##8 b ); \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, const _##vectype##8 b ) \ + { return (_##vectype##8){ a.vec op b.vec }; } \ +static __always_inline _##vectype##8 \ + operator op ( type a, const _##vectype##8 b ); \ +static __always_inline _##vectype##8 \ + operator op ( type a, const _##vectype##8 b ) \ + { return (_##vectype##8){ __vector_##vectype##8(a,a,a,a,a,a,a,a) op b.vec }; } \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, type b ); \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, type b ) \ + { return (_##vectype##8){ a.vec op __vector_##vectype##8(b,b,b,b,b,b,b,b) }; } \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, const _##vectype##16 b ); \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, const _##vectype##16 b ) \ + { return (_##vectype##16){ a.vec op b.vec }; } \ +static __always_inline _##vectype##16 \ + operator op ( type a, const _##vectype##16 b ); \ +static __always_inline _##vectype##16 \ + operator op ( type a, const _##vectype##16 b ) \ + { return (_##vectype##16){ __vector_##vectype##16(a,a,a,a,a,a,a,a,a,a,a,a,a,a,a,a) op b.vec }; } \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, type b ); \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, type b ) \ + { return (_##vectype##16){ a.vec op __vector_##vectype##16(b,b,b,b,b,b,b,b,b,b,b,b,b,b,b,b) }; } + +#define EXPLICIT_BINOP_ALL_TYPE_SIZE(vectype,type,op) \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, const _##vectype##2 b ); \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, const _##vectype##2 b ) \ { \ - _##vectype tmp = { a.s0 op b.s0, a.s1 op b.s1 }; \ + _##vectype##2 tmp = { a.s0 op b.s0, a.s1 op b.s1 }; \ return tmp; \ } \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ) \ +static __always_inline _##vectype##2 \ + operator op ( type a, const _##vectype##2 b ); \ +static __always_inline _##vectype##2 \ + operator op ( type a, const _##vectype##2 b ) \ { \ - _##vectype tmp = { a op b.s0, a op b.s1 }; \ + _##vectype##2 tmp = { a op b.s0, a op b.s1 }; \ return tmp; \ } \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ) \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, type b ); \ +static __always_inline _##vectype##2 \ + operator op ( const _##vectype##2 a, type b ) \ { \ - _##vectype tmp = { a.s0 op b, a.s1 op b }; \ + _##vectype##2 tmp = { a.s0 op b, a.s1 op b }; \ return tmp; \ } \ - -#define EXPLICIT_BINOP_VEC4(type,vectype,op) \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ); \ -static __always_inline _##vectype \ - operator op ( const _##vectype a, const _##vectype b ) \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, const _##vectype##4 b ); \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, const _##vectype##4 b ) \ + { \ + _##vectype##4 tmp={a.s0 op b.s0,a.s1 op b.s1,a.s2 op b.s2,a.s3 op b.s3}; \ + return tmp; \ + } \ +static __always_inline _##vectype##4 operator op ( type a, const _##vectype##4 b );\ +static __always_inline _##vectype##4 \ + operator op ( type a, const _##vectype##4 b ) \ + { \ + _##vectype##4 tmp={a op b.s0,a op b.s1,a op b.s2,a op b.s3}; \ + return tmp; \ + } \ +static __always_inline _##vectype##4 operator op ( const _##vectype##4 a, type b ); \ +static __always_inline _##vectype##4 \ + operator op ( const _##vectype##4 a, type b ) \ + { \ + _##vectype##4 tmp={a.s0 op b,a.s1 op b,a.s2 op b,a.s3 op b}; \ + return tmp; \ + } \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, const _##vectype##8 b ); \ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, const _##vectype##8 b ) \ + { \ + _##vectype##8 tmp={a.s0 op b.s0,a.s1 op b.s1,a.s2 op b.s2,a.s3 op b.s3, \ + a.s4 op b.s4,a.s5 op b.s5,a.s6 op b.s6,a.s7 op b.s7}; \ + return tmp; \ + } \ +static __always_inline _##vectype##8 operator op ( type a, const _##vectype##8 b );\ +static __always_inline _##vectype##8 \ + operator op ( type a, const _##vectype##8 b ) \ + { \ + _##vectype##8 tmp={a op b.s0,a op b.s1,a op b.s2,a op b.s3, \ + a op b.s4,a op b.s5,a op b.s6,a op b.s7}; \ + return tmp; \ + } \ +static __always_inline _##vectype##8 operator op ( const _##vectype##8 a, type b );\ +static __always_inline _##vectype##8 \ + operator op ( const _##vectype##8 a, type b ) \ { \ - _##vectype tmp={a.s0 op b.s0,a.s1 op b.s1,a.s2 op b.s2,a.s3 op b.s3}; \ + _##vectype##8 tmp={a.s0 op b,a.s1 op b,a.s2 op b,a.s3 op b, \ + a.s4 op b,a.s5 op b,a.s6 op b,a.s7 op b}; \ return tmp; \ } \ -static __always_inline _##vectype operator op ( type a, const _##vectype b );\ -static __always_inline _##vectype \ - operator op ( type a, const _##vectype b ) \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, const _##vectype##16 b ); \ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, const _##vectype##16 b ) \ { \ - _##vectype tmp={a op b.s0,a op b.s1,a op b.s2,a op b.s3}; \ + _##vectype##16 tmp={a.s0 op b.s0,a.s1 op b.s1,a.s2 op b.s2,a.s3 op b.s3, \ + a.s4 op b.s4,a.s5 op b.s5,a.s6 op b.s6,a.s7 op b.s7, \ + a.s8 op b.s8,a.s9 op b.s9,a.sA op b.sA,a.sB op b.sB, \ + a.sC op b.sC,a.sD op b.sD,a.sE op b.sE,a.sF op b.sF}; \ + return tmp; \ + } \ +static __always_inline _##vectype##16 operator op ( type a, const _##vectype##16 b );\ +static __always_inline _##vectype##16 \ + operator op ( type a, const _##vectype##16 b ) \ + { \ + _##vectype##16 tmp={a op b.s0,a op b.s1,a op b.s2,a op b.s3, \ + a op b.s4,a op b.s5,a op b.s6,a op b.s7, \ + a op b.s8,a op b.s9,a op b.sA,a op b.sB, \ + a op b.sC,a op b.sD,a op b.sE,a op b.sF}; \ return tmp; \ }\ -static __always_inline _##vectype operator op ( const _##vectype a, type b );\ -static __always_inline _##vectype \ - operator op ( const _##vectype a, type b ) \ +static __always_inline _##vectype##16 operator op ( const _##vectype##16 a, type b );\ +static __always_inline _##vectype##16 \ + operator op ( const _##vectype##16 a, type b ) \ { \ - _##vectype tmp={a.s0 op b,a.s1 op b,a.s2 op b,a.s3 op b}; \ + _##vectype##16 tmp={a.s0 op b,a.s1 op b,a.s2 op b,a.s3 op b, \ + a.s4 op b,a.s5 op b,a.s6 op b,a.s7 op b, \ + a.s8 op b,a.s9 op b,a.sA op b,a.sB op b, \ + a.sC op b,a.sD op b,a.sE op b,a.sF op b}; \ return tmp; \ } +#define INT_TYPE2(m,...) \ +m(char,char,__VA_ARGS__) \ +m(uchar,unsigned char,__VA_ARGS__) \ +m(short,short,__VA_ARGS__) \ +m(ushort,unsigned short,__VA_ARGS__) \ +m(int,int,__VA_ARGS__) \ +m(uint,unsigned int,__VA_ARGS__) \ +m(long,long long int,__VA_ARGS__) \ +m(ulong,unsigned long long int,__VA_ARGS__) + +#define ALL_TYPE2(m,...) \ +INT_TYPE2(m,__VA_ARGS__) \ +m(float,float,__VA_ARGS__) \ +m(double,float,__VA_ARGS__) #if GCC_VERSION < 40600 - -/* binary add */ -EXPLICIT_BINOP_VEC2(char,char2,+) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,+) -EXPLICIT_BINOP_VEC2(int,int2,+) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,+) -EXPLICIT_BINOP_VEC2(long long,long2,+) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,+) -EXPLICIT_BINOP_VEC2(float,float2,+) -EXPLICIT_BINOP_VEC2(double,double2,+) -EXPLICIT_BINOP_VEC4(char,char4,+) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,+) -EXPLICIT_BINOP_VEC4(int,int4,+) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,+) -EXPLICIT_BINOP_VEC4(float,float4,+) - -/* binary sub */ -EXPLICIT_BINOP_VEC2(char,char2,-) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,-) -EXPLICIT_BINOP_VEC2(int,int2,-) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,-) -EXPLICIT_BINOP_VEC2(long long,long2,-) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,-) -EXPLICIT_BINOP_VEC2(float,float2,-) -EXPLICIT_BINOP_VEC2(double,double2,-) -EXPLICIT_BINOP_VEC4(char,char4,-) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,-) -EXPLICIT_BINOP_VEC4(int,int4,-) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,-) -EXPLICIT_BINOP_VEC4(float,float4,-) - -/* binary mult */ -EXPLICIT_BINOP_VEC2(char,char2,*) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,*) -EXPLICIT_BINOP_VEC2(int,int2,*) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,*) -EXPLICIT_BINOP_VEC2(long long,long2,*) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,*) -EXPLICIT_BINOP_VEC2(float,float2,*) -EXPLICIT_BINOP_VEC2(double,double2,*) -EXPLICIT_BINOP_VEC4(char,char4,*) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,*) -EXPLICIT_BINOP_VEC4(int,int4,*) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,*) -EXPLICIT_BINOP_VEC4(float,float4,*) - -/* binary mod */ -EXPLICIT_BINOP_VEC2(char,char2,%) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,%) -EXPLICIT_BINOP_VEC2(int,int2,%) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,%) -EXPLICIT_BINOP_VEC2(long long,long2,%) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,%) -EXPLICIT_BINOP_VEC4(char,char4,%) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,%) -EXPLICIT_BINOP_VEC4(int,int4,%) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,%) - -/* binary div */ -EXPLICIT_BINOP_VEC2(char,char2,/) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,/) -EXPLICIT_BINOP_VEC2(int,int2,/) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,/) -EXPLICIT_BINOP_VEC2(long long,long2,/) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,/) -EXPLICIT_BINOP_VEC2(float,float2,/) -EXPLICIT_BINOP_VEC2(double,double2,/) -EXPLICIT_BINOP_VEC4(char,char4,/) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,/) -EXPLICIT_BINOP_VEC4(int,int4,/) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,/) -EXPLICIT_BINOP_VEC4(float,float4,/) - -/* binary and */ -EXPLICIT_BINOP_VEC2(char,char2,&) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,&) -EXPLICIT_BINOP_VEC2(int,int2,&) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,&) -EXPLICIT_BINOP_VEC2(long long,long2,&) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,&) -EXPLICIT_BINOP_VEC4(char,char4,&) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,&) -EXPLICIT_BINOP_VEC4(int,int4,&) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,&) - -/* binary or */ -EXPLICIT_BINOP_VEC2(char,char2,|) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,|) -EXPLICIT_BINOP_VEC2(int,int2,|) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,|) -EXPLICIT_BINOP_VEC2(long long,long2,|) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,|) -EXPLICIT_BINOP_VEC4(char,char4,|) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,|) -EXPLICIT_BINOP_VEC4(int,int4,|) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,|) - -/* binary xor */ -EXPLICIT_BINOP_VEC2(char,char2,^) -EXPLICIT_BINOP_VEC2(unsigned char,uchar2,^) -EXPLICIT_BINOP_VEC2(int,int2,^) -EXPLICIT_BINOP_VEC2(unsigned int,uint2,^) -EXPLICIT_BINOP_VEC2(long long,long2,^) -EXPLICIT_BINOP_VEC2(unsigned long long,ulong2,^) -EXPLICIT_BINOP_VEC4(char,char4,^) -EXPLICIT_BINOP_VEC4(unsigned char,uchar4,^) -EXPLICIT_BINOP_VEC4(int,int4,^) -EXPLICIT_BINOP_VEC4(unsigned int,uint4,^) - +ALL_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,+) /* binary add */ +ALL_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,-) /* binary sub */ +ALL_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,*) /* binary mult */ +INT_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,%) /* binary mod */ +ALL_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,/) /* binary div */ +INT_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,&) /* binary and */ +INT_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,|) /* binary or */ +INT_TYPE2(EXPLICIT_BINOP_ALL_TYPE_SIZE,^) /* binary xor */ #else - -/* binary add */ -GENERIC_BINOP_VEC2(char,char2,+) -GENERIC_BINOP_VEC2(unsigned char,uchar2,+) -GENERIC_BINOP_VEC2(int,int2,+) -GENERIC_BINOP_VEC2(unsigned int,uint2,+) -GENERIC_BINOP_VEC2(long long,long2,+) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,+) -GENERIC_BINOP_VEC2(float,float2,+) -GENERIC_BINOP_VEC2(double,double2,+) -GENERIC_BINOP_VEC4(char,char4,+) -GENERIC_BINOP_VEC4(unsigned char,uchar4,+) -GENERIC_BINOP_VEC4(int,int4,+) -GENERIC_BINOP_VEC4(unsigned int,uint4,+) -GENERIC_BINOP_VEC4(float,float4,+) - -/* binary sub */ -GENERIC_BINOP_VEC2(char,char2,-) -GENERIC_BINOP_VEC2(unsigned char,uchar2,-) -GENERIC_BINOP_VEC2(int,int2,-) -GENERIC_BINOP_VEC2(unsigned int,uint2,-) -GENERIC_BINOP_VEC2(long long,long2,-) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,-) -GENERIC_BINOP_VEC2(float,float2,-) -GENERIC_BINOP_VEC2(double,double2,-) -GENERIC_BINOP_VEC4(char,char4,-) -GENERIC_BINOP_VEC4(unsigned char,uchar4,-) -GENERIC_BINOP_VEC4(int,int4,-) -GENERIC_BINOP_VEC4(unsigned int,uint4,-) -GENERIC_BINOP_VEC4(float,float4,-) - -/* binary mult */ -GENERIC_BINOP_VEC2(char,char2,*) -GENERIC_BINOP_VEC2(unsigned char,uchar2,*) -GENERIC_BINOP_VEC2(int,int2,*) -GENERIC_BINOP_VEC2(unsigned int,uint2,*) -GENERIC_BINOP_VEC2(long long,long2,*) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,*) -GENERIC_BINOP_VEC2(float,float2,*) -GENERIC_BINOP_VEC2(double,double2,*) -GENERIC_BINOP_VEC4(char,char4,*) -GENERIC_BINOP_VEC4(unsigned char,uchar4,*) -GENERIC_BINOP_VEC4(int,int4,*) -GENERIC_BINOP_VEC4(unsigned int,uint4,*) -GENERIC_BINOP_VEC4(float,float4,*) - -/* binary mod */ -GENERIC_BINOP_VEC2(char,char2,%) -GENERIC_BINOP_VEC2(unsigned char,uchar2,%) -GENERIC_BINOP_VEC2(int,int2,%) -GENERIC_BINOP_VEC2(unsigned int,uint2,%) -GENERIC_BINOP_VEC2(long long,long2,%) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,%) -GENERIC_BINOP_VEC4(char,char4,%) -GENERIC_BINOP_VEC4(unsigned char,uchar4,%) -GENERIC_BINOP_VEC4(int,int4,%) -GENERIC_BINOP_VEC4(unsigned int,uint4,%) - -/* binary div */ -GENERIC_BINOP_VEC2(char,char2,/) -GENERIC_BINOP_VEC2(unsigned char,uchar2,/) -GENERIC_BINOP_VEC2(int,int2,/) -GENERIC_BINOP_VEC2(unsigned int,uint2,/) -GENERIC_BINOP_VEC2(long long,long2,/) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,/) -GENERIC_BINOP_VEC2(float,float2,/) -GENERIC_BINOP_VEC2(double,double2,/) -GENERIC_BINOP_VEC4(char,char4,/) -GENERIC_BINOP_VEC4(unsigned char,uchar4,/) -GENERIC_BINOP_VEC4(int,int4,/) -GENERIC_BINOP_VEC4(unsigned int,uint4,/) -GENERIC_BINOP_VEC4(float,float4,/) - -/* binary and */ -GENERIC_BINOP_VEC2(char,char2,&) -GENERIC_BINOP_VEC2(unsigned char,uchar2,&) -GENERIC_BINOP_VEC2(int,int2,&) -GENERIC_BINOP_VEC2(unsigned int,uint2,&) -GENERIC_BINOP_VEC2(long long,long2,&) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,&) -GENERIC_BINOP_VEC4(char,char4,&) -GENERIC_BINOP_VEC4(unsigned char,uchar4,&) -GENERIC_BINOP_VEC4(int,int4,&) -GENERIC_BINOP_VEC4(unsigned int,uint4,&) - -/* binary or */ -GENERIC_BINOP_VEC2(char,char2,|) -GENERIC_BINOP_VEC2(unsigned char,uchar2,|) -GENERIC_BINOP_VEC2(int,int2,|) -GENERIC_BINOP_VEC2(unsigned int,uint2,|) -GENERIC_BINOP_VEC2(long long,long2,|) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,|) -GENERIC_BINOP_VEC4(char,char4,|) -GENERIC_BINOP_VEC4(unsigned char,uchar4,|) -GENERIC_BINOP_VEC4(int,int4,|) -GENERIC_BINOP_VEC4(unsigned int,uint4,|) - -/* binary xor */ -GENERIC_BINOP_VEC2(char,char2,^) -GENERIC_BINOP_VEC2(unsigned char,uchar2,^) -GENERIC_BINOP_VEC2(int,int2,^) -GENERIC_BINOP_VEC2(unsigned int,uint2,^) -GENERIC_BINOP_VEC2(long long,long2,^) -GENERIC_BINOP_VEC2(unsigned long long,ulong2,^) -GENERIC_BINOP_VEC4(char,char4,^) -GENERIC_BINOP_VEC4(unsigned char,uchar4,^) -GENERIC_BINOP_VEC4(int,int4,^) -GENERIC_BINOP_VEC4(unsigned int,uint4,^) - +ALL_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,+) /* binary add */ +ALL_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,-) /* binary sub */ +ALL_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,*) /* binary mult */ +INT_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,%) /* binary mod */ +ALL_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,/) /* binary div */ +INT_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,&) /* binary and */ +INT_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,|) /* binary or */ +INT_TYPE2(GENERIC_BINOP_ALL_TYPE_SIZE,^) /* binary xor */ #endif @@ -1182,23 +949,203 @@ static __always_inline float4 func##_T( float4 a ) \ return vector_float4( \ func##f(a.x),func##f(a.y),func##f(a.z),func##f(a.w)); \ } \ +static __always_inline float8 func##_T( float8 a ); \ +static __always_inline float8 func##_T( float8 a ) \ + { \ + return vector_float8( \ + func##f(a.s0),func##f(a.s1),func##f(a.s2),func##f(a.s3), \ + func##f(a.s4),func##f(a.s5),func##f(a.s6),func##f(a.s7)); \ + } \ +static __always_inline float16 func##_T( float16 a ); \ +static __always_inline float16 func##_T( float16 a ) \ + { \ + return vector_float16( \ + func##f(a.s0),func##f(a.s1),func##f(a.s2),func##f(a.s3), \ + func##f(a.s4),func##f(a.s5),func##f(a.s6),func##f(a.s7), \ + func##f(a.s8),func##f(a.s9),func##f(a.sA),func##f(a.sB), \ + func##f(a.sC),func##f(a.sD),func##f(a.sE),func##f(a.sF)); \ + } \ static __always_inline double func##_T( double a ); \ -static __always_inline double func##_T( double a ) { return func(a); } \ +static __always_inline double func##_T( double a ) \ + { return func(a); } \ static __always_inline double2 func##_T( double2 a ); \ static __always_inline double2 func##_T( double2 a ) \ - { return vector_double2( func(a.x),func(a.y) ); } + { return vector_double2( func(a.x),func(a.y) ); } \ +static __always_inline double4 func##_T( double4 a ); \ +static __always_inline double4 func##_T( double4 a ) \ + { \ + return vector_double4( \ + func(a.x),func(a.y),func(a.z),func(a.w)); \ + } \ +static __always_inline double8 func##_T( double8 a ) \ + { \ + return vector_double8( \ + func(a.s0),func(a.s1),func(a.s2),func(a.s3), \ + func(a.s4),func(a.s5),func(a.s6),func(a.s7)); \ + } \ +static __always_inline double16 func##_T( double16 a ) \ + { \ + return vector_double16( \ + func(a.s0),func(a.s1),func(a.s2),func(a.s3), \ + func(a.s4),func(a.s5),func(a.s6),func(a.s7), \ + func(a.s8),func(a.s9),func(a.sA),func(a.sB), \ + func(a.sC),func(a.sD),func(a.sE),func(a.sF)); \ + } + +#define __MATH_BUILTIN_2(func) \ +static __always_inline float func##_T( float a, float b ); \ +static __always_inline float func##_T( float a, float b ) \ + { return func##f(a,b); } \ +static __always_inline float2 func##_T( float2 a, float2 b ); \ +static __always_inline float2 func##_T( float2 a, float2 b ) \ + { return vector_float2( func##f(a.s0,b.s0),func##f(a.s1,b.s1) ); } \ +static __always_inline float4 func##_T( float4 a, float4 b ); \ +static __always_inline float4 func##_T( float4 a, float4 b ) \ + { \ + return vector_float4( \ + func##f(a.s0,b.s0),func##f(a.s1,b.s1),func##f(a.s2,b.s2),func##f(a.s3,b.s3)); \ + } \ +static __always_inline float8 func##_T( float8 a, float8 b ); \ +static __always_inline float8 func##_T( float8 a, float8 b ) \ + { \ + return vector_float8( \ + func##f(a.s0,b.s0),func##f(a.s1,b.s1),func##f(a.s2,b.s2),func##f(a.s3,b.s3), \ + func##f(a.s4,b.s4),func##f(a.s5,b.s5),func##f(a.s6,b.s6),func##f(a.s7,b.s7)); \ + } \ +static __always_inline float16 func##_T( float16 a, float16 b ); \ +static __always_inline float16 func##_T( float16 a, float16 b ) \ + { \ + return vector_float16( \ + func##f(a.s0,b.s0),func##f(a.s1,b.s1),func##f(a.s2,b.s2),func##f(a.s3,b.s3), \ + func##f(a.s4,b.s4),func##f(a.s5,b.s5),func##f(a.s6,b.s6),func##f(a.s7,b.s7), \ + func##f(a.s8,b.s8),func##f(a.s9,b.s9),func##f(a.sA,b.sA),func##f(a.sB,b.sB), \ + func##f(a.sC,b.sC),func##f(a.sD,b.sD),func##f(a.sE,b.sE),func##f(a.sF,b.sF)); \ + } \ +static __always_inline double func##_T( double a, double b ); \ +static __always_inline double func##_T( double a, double b ) \ + { return func(a,b); } \ +static __always_inline double2 func##_T( double2 a, double2 b ); \ +static __always_inline double2 func##_T( double2 a, double2 b ) \ + { return vector_double2( func(a.s0,b.s0),func(a.s1,b.s1) ); } \ +static __always_inline double4 func##_T( double4 a, double4 b ); \ +static __always_inline double4 func##_T( double4 a, double4 b ) \ + { \ + return vector_double4( \ + func(a.s0,b.s0),func(a.s1,b.s1),func(a.s2,b.s2),func(a.s3,b.s3)); \ + } \ +static __always_inline double8 func##_T( double8 a, double8 b ); \ +static __always_inline double8 func##_T( double8 a, double8 b ) \ + { \ + return vector_double8( \ + func(a.s0,b.s0),func(a.s1,b.s1),func(a.s2,b.s2),func(a.s3,b.s3), \ + func(a.s4,b.s4),func(a.s5,b.s5),func(a.s6,b.s6),func(a.s7,b.s7)); \ + } \ +static __always_inline double16 func##_T( double16 a, double16 b ); \ +static __always_inline double16 func##_T( double16 a, double16 b ) \ + { \ + return vector_double16( \ + func(a.s0,b.s0),func(a.s1,b.s1),func(a.s2,b.s2),func(a.s3,b.s3), \ + func(a.s4,b.s4),func(a.s5,b.s5),func(a.s6,b.s6),func(a.s7,b.s7), \ + func(a.s8,b.s8),func(a.s9,b.s9),func(a.sA,b.sA),func(a.sB,b.sB), \ + func(a.sC,b.sC),func(a.sD,b.sD),func(a.sE,b.sE),func(a.sF,b.sF)); \ + } + +#define __MATH_BUILTIN_3(func) \ +static __always_inline float func##_T( float a, float b, float c); \ +static __always_inline float func##_T( float a, float b, float c ) \ + { return func##f(a,b,c); } \ +static __always_inline float2 func##_T( float2 a, float2 b, float2 c ); \ +static __always_inline float2 func##_T( float2 a, float2 b, float2 c ) \ + { return vector_float2( func##f(a.s0,b.s0,c.s0),func##f(a.s1,b.s1,c.s1) ); } \ +static __always_inline float4 func##_T( float4 a, float4 b, float4 c ); \ +static __always_inline float4 func##_T( float4 a, float4 b, float4 c ) \ + { \ + return vector_float4( \ + func##f(a.s0,b.s0,c.s0),func##f(a.s1,b.s1,c.s1),func##f(a.s2,b.s2,c.s2),func##f(a.s3,b.s3,c.s3)); \ + } \ +static __always_inline float8 func##_T( float8 a, float8 b, float8 c ); \ +static __always_inline float8 func##_T( float8 a, float8 b, float8 c ) \ + { \ + return vector_float8( \ + func##f(a.s0,b.s0,c.s0),func##f(a.s1,b.s1,c.s1),func##f(a.s2,b.s2,c.s2),func##f(a.s3,b.s3,c.s3), \ + func##f(a.s4,b.s4,c.s4),func##f(a.s5,b.s5,c.s5),func##f(a.s6,b.s6,c.s6),func##f(a.s7,b.s7,c.s7)); \ + } \ +static __always_inline float16 func##_T( float16 a, float16 b, float16 c ); \ +static __always_inline float16 func##_T( float16 a, float16 b, float16 c ) \ + { \ + return vector_float16( \ + func##f(a.s0,b.s0,c.s0),func##f(a.s1,b.s1,c.s1),func##f(a.s2,b.s2,c.s2),func##f(a.s3,b.s3,c.s3), \ + func##f(a.s4,b.s4,c.s4),func##f(a.s5,b.s5,c.s5),func##f(a.s6,b.s6,c.s6),func##f(a.s7,b.s7,c.s7), \ + func##f(a.s8,b.s8,c.s8),func##f(a.s9,b.s9,c.s9),func##f(a.sA,b.sA,c.sA),func##f(a.sB,b.sB,c.sB), \ + func##f(a.sC,b.sC,c.sC),func##f(a.sD,b.sD,c.sD),func##f(a.sE,b.sE,c.sE),func##f(a.sF,b.sF,c.sF)); \ + } \ +static __always_inline double func##_T( double a, double b, double c ); \ +static __always_inline double func##_T( double a, double b, double c ) \ + { return func(a,b,c); } \ +static __always_inline double2 func##_T( double2 a, double2 b, double2 c ); \ +static __always_inline double2 func##_T( double2 a, double2 b, double2 c ) \ + { return vector_double2( func(a.s0,b.s0,c.s0),func(a.s1,b.s1,c.s1) ); } \ +static __always_inline double4 func##_T( double4 a, double4 b, double4 c ); \ +static __always_inline double4 func##_T( double4 a, double4 b, double4 c ) \ + { \ + return vector_double4( \ + func(a.s0,b.s0,c.s0),func(a.s1,b.s1,c.s1),func(a.s2,b.s2,c.s2),func(a.s3,b.s3,c.s3)); \ + } \ +static __always_inline double8 func##_T( double8 a, double8 b, double8 c ); \ +static __always_inline double8 func##_T( double8 a, double8 b, double8 c ) \ + { \ + return vector_double8( \ + func(a.s0,b.s0,c.s0),func(a.s1,b.s1,c.s1),func(a.s2,b.s2,c.s2),func(a.s3,b.s3,c.s3), \ + func(a.s4,b.s4,c.s4),func(a.s5,b.s5,c.s5),func(a.s6,b.s6,c.s6),func(a.s7,b.s7,c.s7)); \ + } \ +static __always_inline double16 func##_T( double16 a, double16 b, double16 c ); \ +static __always_inline double16 func##_T( double16 a, double16 b, double16 c ) \ + { \ + return vector_double16( \ + func(a.s0,b.s0,c.s0),func(a.s1,b.s1,c.s1),func(a.s2,b.s2,c.s2),func(a.s3,b.s3,c.s3), \ + func(a.s4,b.s4,c.s4),func(a.s5,b.s5,c.s5),func(a.s6,b.s6,c.s6),func(a.s7,b.s7,c.s7), \ + func(a.s8,b.s8,c.s8),func(a.s9,b.s9,c.s9),func(a.sA,b.sA,c.sA),func(a.sB,b.sB,c.sB), \ + func(a.sC,b.sC,c.sC),func(a.sD,b.sD,c.sD),func(a.sE,b.sE,c.sE),func(a.sF,b.sF,c.sF)); \ + } __MATH_BUILTIN_1(acos) __MATH_BUILTIN_1(acosh) +static __always_inline float acospif( float a ); +static __always_inline float acospif( float a ) { return acosf(a)*M_1_PI_F; } +static __always_inline double acospi( double a ); +static __always_inline double acospi( double a ) { return acos(a)*M_1_PI; } +__MATH_BUILTIN_1(acospi) __MATH_BUILTIN_1(asin) __MATH_BUILTIN_1(asinh) +static __always_inline float asinpif( float a ); +static __always_inline float asinpif( float a ) { return asinf(a)*M_1_PI_F; } +static __always_inline double asinpi( double a ); +static __always_inline double asinpi( double a ) { return asin(a)*M_1_PI; } +__MATH_BUILTIN_1(asinpi) __MATH_BUILTIN_1(atan) +__MATH_BUILTIN_2(atan2) __MATH_BUILTIN_1(atanh) +static __always_inline float atanpif( float a ); +static __always_inline float atanpif( float a ) { return atanf(a)*M_1_PI_F; } +static __always_inline double atanpi( double a ); +static __always_inline double atanpi( double a ) { return atan(a)*M_1_PI; } +__MATH_BUILTIN_1(atanpi) +static __always_inline float atan2pif( float a, float b ); +static __always_inline float atan2pif( float a, float b ) { return atan2f(a,b)*M_1_PI_F; } +static __always_inline double atan2pi( double a, double b ); +static __always_inline double atan2pi( double a, double b ) { return atan2(a,b)*M_1_PI; } +__MATH_BUILTIN_2(atan2pi) __MATH_BUILTIN_1(cbrt) __MATH_BUILTIN_1(ceil) +__MATH_BUILTIN_2(copysign) __MATH_BUILTIN_1(cos) __MATH_BUILTIN_1(cosh) +static __always_inline float cospif( float a ); +static __always_inline float cospif( float a ) { return cosf(a*M_PI_F); } +static __always_inline double cospi( double a ); +static __always_inline double cospi( double a ) { return cos(a*M_PI); } +__MATH_BUILTIN_1(cospi) __MATH_BUILTIN_1(erfc) __MATH_BUILTIN_1(erf) __MATH_BUILTIN_1(exp) @@ -1210,7 +1157,70 @@ __MATH_BUILTIN_1(exp10) #endif __MATH_BUILTIN_1(expm1) __MATH_BUILTIN_1(fabs) +__MATH_BUILTIN_2(fdim) __MATH_BUILTIN_1(floor) +__MATH_BUILTIN_3(fma) +__MATH_BUILTIN_2(fmax); +static __always_inline float2 fmax_T( float2 a, float b ); +static __always_inline float2 fmax_T( float2 a, float b ) { return fmax_T(a,vector_float2(b,b)); } +static __always_inline float4 fmax_T( float4 a, float b ); +static __always_inline float4 fmax_T( float4 a, float b ) { return fmax_T(a,vector_float4(b,b,b,b)); } +static __always_inline float8 fmax_T( float8 a, float b ); +static __always_inline float8 fmax_T( float8 a, float b ) { return fmax_T(a,vector_float8(b,b,b,b,b,b,b,b)); } +static __always_inline float16 fmax_T( float16 a, float b ); +static __always_inline float16 fmax_T( float16 a, float b ) { return fmax_T(a,vector_float16(b,b,b,b,b,b,b,b,b,b,b,b,b,b,b,b)); } +static __always_inline double2 fmax_T( double2 a, double b ); +static __always_inline double2 fmax_T( double2 a, double b ) { return fmax_T(a,vector_double2(b,b)); } +static __always_inline double4 fmax_T( double4 a, double b ); +static __always_inline double4 fmax_T( double4 a, double b ) { return fmax_T(a,vector_double4(b,b,b,b)); } +static __always_inline double8 fmax_T( double8 a, double b ); +static __always_inline double8 fmax_T( double8 a, double b ) { return fmax_T(a,vector_double8(b,b,b,b,b,b,b,b)); } +static __always_inline double16 fmax_T( double16 a, double b ); +static __always_inline double16 fmax_T( double16 a, double b ) { return fmax_T(a,vector_double16(b,b,b,b,b,b,b,b,b,b,b,b,b,b,b,b)); } +__MATH_BUILTIN_2(fmin); +static __always_inline float2 fmin_T( float2 a, float b ); +static __always_inline float2 fmin_T( float2 a, float b ) { return fmin_T(a,vector_float2(b,b)); } +static __always_inline float4 fmin_T( float4 a, float b ); +static __always_inline float4 fmin_T( float4 a, float b ) { return fmin_T(a,vector_float4(b,b,b,b)); } +static __always_inline float8 fmin_T( float8 a, float b ); +static __always_inline float8 fmin_T( float8 a, float b ) { return fmin_T(a,vector_float8(b,b,b,b,b,b,b,b)); } +static __always_inline float16 fmin_T( float16 a, float b ); +static __always_inline float16 fmin_T( float16 a, float b ) { return fmin_T(a,vector_float16(b,b,b,b,b,b,b,b,b,b,b,b,b,b,b,b)); } +static __always_inline double2 fmin_T( double2 a, double b ); +static __always_inline double2 fmin_T( double2 a, double b ) { return fmin_T(a,vector_double2(b,b)); } +static __always_inline double4 fmin_T( double4 a, double b ); +static __always_inline double4 fmin_T( double4 a, double b ) { return fmin_T(a,vector_double4(b,b,b,b)); } +static __always_inline double8 fmin_T( double8 a, double b ); +static __always_inline double8 fmin_T( double8 a, double b ) { return fmin_T(a,vector_double8(b,b,b,b,b,b,b,b)); } +static __always_inline double16 fmin_T( double16 a, double b ); +static __always_inline double16 fmin_T( double16 a, double b ) { return fmin_T(a,vector_double16(b,b,b,b,b,b,b,b,b,b,b,b,b,b,b,b)); } +__MATH_BUILTIN_2(fmod) +static __always_inline float fract_T( float a, float* b ); +static __always_inline float fract_T( float a, float* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline float2 fract_T( float2 a, float2* b ); +static __always_inline float2 fract_T( float2 a, float2* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline float4 fract_T( float4 a, float4* b ); +static __always_inline float4 fract_T( float4 a, float4* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline float8 fract_T( float8 a, float8* b ); +static __always_inline float8 fract_T( float8 a, float8* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline float16 fract_T( float16 a, float16* b ); +static __always_inline float16 fract_T( float16 a, float16* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline double fract_T( double a, double* b ); +static __always_inline double fract_T( double a, double* b ) { *b = floor(a); return fmin(a-floor(a), 0x1.fffffep-1f); } +static __always_inline double2 fract_T( double2 a, double2* b ); +static __always_inline double2 fract_T( double2 a, double2* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline double4 fract_T( double4 a, double4* b ); +static __always_inline double4 fract_T( double4 a, double4* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline double8 fract_T( double8 a, double8* b ); +static __always_inline double8 fract_T( double8 a, double8* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +static __always_inline double16 fract_T( double16 a, double16* b ); +static __always_inline double16 fract_T( double16 a, double16* b ) { *b = floor_T(a); return fmin_T(a-floor_T(a), 0x1.fffffep-1f); } +// frexp +__MATH_BUILTIN_2(hypot) +// ilogb +// ldexp +__MATH_BUILTIN_1(lgamma) +// lgamma_r __MATH_BUILTIN_1(log) #ifndef __FreeBSD__ __MATH_BUILTIN_1(log2) @@ -1220,110 +1230,84 @@ __MATH_BUILTIN_1(log2) __MATH_BUILTIN_1(log10) __MATH_BUILTIN_1(log1p) __MATH_BUILTIN_1(logb) +static __always_inline float madf( float a, float b, float c ); +static __always_inline float madf( float a, float b, float c ) { return a*b+c; } +static __always_inline double mad( double a, double b, double c ); +static __always_inline double mad( double a, double b, double c ) { return a*b+c; } +__MATH_BUILTIN_3(mad) +static __always_inline float maxmagf( float a, float b ); +static __always_inline float maxmagf( float a, float b ) { + if ( fabsf(a) > fabsf(b) ) return a; + else if ( fabsf(b) > fabsf(a) ) return b; + else return fmaxf(a,b); +} +static __always_inline double maxmag( double a, double b ); +static __always_inline double maxmag( double a, double b ) { + if ( fabs(a) > fabs(b) ) return a; + else if ( fabs(b) > fabs(a) ) return b; + else return fmax(a,b); +} +__MATH_BUILTIN_2(maxmag) +static __always_inline float minmagf( float a, float b ); +static __always_inline float minmagf( float a, float b ) { + if ( fabsf(a) < fabsf(b) ) return a; + else if ( fabsf(b) < fabsf(a) ) return b; + else return fminf(a,b); +} +static __always_inline double minmag( double a, double b ); +static __always_inline double minmag( double a, double b ) { + if ( fabs(a) < fabs(b) ) return a; + else if ( fabs(b) < fabs(a) ) return b; + else return fmin(a,b); +} +__MATH_BUILTIN_2(minmag) +// modf +// nan +__MATH_BUILTIN_2(nextafter) +__MATH_BUILTIN_2(pow) +// pown +// powr +__MATH_BUILTIN_2(remainder) +// remquo __MATH_BUILTIN_1(rint) +// rootn __MATH_BUILTIN_1(round) +static __always_inline float rsqrtf( float a ); +static __always_inline float rsqrtf( float a ) { return 1.0f/sqrtf(a); } +static __always_inline double rsqrt( double a ); +static __always_inline double rsqrt( double a ) { return 1.0/sqrt(a); } +__MATH_BUILTIN_1(rsqrt) __MATH_BUILTIN_1(sin) +// sincos __MATH_BUILTIN_1(sinh) __MATH_BUILTIN_1(sqrt) __MATH_BUILTIN_1(tan) __MATH_BUILTIN_1(tanh) +static __always_inline float tanpif( float a ); +static __always_inline float tanpif( float a ) { return tanf(a*M_PI_F); } +static __always_inline double tanpi( double a ); +static __always_inline double tanpi( double a ) { return tan(a*M_PI); } +__MATH_BUILTIN_1(tanpi) __MATH_BUILTIN_1(tgamma) __MATH_BUILTIN_1(trunc) -/* -static __inline float rsqrt_T( float a ) { return 1.0f/sqrtf(a); } \ -static __inline _float4 rsqrt_T( _float4 a ) \ - { return _vector_float4( \ - 1.0f/sqrtf(a.x),1.0f/sqrtf(a.y),1.0f/sqrtf(a.z),1.0f/sqrtf(a.w) \ - ); } -*/ - -static __always_inline float rsqrt_T( float a ); -static __always_inline float rsqrt_T( float a ) { return 1.0f/sqrtf(a); } - -static __always_inline _float2 rsqrt_T( _float2 a ); -static __always_inline _float2 rsqrt_T( _float2 a ) - { return _vector_float2(1.0f/sqrtf(a.x),1.0f/sqrtf(a.y)); } - -static __always_inline float4 rsqrt_T( _float4 a ); -static __always_inline float4 rsqrt_T( _float4 a ) - { - return _vector_float4( 1.0f/sqrtf(a.x),1.0f/sqrtf(a.y), - 1.0f/sqrtf(a.z),1.0f/sqrtf(a.w)); - } - -static __always_inline double rsqrt_T( double a ); -static __always_inline double rsqrt_T( double a ) { return 1.0f/sqrt(a); } - -static __always_inline _double2 rsqrt_T( _double2 a ); -static __always_inline _double2 rsqrt_T( _double2 a ) - { return _vector_double2(1.0f/sqrt(a.x),1.0f/sqrt(a.y)); } - -template < typename T > - static __always_inline T clamp_T( T a, T b0, T b1 ); -template < typename T > - static __always_inline T clamp_T( T a, T b0, T b1 ) - { return min(max(a,b0),b1); } - -static __always_inline _float2 clamp_T( _float2 a, float b0, float b1); -static __always_inline _float2 clamp_T( _float2 a, float b0, float b1) - { return _vector_float2( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1) ); } - -static __always_inline _float4 - clamp_T( _float4 a, float b0, float b1); -static __always_inline _float4 clamp_T( _float4 a, float b0, float b1) - { - return _vector_float4( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1), - clamp_T(a.z,b0,b1), clamp_T(a.w,b0,b1) ); - } - -static __always_inline _double2 - clamp_T( _double2 a, double b0, double b1); -static __always_inline _double2 clamp_T( _double2 a, double b0, double b1) - { return _vector_double2( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1) ); } - -static __always_inline float dot_T( _float2 a, _float2 b); -static __always_inline float dot_T( _float2 a, _float2 b) - { float2 tmp = a*b; return tmp.x+tmp.y; } - -static __always_inline float dot_T( _float4 a, _float4 b); -static __always_inline float dot_T( _float4 a, _float4 b) - { float4 tmp = a*b; return tmp.x+tmp.y+tmp.z+tmp.w; } - -static __always_inline double dot_T( _double2 a, _double2 b); -static __always_inline double dot_T( _double2 a, _double2 b) - { double2 tmp = a*b; return tmp.x+tmp.y; } - -template < typename T > - static __always_inline T normalize_T( T a ); -template < typename T > - static __always_inline T normalize_T( T a ) { return rsqrt_T(dot_T(a,a))*a; } - -static __always_inline _float4 cross_T( _float4 a, _float4 b); -static __always_inline _float4 cross_T( _float4 a, _float4 b) - { - return _vector_float4(a.y*b.z-a.z*b.y, a.z*b.x-a.x*b.z, - a.x*b.y-a.y*b.x, 0.0f); - } - -/* xxx */ -static __always_inline int sign_T( int a ); -static __always_inline int sign_T( int a ) { return copysignf(1,a); } -static __always_inline float sign_T( float a ); -static __always_inline float sign_T( float a ) { return copysignf(1.0f,a); } - - -#define sqrt(a) sqrt_T(a) #define acos(a) acos_T(a) #define acosh(a) acosh_T(a) +#define acospi(a) acospi_T(a) #define asin(a) asin_T(a) #define asinh(a) asinh_T(a) -#define tan(a) tan_T(a) -#define tanh(a) tanh_T(a) +#define asinpi(a) asinpi_T(a) +#define atan(a) atan_T(a) +#define atan2(a,b) atan2_T(a,b) +#define atanh(a) atanh_T(a) +#define atanpi(a) atanpi_T(a) +#define atan2pi(a,b) atan2pi_T(a) #define cbrt(a) cbrt_T(a) #define ceil(a) ceil_T(a) +#define copysign(a,b) copysign_T(a,b) #define cos(a) cos_T(a) #define cosh(a) cosh_T(a) +#define cospi(a) cospi_T(a) #define erfc(a) erfc_T(a) #define erf(a) erf_T(a) #define exp(a) exp_T(a) @@ -1331,28 +1315,263 @@ static __always_inline float sign_T( float a ) { return copysignf(1.0f,a); } #define exp10(a) exp10_T(a) #define expm1(a) expm1_T(a) #define fabs(a) fabs_T(a) +#define fdim(a,b) fdim_T(a,b) #define floor(a) floor_T(a) +#define fma(a,b,c) fma_T(a,b,c) +#define fmax(a,b) fmax_T(a,b) +#define fmin(a,b) fmin_T(a,b) +#define fmod(a,b) fmod_T(a,b) +#define fract(a,b) fract_T(a,b) +//#define frexp(a,b) frexp_T(a,b) +#define hypot(a,b) hypot_T(a,b) +//#define ilogb(a) ilogb_T(a) +//#define ldexp(a,b) ldexp_T(a,b) +#define lgamma(a) lgamma_T(a) +//#define lgamma_r(a,b) lgamma_r_T(a,b) #define log(a) log_T(a) #define log2(a) log2_T(a) #define log10(a) log10_T(a) #define log1p(a) log1p_T(a) #define logb(a) logb_T(a) +#define mad(a,b,c) mad_T(a,b,c) +#define maxmag(a,b) maxmag_T(a,b) +#define minmag(a,b) minmag_T(a,b) +//#define modf(a,b) modf_T(a,b) +//#define nan(a) nan_T(a) +#define nextafter(a,b) nextafter_T(a,b) +#define pow(a,b) pow_T(a,b) +//#define pown(a,b) pown_T(a,b) +//#define powr(a,b) powr(a,b) +#define remainder(a,b) remainder_T(a,b) +//#define remquo(a,b,c) remquo_T(a,b,c) #define rint(a) rint_T(a) +//#define rootn(a,b) rootn_T(a,b) #define round(a) round_T(a) #define rsqrt(a) rsqrt_T(a) -#define sign(a) sign_T(a) #define sin(a) sin_T(a) +//#define sincos(a,b) sincos_T(a,b) #define sinh(a) sinh_T(a) #define sqrt(a) sqrt_T(a) #define tan(a) tan_T(a) #define tanh(a) tanh_T(a) +#define tanpi(a) tanpi_T(a) #define tgamma(a) tgamma_T(a) -#define tfunc(a) tfunc_T(a) +#define trunc(a) trunc_T(a) + + +/*** common builtin functions [6.11.4] ***/ + +template < typename T > +static __always_inline T clamp_T( T a, T b0, T b1 ); +template < typename T > +static __always_inline T clamp_T( T a, T b0, T b1 ) +{ return min(max(a,b0),b1); } +static __always_inline _float2 clamp_T( _float2 a, float b0, float b1); +static __always_inline _float2 clamp_T( _float2 a, float b0, float b1) +{ return _vector_float2( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1) ); } +static __always_inline _float4 clamp_T( _float4 a, float b0, float b1); +static __always_inline _float4 clamp_T( _float4 a, float b0, float b1) +{ + return _vector_float4( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1), + clamp_T(a.z,b0,b1), clamp_T(a.w,b0,b1) ); +} +static __always_inline _float8 clamp_T( _float8 a, float b0, float b1); +static __always_inline _float8 clamp_T( _float8 a, float b0, float b1) +{ + return _vector_float8( clamp_T(a.s0,b0,b1), clamp_T(a.s1,b0,b1), + clamp_T(a.s2,b0,b1), clamp_T(a.s3,b0,b1), + clamp_T(a.s4,b0,b1), clamp_T(a.s5,b0,b1), + clamp_T(a.s6,b0,b1), clamp_T(a.s7,b0,b1) ); +} +static __always_inline _float16 clamp_T( _float16 a, float b0, float b1); +static __always_inline _float16 clamp_T( _float16 a, float b0, float b1) +{ + return _vector_float16( clamp_T(a.s0,b0,b1), clamp_T(a.s1,b0,b1), + clamp_T(a.s2,b0,b1), clamp_T(a.s3,b0,b1), + clamp_T(a.s4,b0,b1), clamp_T(a.s5,b0,b1), + clamp_T(a.s6,b0,b1), clamp_T(a.s7,b0,b1), + clamp_T(a.s7,b0,b1), clamp_T(a.s9,b0,b1), + clamp_T(a.sA,b0,b1), clamp_T(a.sB,b0,b1), + clamp_T(a.sC,b0,b1), clamp_T(a.sD,b0,b1), + clamp_T(a.sE,b0,b1), clamp_T(a.sF,b0,b1) ); +} +static __always_inline _double2 clamp_T( _double2 a, double b0, double b1); +static __always_inline _double2 clamp_T( _double2 a, double b0, double b1) +{ return _vector_double2( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1) ); } +static __always_inline _double4 clamp_T( _double4 a, double b0, double b1); +static __always_inline _double4 clamp_T( _double4 a, double b0, double b1) +{ + return _vector_double4( clamp_T(a.x,b0,b1), clamp_T(a.y,b0,b1), + clamp_T(a.z,b0,b1), clamp_T(a.w,b0,b1) ); +} +static __always_inline _double8 clamp_T( _double8 a, double b0, double b1); +static __always_inline _double8 clamp_T( _double8 a, double b0, double b1) +{ + return _vector_double8( clamp_T(a.s0,b0,b1), clamp_T(a.s1,b0,b1), + clamp_T(a.s2,b0,b1), clamp_T(a.s3,b0,b1), + clamp_T(a.s4,b0,b1), clamp_T(a.s5,b0,b1), + clamp_T(a.s6,b0,b1), clamp_T(a.s7,b0,b1) ); +} +static __always_inline _double16 clamp_T( _double16 a, double b0, double b1); +static __always_inline _double16 clamp_T( _double16 a, double b0, double b1) +{ + return _vector_double16( clamp_T(a.s0,b0,b1), clamp_T(a.s1,b0,b1), + clamp_T(a.s2,b0,b1), clamp_T(a.s3,b0,b1), + clamp_T(a.s4,b0,b1), clamp_T(a.s5,b0,b1), + clamp_T(a.s6,b0,b1), clamp_T(a.s7,b0,b1), + clamp_T(a.s7,b0,b1), clamp_T(a.s9,b0,b1), + clamp_T(a.sA,b0,b1), clamp_T(a.sB,b0,b1), + clamp_T(a.sC,b0,b1), clamp_T(a.sD,b0,b1), + clamp_T(a.sE,b0,b1), clamp_T(a.sF,b0,b1) ); +} +static __always_inline float degreesf( float a ); +static __always_inline float degreesf( float a ) { return (a*180.0f)/M_PI_F; } +static __always_inline double degrees( double a ); +static __always_inline double degrees( double a ) { return (a*180.0)/M_PI; } +__MATH_BUILTIN_1(degrees) +// max +// min +// mix +static __always_inline float radiansf( float a ); +static __always_inline float radiansf( float a ) { return (a*M_PI_F)/180.0f; } +static __always_inline double radians( double a ); +static __always_inline double radians( double a ) { return (a*M_PI)/180.0; } +__MATH_BUILTIN_1(radians) +static __always_inline float stepf( float a, float b ); +static __always_inline float stepf( float a, float b ) { return ((b +static __always_inline T length_T( T a ); +template < typename T > +static __always_inline T length_T( T a ) { return sqrt_T(dot_T(a,a)); } + +template < typename T > +static __always_inline T distance_T( T a, T b ); +template < typename T > +static __always_inline T distance_T( T a, T b ) { return length_T(b-a); } + +template < typename T > +static __always_inline T normalize_T( T a ); +template < typename T > +static __always_inline T normalize_T( T a ) { return rsqrt_T(dot_T(a,a))*a; } + +// fast_distance +// fast_length +// fast_normalize + #define dot(a,b) dot_T(a,b) -#define normalize(a) normalize_T(a) #define cross(a,b) cross_T(a,b) +#define distance(a,b) distance_T(a,b) +#define length(a) length_T(a) +#define normalize(a) normalize_T(a) +#define fast_distance(a,b) distance_T(a,b) +#define fast_length(a,b) length_T(a,b) +#define fast_normalize(a) normalize_T(a) /*** sampler declarations [6.11.8.1] ***/ @@ -1368,25 +1587,58 @@ typedef int sampler_t; #define read_imagei( _img, _smplr, _xy) \ *((int4*)((intptr_t)_img+128)+((__global size_t*)_img)[0]*(_xy).y+(_xy).x) +#define read_imageui( _img, _smplr, _xy) \ +*((uint4*)((intptr_t)_img+128)+((__global size_t*)_img)[0]*(_xy).y+(_xy).x) + #define read_imagef( _img, _smplr, _xy) \ *((float4*)((intptr_t)_img+128)+((__global size_t*)_img)[0]*(_xy).y+(_xy).x) +/*** integer builtins [6.12.3] ***/ + + /*** builtin extensions for initializing vector data types [non-standard] ***/ -#define __builtin_vector_char2(x,y) vector_char2(x,y) -#define __builtin_vector_char4(x,y,z,w) vector_char4(x,y,z,w) -#define __builtin_vector_uchar2(x,y) vector_uchar2(x,y) -#define __builtin_vector_uchar4(x,y,z,w) vector_uchar4(x,y,z,w) -#define __builtin_vector_int2(x,y) vector_int2(x,y) -#define __builtin_vector_int4(x,y,z,w) vector_int4(x,y,z,w) -#define __builtin_vector_long2(x,y) vector_long2(x,y) -#define __builtin_vector_uint2(x,y) vector_uint2(x,y) -#define __builtin_vector_uint4(x,y,z,w) vector_uint4(x,y,z,w) -#define __builtin_vector_ulong2(x,y) vector_ulong2(x,y) -#define __builtin_vector_float2(x,y) vector_float2(x,y) -#define __builtin_vector_float4(x,y,z,w) vector_float4(x,y,z,w) -#define __builtin_vector_double2(x,y) vector_double2(x,y) +#define __builtin_vector_char2(a,b) vector_char2(a,b) +#define __builtin_vector_char4(a,b,c,d) vector_char4(a,b,c,d) +#define __builtin_vector_char8(a,b,c,d,e,f,g,h) vector_char8(a,b,c,d,e,f,g,h) +#define __builtin_vector_char16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_char16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uchar2(a,b) vector_uchar2(a,b) +#define __builtin_vector_uchar4(a,b,c,d) vector_uchar4(a,b,c,d) +#define __builtin_vector_uchar8(a,b,c,d,e,f,g,h) vector_uchar8(a,b,c,d,e,f,g,h) +#define __builtin_vector_uchar16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_uchar16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_short2(a,b) vector_short2(a,b) +#define __builtin_vector_short4(a,b,c,d) vector_short4(a,b,c,d) +#define __builtin_vector_short8(a,b,c,d,e,f,g,h) vector_short8(a,b,c,d,e,f,g,h) +#define __builtin_vector_short16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_short16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ushort2(a,b) vector_ushort2(a,b) +#define __builtin_vector_ushort4(a,b,c,d) vector_ushort4(a,b,c,d) +#define __builtin_vector_ushort8(a,b,c,d,e,f,g,h) vector_ushort8(a,b,c,d,e,f,g,h) +#define __builtin_vector_ushort16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_ushort16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_int2(a,b) vector_int2(a,b) +#define __builtin_vector_int4(a,b,c,d) vector_int4(a,b,c,d) +#define __builtin_vector_int8(a,b,c,d,e,f,g,h) vector_int8(a,b,c,d,e,f,g,h) +#define __builtin_vector_int16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_int16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uint2(a,b) vector_uint2(a,b) +#define __builtin_vector_uint4(a,b,c,d) vector_uint4(a,b,c,d) +#define __builtin_vector_uint8(a,b,c,d,e,f,g,h) vector_uint8(a,b,c,d,e,f,g,h) +#define __builtin_vector_uint16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_uint16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_long2(a,b) vector_long2(a,b) +#define __builtin_vector_long4(a,b,c,d) vector_long4(a,b,c,d) +#define __builtin_vector_long8(a,b,c,d,e,f,g,h) vector_long8(a,b,c,d,e,f,g,h) +#define __builtin_vector_long16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_long16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ulong2(a,b) vector_ulong2(a,b) +#define __builtin_vector_ulong4(a,b,c,d) vector_ulong4(a,b,c,d) +#define __builtin_vector_ulong8(a,b,c,d,e,f,g,h) vector_ulong8(a,b,c,d,e,f,g,h) +#define __builtin_vector_ulong16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_ulong16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_float2(a,b) vector_float2(a,b) +#define __builtin_vector_float4(a,b,c,d) vector_float4(a,b,c,d) +#define __builtin_vector_float8(a,b,c,d,e,f,g,h) vector_float8(a,b,c,d,e,f,g,h) +#define __builtin_vector_float16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_float16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_double2(a,b) vector_double2(a,b) +#define __builtin_vector_double4(a,b,c,d) vector_double4(a,b,c,d) +#define __builtin_vector_double8(a,b,c,d,e,f,g,h) vector_double8(a,b,c,d,e,f,g,h) +#define __builtin_vector_double16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) vector_double16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) #else @@ -1396,19 +1648,47 @@ typedef int sampler_t; /*** builtin extensions for initializing vector data types [non-standard] ***/ -#define __builtin_vector_char2(x,y) (char2)(x,y) -#define __builtin_vector_char4(x,y,z,w) (char4)(x,y,z,w) -#define __builtin_vector_uchar2(x,y) (uchar2)(x,y) -#define __builtin_vector_uchar4(x,y,z,w) (uchar4)(x,y,z,w) -#define __builtin_vector_int2(x,y) (int2)(x,y) -#define __builtin_vector_int4(x,y,z,w) (int4)(x,y,z,w) -#define __builtin_vector_uint2(x,y) (uint2)(x,y) -#define __builtin_vector_uint4(x,y,z,w) (uint4)(x,y,z,w) -#define __builtin_vector_float2(x,y) (float2)(x,y) -#define __builtin_vector_float4(x,y,z,w) (float4)(x,y,z,w) -#define __builtin_vector_double2(x,y) (double2)(x,y) - -#endif +#define __builtin_vector_char2(a,b) (vector_char2)(a,b) +#define __builtin_vector_char4(a,b,c,d) (vector_char4)(a,b,c,d) +#define __builtin_vector_char8(a,b,c,d,e,f,g,h) (vector_char8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_char16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_char16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uchar2(a,b) (vector_uchar2)(a,b) +#define __builtin_vector_uchar4(a,b,c,d) (vector_uchar4)(a,b,c,d) +#define __builtin_vector_uchar8(a,b,c,d,e,f,g,h) (vector_uchar8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_uchar16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_uchar16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_short2(a,b) (vector_short2)(a,b) +#define __builtin_vector_short4(a,b,c,d) (vector_short4)(a,b,c,d) +#define __builtin_vector_short8(a,b,c,d,e,f,g,h) (vector_short8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_short16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_short16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ushort2(a,b) (vector_ushort2)(a,b) +#define __builtin_vector_ushort4(a,b,c,d) (vector_ushort4)(a,b,c,d) +#define __builtin_vector_ushort8(a,b,c,d,e,f,g,h) (vector_ushort8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_ushort16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_ushort16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_int2(a,b) (vector_int2)(a,b) +#define __builtin_vector_int4(a,b,c,d) (vector_int4)(a,b,c,d) +#define __builtin_vector_int8(a,b,c,d,e,f,g,h) (vector_int8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_int16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_int16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uint2(a,b) (vector_uint2)(a,b) +#define __builtin_vector_uint4(a,b,c,d) (vector_uint4)(a,b,c,d) +#define __builtin_vector_uint8(a,b,c,d,e,f,g,h) (vector_uint8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_uint16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_uint16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_long2(a,b) (vector_long2)(a,b) +#define __builtin_vector_long4(a,b,c,d) (vector_long4)(a,b,c,d) +#define __builtin_vector_long8(a,b,c,d,e,f,g,h) (vector_long8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_long16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_long16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ulong2(a,b) (vector_ulong2)(a,b) +#define __builtin_vector_ulong4(a,b,c,d) (vector_ulong4)(a,b,c,d) +#define __builtin_vector_ulong8(a,b,c,d,e,f,g,h) (vector_ulong8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_ulong16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_ulong16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_float2(a,b) (vector_float2)(a,b) +#define __builtin_vector_float4(a,b,c,d) (vector_float4)(a,b,c,d) +#define __builtin_vector_float8(a,b,c,d,e,f,g,h) (vector_float8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_float16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_float16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_double2(a,b) (vector_double2)(a,b) +#define __builtin_vector_double4(a,b,c,d) (vector_double4)(a,b,c,d) +#define __builtin_vector_double8(a,b,c,d,e,f,g,h) (vector_double8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_double16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (vector_double16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) #endif +#endif // _opencl_lift_h diff --git a/src/libcoprthr/ser_engine.h b/src/libcoprthr/ser_engine.h index abe6ebd2..cf6403e0 100644 --- a/src/libcoprthr/ser_engine.h +++ b/src/libcoprthr/ser_engine.h @@ -31,20 +31,48 @@ /* XXX this brings in the vec types. temporary soln. improve it. -DAR */ #include typedef unsigned char uchar; +typedef unsigned short ushort; typedef unsigned int uint; typedef cl_char2 char2; typedef cl_char4 char4; +typedef cl_char8 char8; +typedef cl_char16 char16; typedef cl_uchar2 uchar2; typedef cl_uchar4 uchar4; +typedef cl_uchar8 uchar8; +typedef cl_uchar16 uchar16; +typedef cl_short2 short2; +typedef cl_short4 short4; +typedef cl_short8 short8; +typedef cl_short16 short16; +typedef cl_ushort2 ushort2; +typedef cl_ushort4 ushort4; +typedef cl_ushort8 ushort8; +typedef cl_ushort16 ushort16; typedef cl_int2 int2; typedef cl_int4 int4; +typedef cl_int8 int8; +typedef cl_int16 int16; typedef cl_uint2 uint2; typedef cl_uint4 uint4; +typedef cl_uint8 uint8; +typedef cl_uint16 uint16; typedef cl_long2 long2; +typedef cl_long4 long4; +typedef cl_long8 long8; +typedef cl_long16 long16; typedef cl_ulong2 ulong2; +typedef cl_ulong4 ulong4; +typedef cl_ulong8 ulong8; +typedef cl_ulong16 ulong16; typedef cl_float2 float2; typedef cl_float4 float4; +typedef cl_float8 float8; +typedef cl_float16 float16; typedef cl_double2 double2; +typedef cl_double4 double4; +typedef cl_double8 double8; +typedef cl_double16 double16; #endif diff --git a/src/libcoprthr/sl_engine.h b/src/libcoprthr/sl_engine.h index 2e2ac1f3..da95eb1d 100644 --- a/src/libcoprthr/sl_engine.h +++ b/src/libcoprthr/sl_engine.h @@ -31,20 +31,48 @@ /* XXX this brings in the vec types. temporary soln. improve it. -DAR */ #include typedef unsigned char uchar; +typedef unsigned short ushort; typedef unsigned int uint; typedef cl_char2 char2; typedef cl_char4 char4; +typedef cl_char8 char8; +typedef cl_char16 char16; typedef cl_uchar2 uchar2; typedef cl_uchar4 uchar4; +typedef cl_uchar8 uchar8; +typedef cl_uchar16 uchar16; +typedef cl_short2 short2; +typedef cl_short4 short4; +typedef cl_short8 short8; +typedef cl_short16 short16; +typedef cl_ushort2 ushort2; +typedef cl_ushort4 ushort4; +typedef cl_ushort8 ushort8; +typedef cl_ushort16 ushort16; typedef cl_int2 int2; typedef cl_int4 int4; +typedef cl_int8 int8; +typedef cl_int16 int16; typedef cl_uint2 uint2; typedef cl_uint4 uint4; +typedef cl_uint8 uint8; +typedef cl_uint16 uint16; typedef cl_long2 long2; +typedef cl_long4 long4; +typedef cl_long8 long8; +typedef cl_long16 long16; typedef cl_ulong2 ulong2; +typedef cl_ulong4 ulong4; +typedef cl_ulong8 ulong8; +typedef cl_ulong16 ulong16; typedef cl_float2 float2; typedef cl_float4 float4; +typedef cl_float8 float8; +typedef cl_float16 float16; typedef cl_double2 double2; +typedef cl_double4 double4; +typedef cl_double8 double8; +typedef cl_double16 double16; #endif diff --git a/src/libstdcl/clforka.h b/src/libstdcl/clforka.h index 97c37315..77a57a73 100644 --- a/src/libstdcl/clforka.h +++ b/src/libstdcl/clforka.h @@ -132,54 +132,118 @@ #ifdef __cplusplus -inline static -void __clarg_set(CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_int a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_uint a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_float a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_int2 a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_uint2 a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_float2 a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_int4 a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_uint4 a) -{ ::clarg_set(cp,krn,ia,(a)); } -inline static -void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, cl_float4 a) -{ ::clarg_set(cp,krn,ia,(a)); } inline static void __clarg_set( CLCONTEXT* cp, cl_kernel krn, unsigned int ia, void* a) { ::clarg_set_global(cp,krn,ia,(a)); } +#define __cpp_clarg_set(t) \ + inline static void __clarg_set(CLCONTEXT* cp, cl_kernel krn, unsigned int ia, t a){ :: clarg_set(cp,krn,ia, (a)); } + +__cpp_clarg_set(cl_char) +__cpp_clarg_set(cl_uchar) +__cpp_clarg_set(cl_short) +__cpp_clarg_set(cl_ushort) +__cpp_clarg_set(cl_int) +__cpp_clarg_set(cl_uint) +__cpp_clarg_set(cl_long) +__cpp_clarg_set(cl_ulong) +__cpp_clarg_set(cl_float) +__cpp_clarg_set(cl_double) +__cpp_clarg_set(cl_char2) +__cpp_clarg_set(cl_uchar2) +__cpp_clarg_set(cl_short2) +__cpp_clarg_set(cl_ushort2) +__cpp_clarg_set(cl_int2) +__cpp_clarg_set(cl_uint2) +__cpp_clarg_set(cl_long2) +__cpp_clarg_set(cl_ulong2) +__cpp_clarg_set(cl_float2) +__cpp_clarg_set(cl_double2) +__cpp_clarg_set(cl_char4) +__cpp_clarg_set(cl_uchar4) +__cpp_clarg_set(cl_short4) +__cpp_clarg_set(cl_ushort4) +__cpp_clarg_set(cl_int4) +__cpp_clarg_set(cl_uint4) +__cpp_clarg_set(cl_long4) +__cpp_clarg_set(cl_ulong4) +__cpp_clarg_set(cl_float4) +__cpp_clarg_set(cl_double4) +__cpp_clarg_set(cl_char8) +__cpp_clarg_set(cl_uchar8) +__cpp_clarg_set(cl_short8) +__cpp_clarg_set(cl_ushort8) +__cpp_clarg_set(cl_int8) +__cpp_clarg_set(cl_uint8) +__cpp_clarg_set(cl_long8) +__cpp_clarg_set(cl_ulong8) +__cpp_clarg_set(cl_float8) +__cpp_clarg_set(cl_double8) +__cpp_clarg_set(cl_char16) +__cpp_clarg_set(cl_uchar16) +__cpp_clarg_set(cl_short16) +__cpp_clarg_set(cl_ushort16) +__cpp_clarg_set(cl_int16) +__cpp_clarg_set(cl_uint16) +__cpp_clarg_set(cl_long16) +__cpp_clarg_set(cl_ulong16) +__cpp_clarg_set(cl_float16) +__cpp_clarg_set(cl_double16) + #else #define __clarg_set(cp,krn,ia,a) \ ({ \ -__iftype_clarg_set(cp,krn,ia,a,cl_int) \ -else __iftype_clarg_set(cp,krn,ia,a,cl_char) \ +__iftype_clarg_set(cp,krn,ia,a,cl_char) \ else __iftype_clarg_set(cp,krn,ia,a,cl_uchar) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_short) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ushort) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_int) \ else __iftype_clarg_set(cp,krn,ia,a,cl_uint) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_long) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ulong) \ else __iftype_clarg_set(cp,krn,ia,a,cl_float) \ else __iftype_clarg_set(cp,krn,ia,a,cl_double) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_char2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uchar2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_short2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ushort2) \ else __iftype_clarg_set(cp,krn,ia,a,cl_int2) \ else __iftype_clarg_set(cp,krn,ia,a,cl_uint2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_long2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ulong2) \ else __iftype_clarg_set(cp,krn,ia,a,cl_float2) \ else __iftype_clarg_set(cp,krn,ia,a,cl_double2) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_char4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uchar4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_short4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ushort4) \ else __iftype_clarg_set(cp,krn,ia,a,cl_int4) \ else __iftype_clarg_set(cp,krn,ia,a,cl_uint4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_long4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ulong4) \ else __iftype_clarg_set(cp,krn,ia,a,cl_float4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_double4) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_char8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uchar8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_short8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ushort8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_int8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uint8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_long8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ulong8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_float8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_double8) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_char16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uchar16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_short16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ushort16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_int16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_uint16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_long16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_ulong16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_float16) \ +else __iftype_clarg_set(cp,krn,ia,a,cl_double16) \ else { fprintf(stderr,"__clarg_set: unrecognized type\n"); exit(-1); } \ }) diff --git a/src/libstdcl/stdcl_kernel_openclcompat.h b/src/libstdcl/stdcl_kernel_openclcompat.h index 492c22ce..0353838b 100644 --- a/src/libstdcl/stdcl_kernel_openclcompat.h +++ b/src/libstdcl/stdcl_kernel_openclcompat.h @@ -27,17 +27,46 @@ #ifndef _STDCL_KERNEL_OPENCLCOMPAT_H #define _STDCL_KERNEL_OPENCLCOMPAT_H -#define __builtin_vector_char2(x,y) (char2)(x,y) -#define __builtin_vector_char4(x,y,z,w) (char4)(x,y,z,w) -#define __builtin_vector_uchar2(x,y) (uchar2)(x,y) -#define __builtin_vector_uchar4(x,y,z,w) (uchar4)(x,y,z,w) -#define __builtin_vector_int2(x,y) (int2)(x,y) -#define __builtin_vector_int4(x,y,z,w) (int4)(x,y,z,w) -#define __builtin_vector_uint2(x,y) (uint2)(x,y) -#define __builtin_vector_uint4(x,y,z,w) (uint4)(x,y,z,w) -#define __builtin_vector_float2(x,y) (float2)(x,y) -#define __builtin_vector_float4(x,y,z,w) (float4)(x,y,z,w) -#define __builtin_vector_double2(x,y) (double2)(x,y) +#define __builtin_vector_short2(a,b) (short2)(a,b) +#define __builtin_vector_short4(a,b,c,d) (short4)(a,b,c,d) +#define __builtin_vector_short8(a,b,c,d,e,f,g,h) (short8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_short16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (short16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ushort2(a,b) (ushort2)(a,b) +#define __builtin_vector_ushort4(a,b,c,d) (ushort4)(a,b,c,d) +#define __builtin_vector_ushort8(a,b,c,d,e,f,g,h) (ushort8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_ushort16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (ushort16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_char2(a,b) (char2)(a,b) +#define __builtin_vector_char4(a,b,c,d) (char4)(a,b,c,d) +#define __builtin_vector_char8(a,b,c,d,e,f,g,h) (char8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_char16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (char16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uchar2(a,b) (uchar2)(a,b) +#define __builtin_vector_uchar4(a,b,c,d) (uchar4)(a,b,c,d) +#define __builtin_vector_uchar8(a,b,c,d,e,f,g,h) (uchar8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_uchar16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (uchar16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_int2(a,b) (int2)(a,b) +#define __builtin_vector_int4(a,b,c,d) (int4)(a,b,c,d) +#define __builtin_vector_int8(a,b,c,d,e,f,g,h) (int8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_int16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (int16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_uint2(a,b) (uint2)(a,b) +#define __builtin_vector_uint4(a,b,c,d) (uint4)(a,b,c,d) +#define __builtin_vector_uint8(a,b,c,d,e,f,g,h) (uint8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_uint16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (uint16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_long2(a,b) (long2)(a,b) +#define __builtin_vector_long4(a,b,c,d) (long4)(a,b,c,d) +#define __builtin_vector_long8(a,b,c,d,e,f,g,h) (long8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_long16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (long16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_ulong2(a,b) (ulong2)(a,b) +#define __builtin_vector_ulong4(a,b,c,d) (ulong4)(a,b,c,d) +#define __builtin_vector_ulong8(a,b,c,d,e,f,g,h) (ulong8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_ulong16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (ulong16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_float2(a,b) (float2)(a,b) +#define __builtin_vector_float4(a,b,c,d) (float4)(a,b,c,d) +#define __builtin_vector_float8(a,b,c,d,e,f,g,h) (float8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_float16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (float16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) +#define __builtin_vector_double2(a,b) (double2)(a,b) +#define __builtin_vector_double4(a,b,c,d) (double4)(a,b,c,d) +#define __builtin_vector_double8(a,b,c,d,e,f,g,h) (double8)(a,b,c,d,e,f,g,h) +#define __builtin_vector_double16(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) (double16)(a,b,c,d,e,f,g,h,i,j,k,l,m,n,o,p) #endif diff --git a/tools/xclnm/xclnm_gram.y b/tools/xclnm/xclnm_gram.y index 1ca4c890..52f695a5 100644 --- a/tools/xclnm/xclnm_gram.y +++ b/tools/xclnm/xclnm_gram.y @@ -74,6 +74,28 @@ %token TYPE_FLOAT_VEC4 %token TYPE_DOUBLE_VEC4 +%token TYPE_INT8_VEC8 +%token TYPE_INT16_VEC8 +%token TYPE_INT32_VEC8 +%token TYPE_INT64_VEC8 +%token TYPE_UINT8_VEC8 +%token TYPE_UINT16_VEC8 +%token TYPE_UINT32_VEC8 +%token TYPE_UINT64_VEC8 +%token TYPE_FLOAT_VEC8 +%token TYPE_DOUBLE_VEC8 + +%token TYPE_INT8_VEC16 +%token TYPE_INT16_VEC16 +%token TYPE_INT32_VEC16 +%token TYPE_INT64_VEC16 +%token TYPE_UINT8_VEC16 +%token TYPE_UINT16_VEC16 +%token TYPE_UINT32_VEC16 +%token TYPE_UINT64_VEC16 +%token TYPE_FLOAT_VEC16 +%token TYPE_DOUBLE_VEC16 + /* qualifiers */ /* generic tokens */ @@ -143,6 +165,28 @@ void yyerror(const char*); %type TYPE_FLOAT_VEC4 %type TYPE_DOUBLE_VEC4 +%type TYPE_INT8_VEC8 +%type TYPE_INT16_VEC8 +%type TYPE_INT32_VEC8 +%type TYPE_INT64_VEC8 +%type TYPE_UINT8_VEC8 +%type TYPE_UINT16_VEC8 +%type TYPE_UINT32_VEC8 +%type TYPE_UINT64_VEC8 +%type TYPE_FLOAT_VEC8 +%type TYPE_DOUBLE_VEC8 + +%type TYPE_INT8_VEC16 +%type TYPE_INT16_VEC16 +%type TYPE_INT32_VEC16 +%type TYPE_INT64_VEC16 +%type TYPE_UINT8_VEC16 +%type TYPE_UINT16_VEC16 +%type TYPE_UINT32_VEC16 +%type TYPE_UINT64_VEC16 +%type TYPE_FLOAT_VEC16 +%type TYPE_DOUBLE_VEC16 + %type ICONST %type TYPE %type SYMBOL @@ -233,6 +277,26 @@ type: TYPE_VOID { $$ = node_create_type(1,1,TYPEID_VOID,0,0); } | TYPE_UINT64_VEC4 { $$ = node_create_type(1,4,TYPEID_ULONG,0,0); } | TYPE_FLOAT_VEC4 { $$ = node_create_type(1,4,TYPEID_FLOAT,0,0); } | TYPE_DOUBLE_VEC4 { $$ = node_create_type(1,4,TYPEID_DOUBLE,0,0); } + | TYPE_INT8_VEC8 { $$ = node_create_type(1,8,TYPEID_CHAR,0,0); } + | TYPE_INT16_VEC8 { $$ = node_create_type(1,8,TYPEID_SHORT,0,0); } + | TYPE_INT32_VEC8 { $$ = node_create_type(1,8,TYPEID_INT,0,0); } + | TYPE_INT64_VEC8 { $$ = node_create_type(1,8,TYPEID_LONG,0,0); } + | TYPE_UINT8_VEC8 { $$ = node_create_type(1,8,TYPEID_UCHAR,0,0); } + | TYPE_UINT16_VEC8 { $$ = node_create_type(1,8,TYPEID_USHORT,0,0); } + | TYPE_UINT32_VEC8 { $$ = node_create_type(1,8,TYPEID_UINT,0,0); } + | TYPE_UINT64_VEC8 { $$ = node_create_type(1,8,TYPEID_ULONG,0,0); } + | TYPE_FLOAT_VEC8 { $$ = node_create_type(1,8,TYPEID_FLOAT,0,0); } + | TYPE_DOUBLE_VEC8 { $$ = node_create_type(1,8,TYPEID_DOUBLE,0,0); } + | TYPE_INT8_VEC16 { $$ = node_create_type(1,16,TYPEID_CHAR,0,0); } + | TYPE_INT16_VEC16 { $$ = node_create_type(1,16,TYPEID_SHORT,0,0); } + | TYPE_INT32_VEC16 { $$ = node_create_type(1,16,TYPEID_INT,0,0); } + | TYPE_INT64_VEC16 { $$ = node_create_type(1,16,TYPEID_LONG,0,0); } + | TYPE_UINT8_VEC16 { $$ = node_create_type(1,16,TYPEID_UCHAR,0,0); } + | TYPE_UINT16_VEC16 { $$ = node_create_type(1,16,TYPEID_USHORT,0,0); } + | TYPE_UINT32_VEC16 { $$ = node_create_type(1,16,TYPEID_UINT,0,0); } + | TYPE_UINT64_VEC16 { $$ = node_create_type(1,16,TYPEID_ULONG,0,0); } + | TYPE_FLOAT_VEC16 { $$ = node_create_type(1,16,TYPEID_FLOAT,0,0); } + | TYPE_DOUBLE_VEC16 { $$ = node_create_type(1,16,TYPEID_DOUBLE,0,0); } | ADDRSPACE_GLOBAL type { $$=node_set_addrspace($2,1); } | ADDRSPACE_LOCAL type { $$=node_set_addrspace($2,3); } | type ptrc { $$=node_set_ptrc($1,$2); } ; diff --git a/tools/xclnm/xclnm_scan.l b/tools/xclnm/xclnm_scan.l index 37566908..54affe38 100644 --- a/tools/xclnm/xclnm_scan.l +++ b/tools/xclnm/xclnm_scan.l @@ -127,11 +127,41 @@ FILE* __xclnm_yyin; "unsigned"{WS}+"short4" { __plb(); ival_return(TYPE_UINT16_VEC4); } "uint4" { __plb(); ival_return(TYPE_UINT32_VEC4); } "unsigned"{WS}+"int4" { __plb(); ival_return(TYPE_UINT32_VEC4); } -"ulong4" { __plb(); ival_return(TYPE_UINT32_VEC4); } -"unsigned"{WS}+"long4" { __plb(); ival_return(TYPE_UINT32_VEC4); } +"ulong4" { __plb(); ival_return(TYPE_UINT64_VEC4); } +"unsigned"{WS}+"long4" { __plb(); ival_return(TYPE_UINT64_VEC4); } "float4" { __plb(); ival_return(TYPE_FLOAT_VEC4); } "double4" { __plb(); ival_return(TYPE_DOUBLE_VEC4); } +"char8" { __plb(); ival_return(TYPE_INT8_VEC8); } +"short8" { __plb(); ival_return(TYPE_INT16_VEC8); } +"int8" { __plb(); ival_return(TYPE_INT32_VEC8); } +"long8" { __plb(); ival_return(TYPE_INT32_VEC8); } +"uchar8" { __plb(); ival_return(TYPE_UINT8_VEC8); } +"unsigned"{WS}+"char8" { __plb(); ival_return(TYPE_UINT8_VEC8); } +"ushort8" { __plb(); ival_return(TYPE_UINT16_VEC8); } +"unsigned"{WS}+"short8" { __plb(); ival_return(TYPE_UINT16_VEC8); } +"uint8" { __plb(); ival_return(TYPE_UINT32_VEC8); } +"unsigned"{WS}+"int8" { __plb(); ival_return(TYPE_UINT32_VEC8); } +"ulong8" { __plb(); ival_return(TYPE_UINT64_VEC8); } +"unsigned"{WS}+"long8" { __plb(); ival_return(TYPE_UINT64_VEC8); } +"float8" { __plb(); ival_return(TYPE_FLOAT_VEC8); } +"double8" { __plb(); ival_return(TYPE_DOUBLE_VEC8); } + +"char16" { __plb(); ival_return(TYPE_INT8_VEC16); } +"short16" { __plb(); ival_return(TYPE_INT16_VEC16); } +"int16" { __plb(); ival_return(TYPE_INT32_VEC16); } +"long16" { __plb(); ival_return(TYPE_INT32_VEC16); } +"uchar16" { __plb(); ival_return(TYPE_UINT8_VEC16); } +"unsigned"{WS}+"char16" { __plb(); ival_return(TYPE_UINT8_VEC16); } +"ushort16" { __plb(); ival_return(TYPE_UINT16_VEC16); } +"unsigned"{WS}+"short16" { __plb(); ival_return(TYPE_UINT16_VEC16); } +"uint16" { __plb(); ival_return(TYPE_UINT32_VEC16); } +"unsigned"{WS}+"int16" { __plb(); ival_return(TYPE_UINT32_VEC16); } +"ulong16" { __plb(); ival_return(TYPE_UINT64_VEC16); } +"unsigned"{WS}+"long16" { __plb(); ival_return(TYPE_UINT64_VEC16); } +"float16" { __plb(); ival_return(TYPE_FLOAT_VEC16); } +"double16" { __plb(); ival_return(TYPE_DOUBLE_VEC16); } + "..." { __plb(); return(VARG); } {DIG}+ { __plb(); __xclnm_lval.ival = atoi(yytext); return(ICONST); } From e547731bbdf0510f1ebe7e9632b37292c59685da Mon Sep 17 00:00:00 2001 From: JAR Date: Sat, 13 Apr 2013 22:24:16 -0400 Subject: [PATCH 2/2] Added VIM syntax highlighting files --- doc/vim_syntax_highlighting/README | 2 + doc/vim_syntax_highlighting/example.vimrc | 12 + doc/vim_syntax_highlighting/syntax/cl.vim | 358 ++++++++++ .../syntax/coprthr_kern.vim | 75 ++ doc/vim_syntax_highlighting/syntax/opencl.vim | 671 ++++++++++++++++++ doc/vim_syntax_highlighting/syntax/stdcl.vim | 113 +++ 6 files changed, 1231 insertions(+) create mode 100644 doc/vim_syntax_highlighting/README create mode 100644 doc/vim_syntax_highlighting/example.vimrc create mode 100644 doc/vim_syntax_highlighting/syntax/cl.vim create mode 100644 doc/vim_syntax_highlighting/syntax/coprthr_kern.vim create mode 100644 doc/vim_syntax_highlighting/syntax/opencl.vim create mode 100644 doc/vim_syntax_highlighting/syntax/stdcl.vim diff --git a/doc/vim_syntax_highlighting/README b/doc/vim_syntax_highlighting/README new file mode 100644 index 00000000..e4a92d55 --- /dev/null +++ b/doc/vim_syntax_highlighting/README @@ -0,0 +1,2 @@ +Move the contents of ./syntax to ~/.vim/syntax +Copy the code in example.vimrc to your ~/.vimrc file diff --git a/doc/vim_syntax_highlighting/example.vimrc b/doc/vim_syntax_highlighting/example.vimrc new file mode 100644 index 00000000..164aa618 --- /dev/null +++ b/doc/vim_syntax_highlighting/example.vimrc @@ -0,0 +1,12 @@ +filetype on +au Syntax opencl source ~/.vim/syntax/opencl.vim +au BufNewFile,BufRead *.c,*.cpp,*.h,*.hpp set ft=opencl +au Syntax stdcl source ~/.vim/syntax/stdcl.vim +au BufNewFile,BufRead *.c,*.cpp,*.h,*.hpp set ft=stdcl +au Syntax cl source ~/.vim/syntax/cl.vim +au BufNewFile,BufRead *.cl,*.clh set ft=cl +au Syntax coprthr_kern source ~/.vim/syntax/coprthr_kern.vim +au BufNewFile,BufRead *.cl set ft=coprthr_kern +set tabstop=5 +syntax on + diff --git a/doc/vim_syntax_highlighting/syntax/cl.vim b/doc/vim_syntax_highlighting/syntax/cl.vim new file mode 100644 index 00000000..19392938 --- /dev/null +++ b/doc/vim_syntax_highlighting/syntax/cl.vim @@ -0,0 +1,358 @@ +" Vim syntax file +" Language: OpenCL +" Maintainer: James Ross +" Version: 0.5 +" Last Change: 2013-03-25 +" +" Description: Syntax highlighting for OpenCL 1.2 + +if exists("b:current_syntax") + finish +endif + +if version < 600 + so :p:h/c.vim +else + runtime! syntax/c.vim + unlet b:current_syntax +endif + +" Enumerated Types +syn keyword clType cl_channel_order cl_channel_type cl_command_queue_info cl_command_queue_properties + +" [6.1.1] Built-in Data Types +syn keyword clType bool +syn keyword clType half +syn keyword clType quad +syn keyword clType char +syn keyword clType uchar +syn keyword clType short +syn keyword clType ushort +syn keyword clType int +syn keyword clType uint +syn keyword clType long +syn keyword clType ulong +syn keyword clType float +syn keyword clType double +syn keyword clType ptrdiff_t +syn keyword clType intptr_t +syn keyword clType uintptr_t + +" [6.1.2] Built-in Vector Data Types +syn match clType "\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" + +" [6.1.3] Other Built-in Data Types +syn keyword clType image2d_t +syn keyword clType image3d_t +syn keyword clType image2d_array_t +syn keyword clType image1d_t +syn keyword clType image1d_buffer_t +syn keyword clType image1d_array_t +syn keyword clType sampler_t +syn keyword clType event_t + +" [6.1.4] Reserved Data Types +syn keyword clType complex +syn keyword clType imaginary +syn match clType "float\(2\|3\|4\|8\|16\)x\(2\|3\|4\|8\|16\)" +syn match clType "double\(2\|3\|4\|8\|16\)x\(2\|3\|4\|8\|16\)" + +" [6.1.7] Vector Component Addressing +syn match clStructure "\.[w-z]\{1,4\}" +syn match clStructure "\.[s][0-9a-fA-F]\{1,16\}" +syn match clStructure "\.\(lo\|hi\|odd\|even\)" + +" [6.2] Conversions and Type Casting +syn match clFunction "convert_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)" +syn match clFunction "convert_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" +syn match clFunction "convert_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)_rt\(e\|p\|z\|n\)" +syn match clFunction "convert_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)_rt\(e\|p\|z\|n\)" +syn match clFunction "as_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)" +syn match clFunction "as_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" +syn match clFunction "as_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)_sat_rt\(e\|p\|z\|n\)" +syn match clFunction "as_\(bool\|half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)_sat_rt\(e\|p\|z\|n\)" + +" [6.5] Address Space Qualifiers +syn keyword clConstant __global +syn keyword clConstant global +syn keyword clConstant __constant +syn keyword clConstant constant +syn keyword clConstant __local +syn keyword clConstant local +syn keyword clConstant __private +syn keyword clConstant private +syn keyword clConstant __const +syn keyword clConstant const +syn keyword clConstant restrict +syn keyword clConstant volatile + + +" [6.6] Access Qualifiers +syn keyword clConstant __read_only +syn keyword clConstant read_only +syn keyword clConstant __write_only +syn keyword clConstant write_only + +" [6.7] Function Qualifiers +syn keyword clConstant __kernel +syn keyword clConstant kernel +syn keyword clConstant __attribute__ +syn keyword clFunction vec_type_hint +syn keyword clFunction work_group_size_hint +syn keyword clFunction reqd_work_group_size + +" [6.10] Preprocessor Directives and Macros +syn keyword clPreProc __FILE__ +syn keyword clPreProc __func__ +syn keyword clPreProc __LINE__ +syn keyword clPreProc __OPENCL_VERSION__ +syn keyword clPreProc __CL_VERSION_1_0__ +syn keyword clPreProc __CL_VERSION_1_1__ +syn keyword clPreProc __CL_VERSION_1_2__ +syn keyword clPreProc __OPENCL_C_VERSION__ +syn keyword clPreProc __ENDIAN_LITTLE__ +syn keyword clPreProc __IMAGE_SUPPORT__ +syn keyword clPreProc __FAST_RELAXED_MATH__ +syn keyword clPreProc FP_FAST_FMA +syn keyword clPreProc FP_FAST_FMAF +syn keyword clPreProc FP_FAST_FMA_HALF +syn keyword clFunction __kernel_exec +" OpenCL Extensions +syn keyword clConstant cl_khr_select_fprounding_mode +syn keyword clConstant cl_khr_global_int32_base_atomics +syn keyword clConstant cl_khr_global_int32_extended_atomics +syn keyword clConstant cl_khr_local_int32_base_atomics +syn keyword clConstant cl_khr_local_int32_extended_atomics +syn keyword clConstant cl_khr_int64_base_atomics +syn keyword clConstant cl_khr_int64_extended_atomics +syn keyword clConstant cl_khr_fp16 +syn keyword clConstant cl_khr_fp64 +syn keyword clConstant cl_khr_3d_image_writes +syn keyword clConstant cl_khr_byte_addressable_store +syn keyword clConstant CL_APPLE_gl_sharing +syn keyword clConstant CL_KHR_gl_sharing + + +" [6.11.1] Specify Type Attributes +syn keyword clFunction aligned +syn keyword clFunction packed +syn keyword clFunction endian + +" [6.12.2] Math Constants +syn keyword clConstant MAXFLOAT +syn keyword clConstant HUGE_VALF +syn keyword clConstant HUGE_VAL +syn keyword clConstant INFINITY +syn keyword clConstant NAN +syn keyword clConstant M_E_F +syn keyword clConstant M_LOG2E_F +syn keyword clConstant M_LOG10E_F +syn keyword clConstant M_LN2_F +syn keyword clConstant M_LN10_F +syn keyword clConstant M_PI_F +syn keyword clConstant M_PI_2_F +syn keyword clConstant M_PI_4_F +syn keyword clConstant M_1_PI_F +syn keyword clConstant M_2_PI_F +syn keyword clConstant M_2_SQRTPI_F +syn keyword clConstant M_SQRT2_F +syn keyword clConstant M_SQRT1_2_F + +" [6.12.1] Work-Item Built-in Functions +syn keyword clFunction get_work_dim +syn keyword clFunction get_global_size +syn keyword clFunction get_global_id +syn keyword clFunction get_local_size +syn keyword clFunction get_local_id +syn keyword clFunction get_num_groups +syn keyword clFunction get_group_id +syn keyword clFunction get_global_size + +" [6.12.2] Math Built-in Functions +syn keyword clFunction acos +syn keyword clFunction acosh +syn keyword clFunction acospi +syn keyword clFunction asin +syn keyword clFunction asinh +syn keyword clFunction asinpi +syn keyword clFunction atan +syn keyword clFunction atan2 +syn keyword clFunction atanh +syn keyword clFunction atanpi +syn keyword clFunction atan2pi +syn keyword clFunction cbrt +syn keyword clFunction ceil +syn keyword clFunction copysign +syn keyword clFunction cos +syn keyword clFunction cosh +syn keyword clFunction cospi +syn keyword clFunction erfc +syn keyword clFunction erf +syn keyword clFunction exp +syn keyword clFunction exp2 +syn keyword clFunction exp10 +syn keyword clFunction expm1 +syn keyword clFunction fabs +syn keyword clFunction fdim +syn keyword clFunction floor +syn keyword clFunction fma +syn keyword clFunction fmax +syn keyword clFunction fmin +syn keyword clFunction fmod +syn keyword clFunction fract +syn keyword clFunction frexp +syn keyword clFunction hypot +syn keyword clFunction ilogb +syn keyword clFunction ldexp +syn keyword clFunction lgamma +syn keyword clFunction lgamma_r +syn keyword clFunction log +syn keyword clFunction log2 +syn keyword clFunction log10 +syn keyword clFunction log1p +syn keyword clFunction logb +syn keyword clFunction mad +syn keyword clFunction modf +syn keyword clFunction nan +syn keyword clFunction nextafter +syn keyword clFunction pow +syn keyword clFunction pown +syn keyword clFunction powr +syn keyword clFunction remainder +syn keyword clFunction remquo +syn keyword clFunction rint +syn keyword clFunction rootn +syn keyword clFunction round +syn keyword clFunction rsqrt +syn keyword clFunction sin +syn keyword clFunction sincos +syn keyword clFunction sinh +syn keyword clFunction sinpi +syn keyword clFunction sqrt +syn keyword clFunction tan +syn keyword clFunction tanh +syn keyword clFunction tanpi +syn keyword clFunction tgamma +syn keyword clFunction trunc +syn match clFunction "\(half\|native\)_\(cos\|divide\|exp\|exp2\|exp10\|log\|log2\|log10\|powr\|recip\|rsqrt\|sin\|sqrt\|tan\)" + +" [6.12.3] Integer Built-in Functions +syn keyword clFunction abs +syn keyword clFunction abs_diff +syn keyword clFunction add_sat +syn keyword clFunction hadd +syn keyword clFunction rhadd +syn keyword clFunction clamp +syn keyword clFunction clz +syn keyword clFunction mad_hi +syn keyword clFunction mad24 +syn keyword clFunction mad_sat +syn keyword clFunction max +syn keyword clFunction min +syn keyword clFunction mul_hi +syn keyword clFunction rotate +syn keyword clFunction sub_sat +syn keyword clFunction popcount +syn keyword clFunction upsample +syn keyword clFunction mad24 +syn keyword clFunction mul24 + +" [6.12.4] Common Built-in Functions +syn keyword clFunction degrees +syn keyword clFunction mix +syn keyword clFunction radians +syn keyword clFunction step +syn keyword clFunction smoothstep +syn keyword clFunction sign + +" [6.12.5] Geometric Built-in Functions +syn keyword clFunction clFunction +syn keyword clFunction cross +syn keyword clFunction dot +syn keyword clFunction distance +syn keyword clFunction length +syn keyword clFunction normalize +syn keyword clFunction fast_distance +syn keyword clFunction fast_length +syn keyword clFunction fast_normalize + +" [6.12.6] Relational Built-in Functions +syn keyword clFunction isequal +syn keyword clFunction isnotequal +syn keyword clFunction isgreater +syn keyword clFunction isgreaterequal +syn keyword clFunction isless +syn keyword clFunction islessequal +syn keyword clFunction islessgreater +syn keyword clFunction isfinite +syn keyword clFunction isinf +syn keyword clFunction isnan +syn keyword clFunction isnormal +syn keyword clFunction isordered +syn keyword clFunction isunordered +syn keyword clFunction signbit +syn keyword clFunction any +syn keyword clFunction all +syn keyword clFunction bitselect +syn keyword clFunction select + +" [6.12.7] Vector Data Load/Store +syn keyword clFunction vload vstore vload_half vstore_half vloada_half vstorea_half +syn match clFunction "\(vload\|vstore\|vload_half\|vstore_half\|vloada_half\|vstorea_half\)\(2\|3\|4\|8\|16\)" + +" [6.12.8, 6.12.9] Synchronization and Explicit Memory Fence Functions +syn keyword clFunction mem_fence +syn keyword clFunction read_mem_fence +syn keyword clFunction write_mem_fence +syn keyword clFunction barrier +syn keyword clConstant CLK_LOCAL_MEM_FENCE +syn keyword clConstant CLK_GLOBAL_MEM_FENCE + +" [6.12.10] Async Copies and Prefetch Functions +syn keyword clFunction async_work_group_copy +syn keyword clFunction async_work_group_strided_copy +syn keyword clFunction wait_group_events +syn keyword clFunction prefetch + +" [6.12.11] Atomic Functions +syn keyword clFunction atom_add +syn keyword clFunction atom_sub +syn keyword clFunction atom_xchg +syn keyword clFunction atom_inc +syn keyword clFunction atom_dec +syn keyword clFunction atom_cmpxchg +syn keyword clFunction atom_min +syn keyword clFunction atom_max +syn keyword clFunction atom_and +syn keyword clFunction atom_or +syn keyword clFunction atom_xor + +" [6.12.12] Miscellaneous Vector Functions +syn keyword clFunction vec_step +syn keyword clFunction shuffle +syn keyword clFunction shuffle2 + +" [6.12.13] printf Function +syn keyword clFunction printf + +" [6.12.14] Image Read and Write Built-in Functions +syn match clFunction "read_image\(f\|i\|ui\|h\)" +syn match clFunction "write_image\(f\|i\|ui\|h\)" + +" [6.12.14.5] Image Query Functions +syn match clFunction "get_image_\(width\|height\|depth\|array_size\|dim\|channel_data_type\|channel_order\)" + +" [6.12.14.1] Sampler Declaration Fields +syn match clConstant "CLK_ADDRESS_\(REPEAT\|CLAMP\|NONE\)" +syn match clConstant "CLK_ADDRESS_\(CLAMP_TO_EDGE\|MIRRORED_REPEAT\)" +syn match clConstant "CLK_NORMALIZED_COORDS_\(TRUE\|FALSE\)" +syn keyword clConstant CLK_FILTER_NEAREST +syn keyword clConstant CLK_FILTER_LINEAR + +let b:current_syntax = "cl" + +hi def link clType Type +hi def link clPreProc PreProc +hi def link clFunction Function +hi def link clConstant Constant +hi def link clStructure Structure diff --git a/doc/vim_syntax_highlighting/syntax/coprthr_kern.vim b/doc/vim_syntax_highlighting/syntax/coprthr_kern.vim new file mode 100644 index 00000000..43abdb76 --- /dev/null +++ b/doc/vim_syntax_highlighting/syntax/coprthr_kern.vim @@ -0,0 +1,75 @@ +" Vim syntax file +" Language: COPRTHR OpenCL kernel code +" Maintainer: James Ross +" Version: 0.2 +" Last Change: 2013-03-25 +" +" Description: Syntax highlighting for COPRTHR OpenCL kernels + + +"if exists("b:current_syntax") +" finish +"endif + +if version < 600 + so :p:h/c.vim +else + runtime! syntax/c.vim + unlet b:current_syntax +endif + + +" Types +syn keyword coprthrType char uchar short ushort int uint long ulong float double +syn match coprthrType "\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" +syn keyword coprthrType image2d_t +syn keyword coprthrType sampler_t + +" Constants +syn match coprthrConstant "CLK_ADDRESS_\(REPEAT\|CLAMP\|NONE\)" +syn match coprthrConstant "CLK_ADDRESS_\(CLAMP_TO_EDGE\|MIRRORED_REPEAT\)" +syn match coprthrConstant "CLK_NORMALIZED_COORDS_\(TRUE\|FALSE\)" +syn keyword coprthrConstant CLK_LOCAL_MEM_FENCE + +" Pre Processor +syn keyword coprthrPreProc global +syn keyword coprthrPreProc local +syn keyword coprthrPreProc constant +syn keyword coprthrPreProc private +syn keyword coprthrPreProc __global +syn keyword coprthrPreProc __local +syn keyword coprthrPreProc __constant +syn keyword coprthrPreProc __private +syn keyword coprthrPreProc __read_only +syn keyword coprthrPreProc __write_only +syn keyword coprthrPreProc __kernel +syn keyword coprthrPreProc restrict +syn keyword coprthrPreProc __restrict +syn match coprthrPreProc "__builtin_vector_\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" +syn match coprthrPreProc "vector_\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" + +" Functions +syn keyword coprthrFunction reqd_work_group_size +syn keyword coprthrFunction get_work_dim +syn keyword coprthrFunction get_global_size +syn keyword coprthrFunction get_global_id +syn keyword coprthrFunction get_local_size +syn keyword coprthrFunction get_local_id +syn keyword coprthrFunction get_num_groups +syn keyword coprthrFunction get_group_id +syn keyword coprthrFunction sqrt acos acosh asin asinh tan tanh cbrt ceil cos cosh erfc erf exp exp2 exp10 expm1 fabs floor log log2 log10 log1p logb rint round rsqrt sin sinh sqrt tan tanh tgamma tfunc clamp dot normalize cross read_imagei read_imageui read_imagef +syn match coprthrFunction "read_image\(f\|i\|ui\|h\)" +syn match coprthrFunction "write_image\(f\|i\|ui\|h\)" + +" Vector Component Addressing +syn match coprthrStructure "\.[w-z]\{1,4\}" +syn match coprthrStructure "\.[s][0-9a-fA-F]\{1,16\}" +syn match coprthrStructure "\.\(lo\|hi\)" + +let b:current_syntax = "coprthr_kern" + +hi def link coprthrType Type +hi def link coprthrPreProc PreProc +hi def link coprthrFunction Function +hi def link coprthrConstant Constant +hi def link coprthrStructure Structure diff --git a/doc/vim_syntax_highlighting/syntax/opencl.vim b/doc/vim_syntax_highlighting/syntax/opencl.vim new file mode 100644 index 00000000..835c0cc4 --- /dev/null +++ b/doc/vim_syntax_highlighting/syntax/opencl.vim @@ -0,0 +1,671 @@ +" Vim syntax file +" Language: C OpenCL +" Maintainer: James Ross +" Version: 0.5 +" Last Change: 2013-03-25 +" +" Description: Syntax highlighting for OpenCL 1.2 + + +if exists("b:current_syntax") + finish +endif + +if version < 600 + so :p:h/c.vim +else + runtime! syntax/c.vim + unlet b:current_syntax +endif + +" Enumerated Types +syn keyword clType cl_channel_order +syn keyword clType cl_channel_type +syn keyword clType cl_command_queue_info +syn keyword clType cl_command_queue_properties +syn keyword clType cl_command_type +syn keyword clType cl_context_info +syn keyword clType cl_context_properties +syn keyword clType cl_device_exec_capabilities +syn keyword clType cl_device_fp_config +syn keyword clType cl_device_info +syn keyword clType cl_device_local_mem_type +syn keyword clType cl_device_mem_cache_type +syn keyword clType cl_device_type +syn keyword clType cl_event_info +syn keyword clType cl_filter_mode +syn keyword clType cl_gl_object_type +syn keyword clType cl_gl_texture_info +syn keyword clType cl_gl_platform_info +syn keyword clType cl_gl_context_info +syn keyword clType cl_GLsync +syn keyword clType cl_GLuint +syn keyword clType cl_GLint +syn keyword clType cl_GLenum +syn keyword clType cl_image_info +syn keyword clType cl_image_format +syn keyword clType cl_kernel_info +syn keyword clType cl_kernel_work_group_info +syn keyword clType cl_map_flags cl_mem_flags +syn keyword clType cl_mem_info +syn keyword clType cl_mem_object_type +syn keyword clType cl_platform_info +syn keyword clType cl_profiling_info +syn keyword clType cl_program_build_info +syn keyword clType cl_sampler_info +syn keyword clType cl_bool +syn keyword clType cl_bitfield +syn keyword clType cl_device_partition_property +syn keyword clType cl_device_partition_property_ext +syn keyword clType cl_device_affinity_domain +syn keyword clType cl_mem_migration_flags +syn keyword clType cl_buffer_create_type +syn keyword clType cl_addressing_mode +syn keyword clType cl_program_info +syn keyword clType cl_program_binary_type +syn keyword clType cl_build_status +syn keyword clType cl_kernel_arg_info +syn keyword clType cl_kernel_arg_address_qualifier +syn keyword clType cl_kernel_arg_access_qualifier +syn keyword clType cl_kernel_arg_type_qualifier +syn keyword clType cl_image_desc +syn keyword clType cl_buffer_region + +" Types +syn keyword clType cl_half +syn keyword clType cl_quad +syn keyword clType cl_char +syn keyword clType cl_uchar +syn keyword clType cl_short +syn keyword clType cl_ushort +syn keyword clType cl_int +syn keyword clType cl_uint +syn keyword clType cl_ulong +syn keyword clType cl_float +syn keyword clType cl_double +syn match clType "cl_\(half\|quad\|char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\|16\)" + +syn match clStructure "\.[w-z]\{1\}" +syn match clStructure "\.[s][0-9a-fA-F]\{1\}" +syn match clStructure "\.\(lo\|hi\)" + +syn keyword clType image2d_t +syn keyword clType image3d_t +syn keyword clType sampler_t +syn keyword clType event_t +syn keyword clType cl_platform_id +syn keyword clType cl_device_id +syn keyword clType cl_context +syn keyword clType cl_command_queue +syn keyword clType cl_mem +syn keyword clType cl_program +syn keyword clType cl_kernel +syn keyword clType cl_event +syn keyword clType cl_sampler + +" Constants {{{ +syn keyword clConstant CL_VERSION_1_0 +syn keyword clConstant CL_VERSION_1_1 +syn keyword clConstant CL_VERSION_1_2 +syn keyword clConstant CL_ADDRESS_NONE +syn keyword clConstant CL_ADDRESS_CLAMP_TO_EDGE +syn keyword clConstant CL_ADDRESS_CLAMP +syn keyword clConstant CL_ADDRESS_REPEAT +syn keyword clConstant CL_ADDRESS_MIRRORED_REPEAT +syn keyword clConstant CL_FALSE +syn keyword clConstant CL_TRUE +syn keyword clConstant CL_BLOCKING +syn keyword clConstant CL_NON_BLOCKING +syn keyword clConstant CL_PLATFORM_PROFILE +syn keyword clConstant CL_PLATFORM_VERSION +syn keyword clConstant CL_PLATFORM_NAME +syn keyword clConstant CL_PLATFORM_VENDOR +syn keyword clConstant CL_PLATFORM_EXTENSIONS +syn keyword clConstant CL_BUILD_ERROR +syn keyword clConstant CL_BUILD_IN_PROGRESS +syn keyword clConstant CL_BUILD_NONE +syn keyword clConstant CL_BUILD_SUCCESS +syn keyword clConstant CL_A +syn keyword clConstant CL_R +syn keyword clConstant CL_RG +syn keyword clConstant CL_RA +syn keyword clConstant CL_RGB +syn keyword clConstant CL_RGBA +syn keyword clConstant CL_BGRA +syn keyword clConstant CL_ARGB +syn keyword clConstant CL_INTENSITY +syn keyword clConstant CL_LUMINANCE +syn keyword clConstant CL_Rx +syn keyword clConstant CL_RGx +syn keyword clConstant CL_RGBx +syn keyword clConstant CL_FLOAT +syn keyword clConstant CL_HALF_FLOAT +syn keyword clConstant CL_SIGNED_INT16 +syn keyword clConstant CL_SIGNED_INT32 +syn keyword clConstant CL_SIGNED_INT8 +syn keyword clConstant CL_SNORM_INT8 +syn keyword clConstant CL_SNORM_INT16 +syn keyword clConstant CL_UNORM_INT_101010 +syn keyword clConstant CL_UNORM_INT16 +syn keyword clConstant CL_UNORM_INT8 +syn keyword clConstant CL_UNORM_SHORT_555 +syn keyword clConstant CL_UNORM_SHORT_565 +syn keyword clConstant CL_UNSIGNED_INT16 +syn keyword clConstant CL_UNSIGNED_INT32 +syn keyword clConstant CL_UNSIGNED_INT8 +syn keyword clConstant CL_QUEUE_CONTEXT +syn keyword clConstant CL_QUEUE_DEVICE +syn keyword clConstant CL_QUEUE_PROPERTIES +syn keyword clConstant CL_QUEUE_REFERENCE_COUNT +syn keyword clConstant CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE +syn keyword clConstant CL_QUEUE_PROFILING_ENABLE +syn keyword clConstant CL_COMMAND_ACQUIRE_GL_OBJECTS +syn keyword clConstant CL_COMMAND_COPY_BUFFER +syn keyword clConstant CL_COMMAND_COPY_BUFFER_TO_IMAGE +syn keyword clConstant CL_COMMAND_COPY_IMAGE +syn keyword clConstant CL_COMMAND_COPY_IMAGE_TO_BUFFER +syn keyword clConstant CL_COMMAND_MAP_BUFFER +syn keyword clConstant CL_COMMAND_MAP_IMAGE +syn keyword clConstant CL_COMMAND_MARKER +syn keyword clConstant CL_COMMAND_NATIVE_KERNEL +syn keyword clConstant CL_COMMAND_NDRANGE_KERNEL +syn keyword clConstant CL_COMMAND_READ_BUFFER +syn keyword clConstant CL_COMMAND_READ_IMAGE +syn keyword clConstant CL_COMMAND_RELEASE_GL_OBJECTS +syn keyword clConstant CL_COMMAND_TASK +syn keyword clConstant CL_COMMAND_UNMAP_MEM_OBJECT +syn keyword clConstant CL_COMMAND_WRITE_BUFFER +syn keyword clConstant CL_COMMAND_WRITE_IMAGE +syn keyword clConstant CL_COMMAND_READ_BUFFER_RECT +syn keyword clConstant CL_COMMAND_WRITE_BUFFER_RECT +syn keyword clConstant CL_COMMAND_COPY_BUFFER_RECT +syn keyword clConstant CL_COMMAND_USER +syn keyword clConstant CL_COMMAND_BARRIER +syn keyword clConstant CL_COMMAND_MIGRATE_MEM_OBJECTS +syn keyword clConstant CL_COMMAND_FILL_BUFFER +syn keyword clConstant CL_COMMAND_FILL_IMAGE +syn keyword clConstant CL_CONTEXT_DEVICES +syn keyword clConstant CL_CONTEXT_PROPERTIES +syn keyword clConstant CL_CONTEXT_NUM_DEVICES +syn keyword clConstant CL_CONTEXT_REFERENCE_COUNT +syn keyword clConstant CL_CONTEXT_PLATFORM +syn keyword clConstant CL_CONTEXT_INTEROP_USER_SYNC +syn keyword clConstant CL_DEVICE_PARTITION_EQUALLY +syn keyword clConstant CL_DEVICE_PARTITION_BY_COUNTS +syn keyword clConstant CL_DEVICE_PARTITION_BY_COUNTS_LIST_END +syn keyword clConstant CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_NUMA +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE +syn keyword clConstant CL_EXEC_KERNEL +syn keyword clConstant CL_EXEC_NATIVE_KERNEL +syn keyword clConstant CL_FP_DENORM +syn keyword clConstant CL_FP_FMA +syn keyword clConstant CL_FP_SOFT_FLOAT +syn keyword clConstant CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT +syn keyword clConstant CL_FP_INF_NAN +syn keyword clConstant CL_FP_ROUND_TO_INF +syn keyword clConstant CL_FP_ROUND_TO_NEAREST +syn keyword clConstant CL_FP_ROUND_TO_ZERO +syn keyword clConstant CL_DEVICE_ADDRESS_BITS +syn keyword clConstant CL_DEVICE_AVAILABLE +syn keyword clConstant CL_DEVICE_COMPILER_AVAILABLE +syn keyword clConstant CL_DEVICE_ENDIAN_LITTLE +syn keyword clConstant CL_DEVICE_ERROR_CORRECTION_SUPPORT +syn keyword clConstant CL_DEVICE_EXECUTION_CAPABILITIES +syn keyword clConstant CL_DEVICE_EXTENSIONS +syn keyword clConstant CL_DEVICE_GLOBAL_MEM_CACHE_SIZE +syn keyword clConstant CL_DEVICE_GLOBAL_MEM_CACHE_TYPE +syn keyword clConstant CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE +syn keyword clConstant CL_DEVICE_GLOBAL_MEM_SIZE +syn keyword clConstant CL_DEVICE_IMAGE_SUPPORT +syn keyword clConstant CL_DEVICE_IMAGE2D_MAX_HEIGHT +syn keyword clConstant CL_DEVICE_IMAGE2D_MAX_WIDTH +syn keyword clConstant CL_DEVICE_IMAGE3D_MAX_DEPTH +syn keyword clConstant CL_DEVICE_IMAGE3D_MAX_HEIGHT +syn keyword clConstant CL_DEVICE_IMAGE3D_MAX_WIDTH +syn keyword clConstant CL_DEVICE_LOCAL_MEM_SIZE +syn keyword clConstant CL_DEVICE_LOCAL_MEM_TYPE +syn keyword clConstant CL_DEVICE_MAX_CLOCK_FREQUENCY +syn keyword clConstant CL_DEVICE_MAX_COMPUTE_UNITS +syn keyword clConstant CL_DEVICE_MAX_CONSTANT_ARGS +syn keyword clConstant CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE +syn keyword clConstant CL_DEVICE_MAX_MEM_ALLOC_SIZE +syn keyword clConstant CL_DEVICE_MAX_PARAMETER_SIZE +syn keyword clConstant CL_DEVICE_MAX_READ_IMAGE_ARGS +syn keyword clConstant CL_DEVICE_MAX_SAMPLERS +syn keyword clConstant CL_DEVICE_MAX_WORK_GROUP_SIZE +syn keyword clConstant CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS +syn keyword clConstant CL_DEVICE_MAX_WORK_ITEM_SIZES +syn keyword clConstant CL_DEVICE_MAX_WRITE_IMAGE_ARGS +syn keyword clConstant CL_DEVICE_MEM_BASE_ADDR_ALIGN +syn keyword clConstant CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE +syn keyword clConstant CL_DEVICE_NAME +syn keyword clConstant CL_DEVICE_PLATFORM +syn keyword clConstant CL_DEVICE_DOUBLE_FP_CONFIG +syn keyword clConstant CL_DEVICE_HALF_FP_CONFIG +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF +syn keyword clConstant CL_DEVICE_HOST_UNIFIED_MEMORY +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_INT +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE +syn keyword clConstant CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF +syn keyword clConstant CL_DEVICE_OPENCL_C_VERSION +syn keyword clConstant CL_DEVICE_LINKER_AVAILABLE +syn keyword clConstant CL_DEVICE_BUILT_IN_KERNELS +syn keyword clConstant CL_DEVICE_IMAGE_MAX_BUFFER_SIZE +syn keyword clConstant CL_DEVICE_IMAGE_MAX_ARRAY_SIZE +syn keyword clConstant CL_DEVICE_PARENT_DEVICE +syn keyword clConstant CL_DEVICE_PARENT_DEVICE_EXT +syn keyword clConstant CL_DEVICE_PARTITION_MAX_SUB_DEVICES +syn keyword clConstant CL_DEVICE_PARTITION_PROPERTIES +syn keyword clConstant CL_DEVICE_PARTITION_PROPERTIES_EXT +syn keyword clConstant CL_DEVICE_PARTITION_AFFINITY_DOMAIN +syn keyword clConstant CL_DEVICE_AFFINITY_DOMAINS_EXT +syn keyword clConstant CL_DEVICE_PARTITION_TYPE +syn keyword clConstant CL_DEVICE_PARTITION_TYPES_EXT +syn keyword clConstant CL_DEVICE_PARTITION_STYLE_EXT +syn keyword clConstant CL_DEVICE_REFERENCE_COUNT +syn keyword clConstant CL_DEVICE_REFERENCE_COUNT_EXT +syn keyword clConstant CL_DEVICE_PREFERRED_INTEROP_USER_SYNC +syn keyword clConstant CL_DEVICE_PRINTF_BUFFER_SIZE +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG +syn keyword clConstant CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT +syn keyword clConstant CL_DEVICE_PROFILE +syn keyword clConstant CL_DEVICE_PROFILING_TIMER_RESOLUTION +syn keyword clConstant CL_DEVICE_QUEUE_PROPERTIES +syn keyword clConstant CL_DEVICE_SINGLE_FP_CONFIG +syn keyword clConstant CL_DEVICE_TYPE +syn keyword clConstant CL_DEVICE_VENDOR_ID +syn keyword clConstant CL_DEVICE_VENDOR +syn keyword clConstant CL_DEVICE_VERSION +syn keyword clConstant CL_DRIVER_VERSION +syn keyword clConstant CL_GLOBAL +syn keyword clConstant CL_LOCAL +syn keyword clConstant CL_NONE +syn keyword clConstant CL_READ_ONLY_CACHE +syn keyword clConstant CL_READ_WRITE_CACHE +syn keyword clConstant CL_DEVICE_TYPE_DEFAULT +syn keyword clConstant CL_DEVICE_TYPE_CPU +syn keyword clConstant CL_DEVICE_TYPE_GPU +syn keyword clConstant CL_DEVICE_TYPE_ACCELERATOR +syn keyword clConstant CL_DEVICE_TYPE_CUSTOM +syn keyword clConstant CL_DEVICE_TYPE_ALL +syn keyword clConstant CL_EVENT_COMMAND_QUEUE +syn keyword clConstant CL_EVENT_COMMAND_TYPE +syn keyword clConstant CL_EVENT_REFERENCE_COUNT +syn keyword clConstant CL_EVENT_COMMAND_EXECUTION_STATUS +syn keyword clConstant CL_EVENT_CONTEXT +syn keyword clConstant CL_FILTER_NEAREST +syn keyword clConstant CL_FILTER_LINEAR +syn keyword clConstant CL_SAMPLER_REFERENCE_COUNT +syn keyword clConstant CL_SAMPLER_CONTEXT +syn keyword clConstant CL_SAMPLER_NORMALIZED_COORDS +syn keyword clConstant CL_SAMPLER_ADDRESSING_MODE +syn keyword clConstant CL_SAMPLER_FILTER_MODE +syn keyword clConstant CL_PROGRAM_REFERENCE_COUNT +syn keyword clConstant CL_PROGRAM_CONTEXT +syn keyword clConstant CL_PROGRAM_NUM_DEVICES +syn keyword clConstant CL_PROGRAM_DEVICES +syn keyword clConstant CL_PROGRAM_SOURCE +syn keyword clConstant CL_PROGRAM_BINARY_SIZES +syn keyword clConstant CL_PROGRAM_BINARIES +syn keyword clConstant CL_PROGRAM_NUM_KERNELS +syn keyword clConstant CL_PROGRAM_KERNEL_NAMES +syn keyword clConstant CL_PROGRAM_BUILD_STATUS +syn keyword clConstant CL_PROGRAM_BUILD_OPTIONS +syn keyword clConstant CL_PROGRAM_BUILD_LOG +syn keyword clConstant CL_PROGRAM_BINARY_TYPE +syn keyword clConstant CL_PROGRAM_BINARY_TYPE_NONE +syn keyword clConstant CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT +syn keyword clConstant CL_PROGRAM_BINARY_TYPE_LIBRARY +syn keyword clConstant CL_PROGRAM_BINARY_TYPE_EXECUTABLE +syn keyword clConstant CL_GL_OBJECT_BUFFER +syn keyword clConstant CL_GL_OBJECT_TEXTURE2D +syn keyword clConstant CL_GL_OBJECT_TEXTURE3D +syn keyword clConstant CL_GL_OBJECT_RENDERBUFFER +syn keyword clConstant CL_GL_OBJECT_TEXTURE2D_ARRAY +syn keyword clConstant CL_GL_OBJECT_TEXTURE1D +syn keyword clConstant CL_GL_OBJECT_TEXTURE1D_ARRAY +syn keyword clConstant CL_GL_OBJECT_TEXTURE_BUFFER +syn keyword clConstant CL_GL_TEXTURE_TARGET +syn keyword clConstant CL_GL_MIPMAP_LEVEL +syn keyword clConstant CL_IMAGE_FORMAT +syn keyword clConstant CL_IMAGE_ELEMENT_SIZE +syn keyword clConstant CL_IMAGE_ROW_PITCH +syn keyword clConstant CL_IMAGE_SLICE_PITCH +syn keyword clConstant CL_IMAGE_WIDTH +syn keyword clConstant CL_IMAGE_HEIGHT +syn keyword clConstant CL_IMAGE_DEPTH +syn keyword clConstant CL_IMAGE_ARRAY_SIZE +syn keyword clConstant CL_IMAGE_BUFFER +syn keyword clConstant CL_IMAGE_NUM_MIP_LEVELS +syn keyword clConstant CL_IMAGE_NUM_SAMPLES +syn keyword clConstant CL_KERNEL_FUNCTION_NAME +syn keyword clConstant CL_KERNEL_NUM_ARGS +syn keyword clConstant CL_KERNEL_REFERENCE_COUNT +syn keyword clConstant CL_KERNEL_CONTEXT +syn keyword clConstant CL_KERNEL_PROGRAM +syn keyword clConstant CL_KERNEL_ATTRIBUTES +syn keyword clConstant CL_KERNEL_WORK_GROUP_SIZE +syn keyword clConstant CL_KERNEL_COMPILE_WORK_GROUP_SIZE +syn keyword clConstant CL_KERNEL_LOCAL_MEM_SIZE +syn keyword clConstant CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE +syn keyword clConstant CL_KERNEL_PRIVATE_MEM_SIZE +syn keyword clConstant CL_KERNEL_GLOBAL_WORK_SIZE +syn keyword clConstant CL_KERNEL_ARG_ADDRESS_QUALIFIER +syn keyword clConstant CL_KERNEL_ARG_ACCESS_QUALIFIER +syn keyword clConstant CL_KERNEL_ARG_TYPE_NAME +syn keyword clConstant CL_KERNEL_ARG_TYPE_QUALIFIER +syn keyword clConstant CL_KERNEL_ARG_NAME +syn keyword clConstant CL_KERNEL_ARG_ADDRESS_GLOBAL +syn keyword clConstant CL_KERNEL_ARG_ADDRESS_LOCAL +syn keyword clConstant CL_KERNEL_ARG_ADDRESS_CONSTANT +syn keyword clConstant CL_KERNEL_ARG_ADDRESS_PRIVATE +syn keyword clConstant CL_KERNEL_ARG_ACCESS_READ_ONLY +syn keyword clConstant CL_KERNEL_ARG_ACCESS_WRITE_ONLY +syn keyword clConstant CL_KERNEL_ARG_ACCESS_READ_WRITE +syn keyword clConstant CL_KERNEL_ARG_ACCESS_NONE +syn keyword clConstant CL_KERNEL_ARG_TYPE_NONE +syn keyword clConstant CL_KERNEL_ARG_TYPE_CONST +syn keyword clConstant CL_KERNEL_ARG_TYPE_RESTRICT +syn keyword clConstant CL_KERNEL_ARG_TYPE_VOLATILE +syn keyword clConstant CL_MAP_READ +syn keyword clConstant CL_MAP_WRITE +syn keyword clConstant CL_MAP_WRITE_INVALIDATE_REGION +syn keyword clConstant CL_MEM_READ_WRITE +syn keyword clConstant CL_MEM_WRITE_ONLY +syn keyword clConstant CL_MEM_READ_ONLY +syn keyword clConstant CL_MEM_USE_HOST_PTR +syn keyword clConstant CL_MEM_ALLOC_HOST_PTR +syn keyword clConstant CL_MEM_COPY_HOST_PTR +syn keyword clConstant CL_MEM_HOST_WRITE_ONLY +syn keyword clConstant CL_MEM_HOST_READ_ONLY +syn keyword clConstant CL_MEM_HOST_NO_ACCESS +syn keyword clConstant CL_MIGRATE_MEM_OBJECT_HOST +syn keyword clConstant CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED +syn keyword clConstant CL_MEM_TYPE +syn keyword clConstant CL_MEM_FLAGS +syn keyword clConstant CL_MEM_SIZE +syn keyword clConstant CL_MEM_HOST_PTR +syn keyword clConstant CL_MEM_MAP_COUNT +syn keyword clConstant CL_MEM_REFERENCE_COUNT +syn keyword clConstant CL_MEM_CONTEXT +syn keyword clConstant CL_MEM_ASSOCIATED_MEMOBJECT +syn keyword clConstant CL_MEM_OFFSET +syn keyword clConstant CL_MEM_OBJECT_BUFFER +syn keyword clConstant CL_MEM_OBJECT_IMAGE2D +syn keyword clConstant CL_MEM_OBJECT_IMAGE3D +syn keyword clConstant CL_MEM_OBJECT_IMAGE2D_ARRAY +syn keyword clConstant CL_MEM_OBJECT_IMAGE1D +syn keyword clConstant CL_MEM_OBJECT_IMAGE1D_ARRAY +syn keyword clConstant CL_MEM_OBJECT_IMAGE1D_BUFFER +syn keyword clConstant CL_COMPLETE +syn keyword clConstant CL_RUNNING +syn keyword clConstant CL_SUBMITTED +syn keyword clConstant CL_QUEUED +syn keyword clConstant CL_BUFFER_CREATE_TYPE_REGION +syn keyword clConstant CL_PROFILING_COMMAND_QUEUED +syn keyword clConstant CL_PROFILING_COMMAND_SUBMIT +syn keyword clConstant CL_PROFILING_COMMAND_START +syn keyword clConstant CL_PROFILING_COMMAND_END +syn keyword clConstant CL_CHAR_BIT +syn keyword clConstant CL_SCHAR_MAX +syn keyword clConstant CL_SCHAR_MIN +syn keyword clConstant CL_CHAR_MAX +syn keyword clConstant CL_CHAR_MIN +syn keyword clConstant CL_UCHAR_MAX +syn keyword clConstant CL_SHRT_MAX +syn keyword clConstant CL_SHRT_MIN +syn keyword clConstant CL_USHRT_MAX +syn keyword clConstant CL_INT_MAX +syn keyword clConstant CL_INT_MIN +syn keyword clConstant CL_UINT_MAX +syn keyword clConstant CL_LONG_MAX +syn keyword clConstant CL_LONG_MIN +syn keyword clConstant CL_ULONG_MAX +syn keyword clConstant CL_FLT_DIG +syn keyword clConstant CL_FLT_MANT_DIG +syn keyword clConstant CL_FLT_MAX_10_EXP +syn keyword clConstant CL_FLT_MAX_EXP +syn keyword clConstant CL_FLT_MIN_10_EXP +syn keyword clConstant CL_FLT_MIN_EXP +syn keyword clConstant CL_FLT_RADIX +syn keyword clConstant CL_FLT_MAX +syn keyword clConstant CL_FLT_MIN +syn keyword clConstant CL_FLT_EPSILON +syn keyword clConstant CL_DBL_DIG +syn keyword clConstant CL_DBL_MANT_DIG +syn keyword clConstant CL_DBL_MAX_10_EXP +syn keyword clConstant CL_DBL_MAX_EXP +syn keyword clConstant CL_DBL_MIN_10_EXP +syn keyword clConstant CL_DBL_MIN_EXP +syn keyword clConstant CL_DBL_RADIX +syn keyword clConstant CL_DBL_MAX +syn keyword clConstant CL_DBL_MIN +syn keyword clConstant CL_DBL_EPSILON +syn keyword clConstant CL_M_E +syn keyword clConstant CL_M_LOG2E +syn keyword clConstant CL_M_LOG10E +syn keyword clConstant CL_M_LN2 +syn keyword clConstant CL_M_LN10 +syn keyword clConstant CL_M_PI +syn keyword clConstant CL_M_PI_2 +syn keyword clConstant CL_M_PI_4 +syn keyword clConstant CL_M_1_PI +syn keyword clConstant CL_M_2_PI +syn keyword clConstant CL_M_2_SQRTPI +syn keyword clConstant CL_M_SQRT2 +syn keyword clConstant CL_M_SQRT1_2 +syn keyword clConstant CL_M_E_F +syn keyword clConstant CL_M_LOG2E_F +syn keyword clConstant CL_M_LOG10E_F +syn keyword clConstant CL_M_LN2_F +syn keyword clConstant CL_M_LN10_F +syn keyword clConstant CL_M_PI_F +syn keyword clConstant CL_M_PI_2_F +syn keyword clConstant CL_M_PI_4_F +syn keyword clConstant CL_M_1_PI_F +syn keyword clConstant CL_M_2_PI_F +syn keyword clConstant CL_M_2_SQRTPI_F +syn keyword clConstant CL_M_SQRT2_F +syn keyword clConstant CL_M_SQRT1_2_F +syn keyword clConstant CL_NAN +syn keyword clConstant CL_HUGE_VALF +syn keyword clConstant CL_HUGE_VAL +syn keyword clConstant CL_MAXFLOAT +syn keyword clConstant CL_INFINITY +" }}} + +" Functions +syn keyword clFunction clGetPlatformIDs +syn keyword clFunction clGetPlatformInfo +syn keyword clFunction clGetDeviceIDs +syn keyword clFunction clGetDeviceInfo +syn keyword clFunction clCreateSubDevices +syn keyword clFunction clCreateSubDevicesEXT +syn keyword clFunction clRetainDevice +syn keyword clFunction clReleaseDevice +syn keyword clFunction clCreateContext +syn keyword clFunction clCreateContextFromType +syn keyword clFunction clRetainContext +syn keyword clFunction clReleaseContext +syn keyword clFunction clGetContextInfo +syn keyword clFunction clGetGLContextInfoKHR +syn keyword clFunction clCreateCommandQueue +syn keyword clFunction clRetainCommandQueue +syn keyword clFunction clReleaseCommandQueue +syn keyword clFunction clGetCommandQueueInfo +syn keyword clFunction clSetCommandQueueProperty +syn keyword clFunction clCreateBuffer +syn keyword clFunction clCreateSubBuffer +syn keyword clFunction clCreateImage +syn keyword clFunction clCreateFromGLTexture +syn keyword clFunction clCreateFromGLRenderbuffer +syn keyword clFunction clEnqueueReadBuffer +syn keyword clFunction clEnqueueReadBufferRect +syn keyword clFunction clEnqueueWriteBuffer +syn keyword clFunction clEnqueueWriteBufferRect +syn keyword clFunction clEnqueueFillBuffer +syn keyword clFunction clEnqueueCopyBuffer +syn keyword clFunction clEnqueueCopyBufferRect +syn keyword clFunction clRetainMemObject +syn keyword clFunction clReleaseMemObject +syn keyword clFunction clCreateImage2D +syn keyword clFunction clCreateImage3D +syn keyword clFunction clGetSupportedImageFormats +syn keyword clFunction clEnqueueReadImage +syn keyword clFunction clEnqueueWriteImage +syn keyword clFunction clEnqueueFillImage +syn keyword clFunction clEnqueueCopyImage +syn keyword clFunction clEnqueueCopyImageToBuffer +syn keyword clFunction clEnqueueCopyBufferToImage +syn keyword clFunction clEnqueueMapBuffer +syn keyword clFunction clEnqueueMapImage +syn keyword clFunction clEnqueueUnmapMemObject +syn keyword clFunction clEnqueueMigrateMemObjects +syn keyword clFunction clGetMemObjectInfo +syn keyword clFunction clGetImageInfo +syn keyword clFunction clSetMemObjectDestructorCallback +syn keyword clFunction clCreateSampler +syn keyword clFunction clRetainSampler +syn keyword clFunction clReleaseSampler +syn keyword clFunction clGetSamplerInfo +syn keyword clFunction clCreateProgramWithSource +syn keyword clFunction clCreateProgramWithBinary +syn keyword clFunction clCreateProgramWithBuiltInKernels +syn keyword clFunction clRetainProgram +syn keyword clFunction clReleaseProgram +syn keyword clFunction clUnloadCompiler +syn keyword clFunction clBuildProgram +syn keyword clFunction clCompileProgram +syn keyword clFunction clLinkProgram +syn keyword clFunction clUnloadPlatformCompiler +syn keyword clFunction clGetProgramInfo +syn keyword clFunction clGetProgramBuildInfo +syn keyword clFunction clCreateKernel +syn keyword clFunction clCreateKernelsInProgram +syn keyword clFunction clRetainKernel +syn keyword clFunction clReleaseKernel +syn keyword clFunction clSetKernelArg +syn keyword clFunction clGetKernelInfo +syn keyword clFunction clGetKernelArgInfo +syn keyword clFunction clGetKernelWorkGroupInfo +syn keyword clFunction clEnqueueNDRangeKernel +syn keyword clFunction clEnqueueTask +syn keyword clFunction clEnqueueNativeKernel +syn keyword clFunction clEnqueueMarkerWithWaitList +syn keyword clFunction clEnqueueBarrierWithWaitList +syn keyword clFunction clSetPrintfCallback +syn keyword clFunction clGetExtensionFunctionAddressForPlatform +syn keyword clFunction clGetExtensionFunctionAddress +syn keyword clFunction clWaitForEvents +syn keyword clFunction clGetEventInfo +syn keyword clFunction clGetEventProfileInfo +syn keyword clFunction clCreateUserEvent +syn keyword clFunction clRetainEvent +syn keyword clFunction clReleaseEvent +syn keyword clFunction clSetUserEventStatus +syn keyword clFunction clSetEventCallback +syn keyword clFunction clEnqueueMarker +syn keyword clFunction clEnqueueWaitForEvents +syn keyword clFunction clEnqueueBarrier +syn keyword clFunction clGetEventProfilingInfo +syn keyword clFunction clFlush +syn keyword clFunction clFinish +syn keyword clFunction clCreateFromGLBuffer +syn keyword clFunction clCreateFromGLRenderBuffer +syn keyword clFunction clCreateFromGLTexture2D +syn keyword clFunction clCreateFromGLTexture3D +syn keyword clFunction clEnqueueAcquireGLObjects +syn keyword clFunction clEnqueueReleaseGLObjects +syn keyword clFunction clGetGLObjectInfo +syn keyword clFunction clGetGLTextureInfo + +" Error Codes +syn keyword clErr CL_SUCCESS +syn keyword clErr CL_DEVICE_NOT_FOUND +syn keyword clErr CL_DEVICE_NOT_AVAILABLE +syn keyword clErr CL_COMPILER_NOT_AVAILABLE +syn keyword clErr CL_MEM_OBJECT_ALLOCATION_FAILURE +syn keyword clErr CL_OUT_OF_RESOURCES +syn keyword clErr CL_OUT_OF_HOST_MEMORY +syn keyword clErr CL_PROFILING_INFO_NOT_AVAILABLE +syn keyword clErr CL_MEM_COPY_OVERLAP +syn keyword clErr CL_IMAGE_FORMAT_MISMATCH +syn keyword clErr CL_IMAGE_FORMAT_NOT_SUPPORTED +syn keyword clErr CL_BUILD_PROGRAM_FAILURE +syn keyword clErr CL_MAP_FAILURE +syn keyword clErr CL_MISALIGNED_SUB_BUFFER_OFFSET +syn keyword clErr CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST +syn keyword clErr CL_COMPILE_PROGRAM_FAILURE +syn keyword clErr CL_LINKER_NOT_AVAILABLE +syn keyword clErr CL_LINK_PROGRAM_FAILURE +syn keyword clErr CL_DEVICE_PARTITION_FAILED +syn keyword clErr CL_KERNEL_ARG_INFO_NOT_AVAILABLE +syn keyword clErr CL_INVALID_VALUE +syn keyword clErr CL_INVALID_DEVICE_TYPE +syn keyword clErr CL_INVALID_PLATFORM +syn keyword clErr CL_INVALID_DEVICE +syn keyword clErr CL_INVALID_CONTEXT +syn keyword clErr CL_INVALID_QUEUE_PROPERTIES +syn keyword clErr CL_INVALID_COMMAND_QUEUE +syn keyword clErr CL_INVALID_HOST_PTR +syn keyword clErr CL_INVALID_MEM_OBJECT +syn keyword clErr CL_INVALID_IMAGE_FORMAT_DESCRIPTOR +syn keyword clErr CL_INVALID_IMAGE_SIZE +syn keyword clErr CL_INVALID_SAMPLER +syn keyword clErr CL_INVALID_BINARY +syn keyword clErr CL_INVALID_BUILD_OPTIONS +syn keyword clErr CL_INVALID_PROGRAM +syn keyword clErr CL_INVALID_PROGRAM_EXECUTABLE +syn keyword clErr CL_INVALID_KERNEL_NAME +syn keyword clErr CL_INVALID_KERNEL_DEFINITION +syn keyword clErr CL_INVALID_KERNEL +syn keyword clErr CL_INVALID_ARG_INDEX +syn keyword clErr CL_INVALID_ARG_VALUE +syn keyword clErr CL_INVALID_ARG_SIZE +syn keyword clErr CL_INVALID_KERNEL_ARGS +syn keyword clErr CL_INVALID_WORK_DIMENSION +syn keyword clErr CL_INVALID_WORK_GROUP_SIZE +syn keyword clErr CL_INVALID_WORK_ITEM_SIZE +syn keyword clErr CL_INVALID_GLOBAL_OFFSET +syn keyword clErr CL_INVALID_EVENT_WAIT_LIST +syn keyword clErr CL_INVALID_EVENT +syn keyword clErr CL_INVALID_OPERATION +syn keyword clErr CL_INVALID_GL_OBJECT +syn keyword clErr CL_INVALID_BUFFER_SIZE +syn keyword clErr CL_INVALID_MIP_LEVEL +syn keyword clErr CL_INVALID_GLOBAL_WORK_SIZE +syn keyword clErr CL_INVALID_PROPERTY +syn keyword clErr CL_INVALID_IMAGE_DESCRIPTOR +syn keyword clErr CL_INVALID_COMPILER_OPTIONS +syn keyword clErr CL_INVALID_LINKER_OPTIONS +syn keyword clErr CL_INVALID_DEVICE_PARTITION_COUNT +syn keyword clErr CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR +syn keyword clErr CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR +syn keyword clErr CL_DEVICES_FOR_GL_CONTEXT_KHR +syn keyword clErr CL_GL_CONTEXT_KHR +syn keyword clErr CL_EGL_DISPLAY_KHR +syn keyword clErr CL_GLX_DISPLAY_KHR +syn keyword clErr CL_WGL_HDC_KHR +syn keyword clErr CL_CGL_SHAREGROUP_KHR + +let b:current_syntax = "opencl" + +hi def link clType Type +hi def link clStructure Structure +hi def link clFunction Function +hi def link clMacro PreProc +hi def link clConstant Constant +hi def link clErr Constant diff --git a/doc/vim_syntax_highlighting/syntax/stdcl.vim b/doc/vim_syntax_highlighting/syntax/stdcl.vim new file mode 100644 index 00000000..eb61fe1f --- /dev/null +++ b/doc/vim_syntax_highlighting/syntax/stdcl.vim @@ -0,0 +1,113 @@ +" Vim syntax file +" Language: C STDCL +" Maintainer: James Ross +" Version: 0.2 +" Last Change: 2013-03-25 +" +" Description: Syntax highlighting for STDCL + + +"if exists("b:current_syntax") +" finish +"endif + +if version < 600 + so :p:h/c.vim +else + runtime! syntax/c.vim + unlet b:current_syntax +endif + +" Types +syn match stdclType "cl_\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)" +syn match stdclType "cl_\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)\(2\|3\|4\|8\)" +syn match stdclType "cl_\(char\|uchar\|short\|ushort\|int\|uint\|long\|ulong\|float\|double\)16" +syn keyword stdclType cldev_info +syn keyword stdclType clmulti_array +syn keyword stdclType clndrange_struct +syn keyword stdclType clstat_info +syn keyword stdclType clvector + +" Constants {{{ +syn keyword stdclConstant __STDCL__ +syn keyword stdclConstant stddev +syn keyword stdclConstant stdcpu +syn keyword stdclConstant stdgpu +syn keyword stdclConstant stdrpu +syn keyword stdclConstant stdacc +syn keyword stdclConstant stdnpu +syn keyword stdclConstant STDCL_EVENTLIST_MAX +syn keyword stdclConstant CL_DEVICE_TYPE_RPU +syn keyword stdclConstant CLCONTEXT +syn keyword stdclConstant CONTEXT +syn match stdclConstant "CLLD_\(DEFAULT\|LAZY\|NOW\|NOBUILD\|GLOBAL\)" +syn match stdclConstant "CL_MEM_\(RO\|WO\|RW\|HOST\|DEVICE\|NOCOPY\|DETACHED\|NOFORCE\|IMAGE2D\|CLBUF\|GLBUF\|GLTEX2D\|GLTEX3D\|GLRBUF\)" +syn match stdclConstant "CL_MCTL_\(GET_STATUS\|GET_DEVNUM\|SET_DEVNUM\|MARK_CLEAN\|SET_IMAGE2D\|SET_USRFLAGS\|CLR_USRFLAGS\)" +syn keyword stdclConstant CLMEM_MAGIC +syn keyword stdclConstant CL_FAST +syn keyword stdclConstant CL_EVENT_WAIT +syn keyword stdclConstant CL_EVENT_NOWAIT +syn keyword stdclConstant CL_EVENT_RELEASE +syn keyword stdclConstant CL_EVENT_NORELEASE +syn keyword stdclConstant CL_KERNEL_EVENT +syn keyword stdclConstant CL_MEM_EVENT +syn keyword stdclConstant CL_ALL_EVENT + +" }}} + +" Functions +syn keyword stdclFunction clcontext_create +syn keyword stdclFunction clcontext_create_stdnpu +syn keyword stdclFunction clcontext_destroy +syn keyword stdclFunction clgetndev +syn keyword stdclFunction clstat +syn keyword stdclFunction clgetdevinfo +syn keyword stdclFunction clfreport_devinfo +syn keyword stdclFunction clload +syn keyword stdclFunction clloadb +syn keyword stdclFunction clbuild +syn keyword stdclFunction clopen +syn keyword stdclFunction clsopen +syn keyword stdclFunction clsym +syn keyword stdclFunction clclose +syn keyword stdclFunction clerror +syn keyword stdclFunction clsizeofmem +syn keyword stdclFunction clmalloc +syn keyword stdclFunction clfree +syn keyword stdclFunction clmattach +syn keyword stdclFunction clmdetach +syn keyword stdclFunction clmctl_va +syn keyword stdclFunction clmrealloc +syn keyword stdclFunction clmsync +syn keyword stdclFunction clmcopy +syn keyword stdclFunction clmemptr +syn keyword stdclFunction clglmalloc +syn keyword stdclFunction clglmsync +syn keyword stdclFunction clmctl +syn keyword stdclFunction clarg_set_global +syn keyword stdclFunction clarg_set +syn keyword stdclFunction clarg_set_local +syn keyword stdclFunction clforka +syn keyword stdclFunction clfork +syn keyword stdclFunction clwait +syn keyword stdclFunction clwaitev +syn keyword stdclFunction clflush +syn keyword stdclFunction oclperror +syn keyword stdclFunction clperror +syn keyword stdclFunction oclperror_str +syn keyword stdclFunction clperror_str + +" Macros +syn keyword stdclMacro clndrange_init1d +syn keyword stdclMacro clndrange_init2d +syn keyword stdclMacro clndrange_init3d +syn keyword stdclMacro clndrange_set1d +syn keyword stdclMacro clndrange_set2d +syn keyword stdclMacro clndrange_set3d + +let b:current_syntax = "stdcl" + +hi def link stdclType Type +hi def link stdclFunction Function +hi def link stdclMacro PreProc +hi def link stdclConstant Constant