local funcstring = [[ get_work_dim() Returns the number of dimensions in use get_global_size(uint dimindx) Returns the number of global work-items specified for dimension identified by dimindx get_global_id(uint dimindx) Returns the unique global work-item ID value for dimension identified by dimindx get_local_size(uint dimindx) Returns the number of local work-items specified in dimension identified by dimindx get_local_id(uint dimindx) Returns the unique local work-item ID i.e. a work-item within a specific work-group for dimension identified by dimindx. get_num_groups(uint dimindx) Returns the number of work-groups that will execute a kernel for dimension identified by dimindx get_group_id(uint dimindx) Returns the work-group ID acos(gentype) Arc cosine function acosh(gentype) Inverse hyperbolic cosine acospi(gentype) Compute acos (x) / PI asin(gentype) Arc sine function asinh(gentype) Inverse hyperbolic sine asinpi(gentype x) Compute asin (x) / PI atan(gentype y_over_x) Arc tangent function atan2(gentype y, gentype x) Arc tangent of y / x atanh(gentype) Hyperbolic arc tangent. atanpi(gentype x) Compute atan (x) / PI atan2pi(gentype y, gentype x) Compute atan2 (y, x) / PI cbrt(gentype) Compute cube-root ceil(gentype) Round to integral value using the round to +ve infinity rounding mode copysign(gentype x, gentype y) Returns x with its sign changed to match the sign of y cos(gentype) Compute cosine cosh(gentype) Compute hyperbolic consine cospi(gentype x) Compute cos (PI*x) erfc(gentype) Complementary error function erf(gentype) Error function encountered in integrating the normal distribution exp(gentype x) Compute the base- e exponential of x exp2(gentype) Exponential base 2 function exp10(gentype) Exponential base 10 function expm1(gentype x) Compute e^x - 1.0 fabs(gentype) Compute absolute value of a floating-point number fdim(gentype x, gentype y) x - y if x > y, +0 if x is less than or equal to y floor(gentype) Round to integral value using the round to –ve infinity rounding mode fma(gentype a, gentype b, gentype c) Returns the correctly rounded floating-point representation of the sum of c with the infinitely precise product of a and b fmax(gentype x, gentype y) Returns y if x < y, otherwise it returns x fmin(gentype x, gentype y) Returns y if y < x, otherwise it returns x fmod(gentype x, gentype y) Modulus. Returns x – y * trunc (x/y) fract(gentype x, gentype *iptr) Returns fmin( x – floor (x), 0x1.fffffep-1f ). frexp(gentype x, intn *exp) Extract mantissa and exponent from x hypot(gentype x, gentype y) Compute the value of the square root of x2+y2 ilogb(gentype x) Return the exponent as an integer value ldexp(gentype x, intn n) Multiply x by 2 to the power n lgamma(gentype x) Returns the natural logarithm of the absolute value of the gamma function lgamma_r(gentype x, intn *signp) Returns the natural logarithm of the absolute value of the gamma function log(gentype) Compute natural logarithm log2(gentype) Compute a base 2 logarithm log10(gentype) Compute a base 10 logarithm log1p(gentype x) Compute loge(1.0 + x) logb(gentype x) Compute the exponent of x, which is the integral part of logr|x| mad(gentype a, gentype b, gentype c) Approximates a * b + c. modf(gentype x, gentype *iptr) Decompose a floating-point number nan(uintn nancode) Returns a quiet NaN nextafter(gentype x, gentype y) Computes the next representable single-precision floating-point value following x in the direction of y. pow(gentype x, gentype y) Compute x to the power y pown(gentype x, intn y) Compute x to the power y, where y is an integer powr(gentype x, gentype y) Compute x to the power y, where x is >= 0 remainder(gentype x, gentype y) r = x - n*y, where n is the integer nearest the exact value of x/y remquo(gentype x, gentype y, intn *quo) r = x - n*y, where n is the integer nearest the exact value of x/y rint(gentype) Round to integral value (using round to nearest even rounding mode) rootn(gentype x, intn y) Compute x to the power 1/y round(gentype x) Return the integral value nearest to x rounding halfway cases away from zero rsqrt(gentype) Compute inverse square root sin(gentype) Compute sine sincos(gentype x, gentype *cosval) Compute sine and cosine of x sinh(gentype) Compute hyperbolic sine. sinpi(gentype x) Compute sin (PI*x) sqrt(gentype) Compute square root tan(gentype) Compute tangent tanh(gentype) Compute hyperbolic tangent tanpi(gentype x) Compute tan (PI*x) tgamma(gentype) Compute the gamma function trunc(gentype) Round to integral value using the round to zero abs(gentype x) Returns |x| abs_diff(gentype x, gentype y) Returns |x – y| without modulo overflow add_sat(gentype x, gentype y) Returns x + y and saturates the result hadd(gentype x, gentype y) Returns (x + y) >> 1 rhadd(gentype x, gentype y) Returns (x + y + 1) >> 1 clz(gentype x) Returns the number of leading 0-bits in x, starting at the most significant bit position. mad_hi(gentype a, gentype b, gentype c) Returns mul_hi(a, b) + c mad_sat(gentype a, gentype b, gentype c) Returns a * b + c and saturates the result max(gentype x, gentype y) Returns y if x < y, otherwise it returns x min(gentype x, gentype y) Returns y if y < x, otherwise it returns x mul_hi(gentype x, gentype y) Computes x * y and returns the high half of the product of x and y rotate(gentype v, gentype i) sub_sat(gentype x, gentype y) Returns x - y and saturates the result upsample(charn hi, ucharn lo) result[i] = ((short)hi[i] << 8) | lo[i] mad24(gentype x, gentype y, gentype z) mul24(gentype x, gentype y) clamp(gentype x, gentype minval, gentype maxval) Returns fmin(fmax(x, minval), maxval) degrees(gentype radians) Converts radians to degrees max(gentype x, gentype y) min(gentype x, gentype y) mix(gentype x, gentype y, gentype a) Returns the linear blend of x&y: x + (y – x) * a radians(gentype degrees) Converts degrees to radians step(gentype edge, gentype x) Returns 0.0 if x < edge, otherwise it returns 1.0 smoothstep(genType edge0, genType edge1, genType x) sign(gentype x) cross(float4 p0, float4 p1) Returns the cross product of p0.xyz and p1.xyz. dot(gentype p0, gentype p1) Compute dot product distance(gentype p0, gentype p1) Returns the distance between p0 and p1 length(gentype p) Return the length of vecto normalize(gentype p) Returns a vector in the same direction as p but with length of 1. fast_distance(gentype p0, gentype p1) Returns fast_length(p0 – p1). fast_length(gentype p) Returns the length of vector fast_normalize(gentype p) Returns a vector in the same direction as p but with length of 1. read_imagef(image2d_t image, sampler_t sampler, int2 coord) read_imagei(image2d_t image, sampler_t sampler, int2 coord) read_imageui(image2d_t image, sampler_t sampler, int2 coord) write_imagef(image2d_t image, int2 coord, float4 color) write_imagei(image2d_t image, int2 coord, int4 color) write_imageui(image2d_t image, int2 coord, unsigned int4 color) get_image_width(image2d_t image) get_image_width(image3d_t image) get_image_height(image2d_t image) get_image_height(image3d_t image) get_image_channel_data_type(image2d_t image) get_image_channel_data_type(image3d_t image) get_image_channel_order(image2d_t image) get_image_channel_order(image3d_t image) get_image_dim(image2d_t image) get_image_dim(image3d_t image) barrier(cl_mem_fence_flags flags) All work-items in a work-group executing the kernel must execute this function before any are allowed to continue execution beyond the barrier. mem_fence(cl_mem_fence_flags flags) Orders loads and stores of a work-item executing a kernel. read_mem_fence(cl_mem_fence_flags flags) Read memory barrier that orders only loads. write_mem_fence(cl_mem_fence_flags flags) Write memory barrier that orders only stores. async_work_group_copy(gentype *dst, const gentype *src, size_t num_elements, event_t event) Perform an async copy of num_elements gentype elements from src to dst. wait_group_events(int num_events, event_t *event_list) Wait for events that identify the async_work_group_copy operations to complete. prefetch(const __global gentype *p, size_t num_elements) Prefetch num_elements * sizeof(gentype) bytes into the global cache. vload2(size_t offset, const type *p) Read vector data from memory vload4(size_t offset, const type *p) Read vector data from memory vload8(size_t offset, const type *p) Read vector data from memory vload16(size_t offset, const type *p) Read vector data from memory vstore2(type2 data, size_t offset, type *p) Write vector data to memory vstore4(type4 data, size_t offset, type *p) Write vector data to memory vstore8(type8 data, size_t offset, type *p) Write vector data to memory vstore16(type16 data, size_t offset, type *p) Write vector data to memory ]] local function fn (description) local description2,returns,args = description:match("(.+)%-%s*(%b())%s*(%b())") if not description2 then return {type="function",description=description, returns="(?)"} end return {type="function",description=description2, returns=returns:gsub("^%s+",""):gsub("%s+$",""), args = args} end local function val (description) return {type="value",description = description} end -- docs local api = { } local convtypes = [[bool char uchar short ushort int uint long ulong float double]] local convout = {} for i in convtypes:gmatch("([%w_]+)") do local suffix = {"","_rte","_rtz","_rtp","_rtn"} for k,t in ipairs(suffix) do table.insert(convout,"convert_"..i..t) table.insert(convout,"convert_"..i.."_sat"..t) local vectors = {2,4,8,16} for n,v in ipairs(vectors) do table.insert(convout,"convert_"..i..v..t) table.insert(convout,"convert_"..i..v.."_sat"..t) end end end convout = table.concat(convout, " ") local astypes = [[int uint uchar ushort float double size_t ptrdiff_t intptr_t uintptr_t long ulong char short unsigned float2 float4 float8 float16 double2 double4 double8 double16 char2 char4 char8 char16 uchar2 uchar4 uchar8 uchar16 short2 short4 short8 short16 ushort2 ushort4 ushort8 ushort16 int2 int4 int8 int16 uint2 uint4 uint8 uint16 long2 long4 long8 long16 ulong2 ulong4 ulong8 ulong16]] local astypeout = {} for i in astypes:gmatch("([%w_]+)") do table.insert(astypeout, "as_"..i) end astypeout = table.concat(astypeout, " ") local keyw = astypeout.." "..convout.." "..[[ int uint uchar ushort half float bool double size_t ptrdiff_t intptr_t uintptr_t void long ulong char short unsigned half2 half4 half8 half16 float2 float4 float8 float16 double2 double4 double8 double16 char2 char4 char8 char16 uchar2 uchar4 uchar8 uchar16 short2 short4 short8 short16 ushort2 ushort4 ushort8 ushort16 int2 int4 int8 int16 uint2 uint4 uint8 uint16 long2 long4 long8 long16 ulong2 ulong4 ulong8 ulong16 image2d_t image3d_t sampler_t event_t cl_image_format struct typedef void const return switch case for do while if else break continue volatile CLK_A CLK_R CLK_RG CLK_RGB CLK_RGBA CLK_ARGB CLK_BGRA CLK_INTENSITY CLK_LUMINANCE MAXFLOAT HUGE_VALF INFINITY NAN CLK_LOCAL_MEM_FENCE CLK_GLOBAL_MEM_FENCE CLK_SNORM_INT8 CLK_SNORM_INT16 CLK_UNORM_INT8 CLK_UNORM_INT16 CLK_UNORM_SHORT_565 CLK_UNORM_SHORT_555 CLK_UNORM_SHORT_101010 CLK_SIGNED_INT8 CLK_SIGNED_INT16 CLK_SIGNED_INT32 CLK_UNSIGNED_INT8 CLK_UNSIGNED_INT16 CLK_UNSIGNED_INT32 CLK_HALF_FLOAT CLK_FLOAT __FILE__ __LINE__ __OPENCL_VERSION__ __ENDIAN_LITTLE__ __ROUNDING_MODE__ __IMAGE_SUPPORT__ __FAST_RELAXED_MATH__ __kernel kernel __attribute__ __read_only __write_only read_only write_only __constant constant __local local __global global __private private vec_type_hint work_group_size_hint reqd_work_group_size aligned packed endian host device async_work_group_copy wait_group_events prefetch clamp min max degrees radians sign smoothstep step mix mem_fence read_mem_fence write_mem_fence cross prod distance dot length normalize fast_distance fast_length fast_normalize read_image write_image get_image_width get_image_height get_image_depth get_image_channel_data_type get_image_channel_order get_image_dim abs abs_diff add_sat clz hadd mad24 mad_hi mad_sat mul24 mul_hi rhadd rotate sub_sat upsample read_imagei write_imagei read_imageui write_imageui read_imagef write_imagef isequal isnotequal isgreater isgreaterequal isless islessequal islessgreater isfinite isinf isnan isnormal isordered isunordered signbit any all bitselect select acos acosh acospi asin asinh asinpi atan atan2 atanh atanpi atan2pi cbrt ceil copysign cos half_cos native_cos cosh cospi half_divide native_divide erf erfc exp half_exp native_exp exp2 half_exp2 native_exp2 exp10 half_exp10 native_exp10 expm1 fabs fdim floor fma fmax fmin fmod fract frexp hypot ilogb ldexp lgamma lgamma_r log half_log native_log log2 half_log2 native_log2 log10 half_log10 native_log10 log1p logb mad modf nan nextafter pow pown powr half_powr native_powr half_recip native_recip remainder remquo rint round rootn rsqrt half_rsqrt native_rsqrt sin half_sin native_sin sincos sinh sinpi sqrt half_sqrt native_sqrt tan half_tan native_tan tanh tanpi tgamma trunc barrier vload2 vload4 vload8 vload16 vload_half vload_half2 vload_half4 vload_half8 vload_half16 vloada_half4 vloada_half8 vloada_half16 vstore2 vstore4 vstore8 vstore16 vstore_half vstore_half2 vstore_half4 vstore_half8 vstore_half16 vstorea_half4 vstorea_half8 vstorea_half16 get_global_id get_global_size get_group_id get_local_id get_local_size get_num_groups get_work_dim ]] -- keywords - shouldn't be left out for w in keyw:gmatch("([a-zA-Z_0-9]+)") do api[w] = {type="keyword"} end return api