[OpenCL] Fix lang mode predefined macros for C++ mode

[Anastasia]

Summary

We should only avoid adding __OPENCL_C_VERSION__, all other predefined macros about the language mode are still valid.

This change also fixes the language version check in the headers.

[svenvh]

LGTM!

cfe/trunk/lib/Frontend/InitPreprocessor.cpp

if (LangOpts.CPlusPlus) {
         Builder.defineMacro("__OPENCL_C_VERSION__", "120");
         break;
       case 200:
         Builder.defineMacro("__OPENCL_C_VERSION__", "200");
         break;
       default:
         llvm_unreachable("Unsupported OpenCL version");
       }
-      Builder.defineMacro("CL_VERSION_1_0", "100");
-      Builder.defineMacro("CL_VERSION_1_1", "110");
-      Builder.defineMacro("CL_VERSION_1_2", "120");
-      Builder.defineMacro("CL_VERSION_2_0", "200");
+    }
 
-      if (TI.isLittleEndian())
-        Builder.defineMacro("__ENDIAN_LITTLE__");
 
-      if (LangOpts.FastRelaxedMath)
-        Builder.defineMacro("__FAST_RELAXED_MATH__");
-    }
-  }
   // Not "standard" per se, but available even with the -undef flag.
   if (LangOpts.AsmPreprocessor)
     Builder.defineMacro("__ASSEMBLER__");
   if (LangOpts.CUDA && !LangOpts.HIP)
     Builder.defineMacro("__CUDA__");
   if (LangOpts.HIP) {
     Builder.defineMacro("__HIP__");
     Builder.defineMacro("__HIPCC__");

cfe/trunk/lib/Headers/opencl-c-base.h

 #endif
 typedef double double2 __attribute__((ext_vector_type(2)));
 typedef double double3 __attribute__((ext_vector_type(3)));
 typedef double double4 __attribute__((ext_vector_type(4)));
 typedef double double8 __attribute__((ext_vector_type(8)));
 typedef double double16 __attribute__((ext_vector_type(16)));
 #endif
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #define NULL ((void*)0)
 #endif
 
 /**
  * Value of maximum non-infinite single-precision floating-point
  * number.
  */
 #define MAXFLOAT 0x1.fffffep127f

 #define CLK_LOCAL_MEM_FENCE    0x01
 
 /**
  * Queue a memory fence to ensure correct
  * ordering of memory operations to global memory
  */
 #define CLK_GLOBAL_MEM_FENCE   0x02
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 typedef enum memory_scope {
   memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
   memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
   memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
   memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
 #if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
   memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
 #endif
 } memory_scope;
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
-
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
 /**
  * Queue a memory fence to ensure correct ordering of memory
  * operations between work-items of a work-group to
  * image memory.
  */
 #define CLK_IMAGE_MEM_FENCE  0x04
 
 #ifndef ATOMIC_VAR_INIT
 #define ATOMIC_VAR_INIT(x) (x)
 #endif //ATOMIC_VAR_INIT
 #define ATOMIC_FLAG_INIT 0
 
 // enum values aligned with what clang uses in EmitAtomicExpr()
 typedef enum memory_order
 {
   memory_order_relaxed = __ATOMIC_RELAXED,
   memory_order_acquire = __ATOMIC_ACQUIRE,
   memory_order_release = __ATOMIC_RELEASE,
   memory_order_acq_rel = __ATOMIC_ACQ_REL,
   memory_order_seq_cst = __ATOMIC_SEQ_CST
 } memory_order;
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.3, v1.2 s6.12.14, v2.0 s6.13.14 - Image Read and Write Functions
 
 // These values need to match the runtime equivalent
 //
 // Addressing Mode.
 //
 #define CLK_ADDRESS_NONE                0

 #define CLK_sRGB              0x10BF
 #define CLK_sRGBx             0x10C0
 #define CLK_sRGBA             0x10C1
 #define CLK_sBGRA             0x10C2
 #define CLK_ABGR              0x10C3
 #endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
 
 // OpenCL v2.0 s6.13.16 - Pipe Functions
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #define CLK_NULL_RESERVE_ID (__builtin_astype(((void*)(__SIZE_MAX__)), reserve_id_t))
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
-
 
 // OpenCL v2.0 s6.13.17 - Enqueue Kernels
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
-
 #define CL_COMPLETE                                 0x0
 #define CL_RUNNING                                  0x1
 #define CL_SUBMITTED                                0x2
 #define CL_QUEUED                                   0x3
 
 #define CLK_SUCCESS                                 0
 #define CLK_ENQUEUE_FAILURE                         -101
 #define CLK_INVALID_QUEUE                           -102

 
 typedef struct {
   unsigned int workDimension;
   size_t globalWorkOffset[MAX_WORK_DIM];
   size_t globalWorkSize[MAX_WORK_DIM];
   size_t localWorkSize[MAX_WORK_DIM];
 } ndrange_t;
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 #ifdef cl_intel_device_side_avc_motion_estimation
 #pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : begin
 
 #define CLK_AVC_ME_MAJOR_16x16_INTEL 0x0
 #define CLK_AVC_ME_MAJOR_16x8_INTEL 0x1
 #define CLK_AVC_ME_MAJOR_8x16_INTEL 0x2
 #define CLK_AVC_ME_MAJOR_8x8_INTEL 0x3

cfe/trunk/lib/Headers/opencl-c.h

 //===--- opencl-c.h - OpenCL C language builtin function header -----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OPENCL_H_
 #define _OPENCL_H_
 
 #include "opencl-c-base.h"
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifndef cl_khr_depth_images
 #define cl_khr_depth_images
 #endif //cl_khr_depth_images
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 #if __OPENCL_C_VERSION__ < CL_VERSION_2_0
 #ifdef cl_khr_3d_image_writes
 #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
 #endif //cl_khr_3d_image_writes
 #endif //__OPENCL_C_VERSION__ < CL_VERSION_2_0
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 #pragma OPENCL EXTENSION cl_intel_planar_yuv : begin
 #pragma OPENCL EXTENSION cl_intel_planar_yuv : end
-#endif // __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 
 #define __ovld __attribute__((overloadable))
 #define __conv __attribute__((convergent))
 
 // Optimizations
 #define __purefn __attribute__((pure))
 #define __cnfn __attribute__((const))
 

  * global_work_offset argument to
  * clEnqueueNDRangeKernel.
  * Valid values of dimindx are 0 to get_work_dim() - 1.
  * For other values, get_global_offset() returns 0.
  * For clEnqueueTask, this returns 0.
  */
 size_t __ovld __cnfn get_global_offset(uint dimindx);
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 size_t __ovld get_enqueued_local_size(uint dimindx);
 size_t __ovld get_global_linear_id(void);
 size_t __ovld get_local_linear_id(void);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.2, v1.2 s6.12.2, v2.0 s6.13.2 - Math functions
 
 /**
  * Arc cosine function.
  */
 float __ovld __cnfn acos(float);
 float2 __ovld __cnfn acos(float2);

 half8 __ovld __cnfn fmod(half8 x, half8 y);
 half16 __ovld __cnfn fmod(half16 x, half16 y);
 #endif //cl_khr_fp16
 
 /**
  * Returns fmin(x - floor (x), 0x1.fffffep-1f ).
  * floor(x) is returned in iptr.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld fract(float x, float *iptr);
 float2 __ovld fract(float2 x, float2 *iptr);
 float3 __ovld fract(float3 x, float3 *iptr);
 float4 __ovld fract(float4 x, float4 *iptr);
 float8 __ovld fract(float8 x, float8 *iptr);
 float16 __ovld fract(float16 x, float16 *iptr);
 #ifdef cl_khr_fp64
 double __ovld fract(double x, double *iptr);

 half16 __ovld fract(half16 x, __local half16 *iptr);
 half __ovld fract(half x, __private half *iptr);
 half2 __ovld fract(half2 x, __private half2 *iptr);
 half3 __ovld fract(half3 x, __private half3 *iptr);
 half4 __ovld fract(half4 x, __private half4 *iptr);
 half8 __ovld fract(half8 x, __private half8 *iptr);
 half16 __ovld fract(half16 x, __private half16 *iptr);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Extract mantissa and exponent from x. For each
  * component the mantissa returned is a float with
  * magnitude in the interval [1/2, 1) or 0. Each
  * component of x equals mantissa returned * 2^exp.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld frexp(float x, int *exp);
 float2 __ovld frexp(float2 x, int2 *exp);
 float3 __ovld frexp(float3 x, int3 *exp);
 float4 __ovld frexp(float4 x, int4 *exp);
 float8 __ovld frexp(float8 x, int8 *exp);
 float16 __ovld frexp(float16 x, int16 *exp);
 #ifdef cl_khr_fp64
 double __ovld frexp(double x, int *exp);

 half16 __ovld frexp(half16 x, __local int16 *exp);
 half __ovld frexp(half x, __private int *exp);
 half2 __ovld frexp(half2 x, __private int2 *exp);
 half3 __ovld frexp(half3 x, __private int3 *exp);
 half4 __ovld frexp(half4 x, __private int4 *exp);
 half8 __ovld frexp(half8 x, __private int8 *exp);
 half16 __ovld frexp(half16 x, __private int16 *exp);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Compute the value of the square root of x^2 + y^2
  * without undue overflow or underflow.
  */
 float __ovld __cnfn hypot(float x, float y);
 float2 __ovld __cnfn hypot(float2 x, float2 y);
 float3 __ovld __cnfn hypot(float3 x, float3 y);

 half __ovld __cnfn lgamma(half x);
 half2 __ovld __cnfn lgamma(half2 x);
 half3 __ovld __cnfn lgamma(half3 x);
 half4 __ovld __cnfn lgamma(half4 x);
 half8 __ovld __cnfn lgamma(half8 x);
 half16 __ovld __cnfn lgamma(half16 x);
 #endif //cl_khr_fp16
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld lgamma_r(float x, int *signp);
 float2 __ovld lgamma_r(float2 x, int2 *signp);
 float3 __ovld lgamma_r(float3 x, int3 *signp);
 float4 __ovld lgamma_r(float4 x, int4 *signp);
 float8 __ovld lgamma_r(float8 x, int8 *signp);
 float16 __ovld lgamma_r(float16 x, int16 *signp);
 #ifdef cl_khr_fp64
 double __ovld lgamma_r(double x, int *signp);

 half16 __ovld lgamma_r(half16 x, __local int16 *signp);
 half __ovld lgamma_r(half x, __private int *signp);
 half2 __ovld lgamma_r(half2 x, __private int2 *signp);
 half3 __ovld lgamma_r(half3 x, __private int3 *signp);
 half4 __ovld lgamma_r(half4 x, __private int4 *signp);
 half8 __ovld lgamma_r(half8 x, __private int8 *signp);
 half16 __ovld lgamma_r(half16 x, __private int16 *signp);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Compute natural logarithm.
  */
 float __ovld __cnfn log(float);
 float2 __ovld __cnfn log(float2);
 float3 __ovld __cnfn log(float3);
 float4 __ovld __cnfn log(float4);

 
 /**
  * Decompose a floating-point number. The modf
  * function breaks the argument x into integral and
  * fractional parts, each of which has the same sign as
  * the argument. It stores the integral part in the object
  * pointed to by iptr.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld modf(float x, float *iptr);
 float2 __ovld modf(float2 x, float2 *iptr);
 float3 __ovld modf(float3 x, float3 *iptr);
 float4 __ovld modf(float4 x, float4 *iptr);
 float8 __ovld modf(float8 x, float8 *iptr);
 float16 __ovld modf(float16 x, float16 *iptr);
 #ifdef cl_khr_fp64
 double __ovld modf(double x, double *iptr);

 half16 __ovld modf(half16 x, __local half16 *iptr);
 half __ovld modf(half x, __private half *iptr);
 half2 __ovld modf(half2 x, __private half2 *iptr);
 half3 __ovld modf(half3 x, __private half3 *iptr);
 half4 __ovld modf(half4 x, __private half4 *iptr);
 half8 __ovld modf(half8 x, __private half8 *iptr);
 half16 __ovld modf(half16 x, __private half16 *iptr);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Returns a quiet NaN. The nancode may be placed
  * in the significand of the resulting NaN.
  */
 float __ovld __cnfn nan(uint nancode);
 float2 __ovld __cnfn nan(uint2 nancode);
 float3 __ovld __cnfn nan(uint3 nancode);

  * to x/y, n shall be the even one. If r is zero, it is
  * given the same sign as x. This is the same value
  * that is returned by the remainder function.
  * remquo also calculates the lower seven bits of the
  * integral quotient x/y, and gives that value the same
  * sign as x/y. It stores this signed value in the object
  * pointed to by quo.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld remquo(float x, float y, int *quo);
 float2 __ovld remquo(float2 x, float2 y, int2 *quo);
 float3 __ovld remquo(float3 x, float3 y, int3 *quo);
 float4 __ovld remquo(float4 x, float4 y, int4 *quo);
 float8 __ovld remquo(float8 x, float8 y, int8 *quo);
 float16 __ovld remquo(float16 x, float16 y, int16 *quo);
 #ifdef cl_khr_fp64
 double __ovld remquo(double x, double y, int *quo);

 half16 __ovld remquo(half16 x, half16 y, __local int16 *quo);
 half __ovld remquo(half x, half y, __private int *quo);
 half2 __ovld remquo(half2 x, half2 y, __private int2 *quo);
 half3 __ovld remquo(half3 x, half3 y, __private int3 *quo);
 half4 __ovld remquo(half4 x, half4 y, __private int4 *quo);
 half8 __ovld remquo(half8 x, half8 y, __private int8 *quo);
 half16 __ovld remquo(half16 x, half16 y, __private int16 *quo);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 /**
  * Round to integral value (using round to nearest
  * even rounding mode) in floating-point format.
  * Refer to section 7.1 for description of rounding
  * modes.
  */
 float __ovld __cnfn rint(float);
 float2 __ovld __cnfn rint(float2);

 half16 __ovld __cnfn sin(half16);
 #endif //cl_khr_fp16
 
 /**
  * Compute sine and cosine of x. The computed sine
  * is the return value and computed cosine is returned
  * in cosval.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld sincos(float x, float *cosval);
 float2 __ovld sincos(float2 x, float2 *cosval);
 float3 __ovld sincos(float3 x, float3 *cosval);
 float4 __ovld sincos(float4 x, float4 *cosval);
 float8 __ovld sincos(float8 x, float8 *cosval);
 float16 __ovld sincos(float16 x, float16 *cosval);
 #ifdef cl_khr_fp64
 double __ovld sincos(double x, double *cosval);

 half16 __ovld sincos(half16 x, __local half16 *cosval);
 half __ovld sincos(half x, __private half *cosval);
 half2 __ovld sincos(half2 x, __private half2 *cosval);
 half3 __ovld sincos(half3 x, __private half3 *cosval);
 half4 __ovld sincos(half4 x, __private half4 *cosval);
 half8 __ovld sincos(half8 x, __private half8 *cosval);
 half16 __ovld sincos(half16 x, __private half16 *cosval);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Compute hyperbolic sine.
  */
 float __ovld __cnfn sinh(float);
 float2 __ovld __cnfn sinh(float2);
 float3 __ovld __cnfn sinh(float3);
 float4 __ovld __cnfn sinh(float4);

 long16 __ovld __cnfn clz(long16 x);
 ulong16 __ovld __cnfn clz(ulong16 x);
 
 /**
  * Returns the count of trailing 0-bits in x. If x is 0,
  * returns the size in bits of the type of x or
  * component type of x, if x is a vector.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 char __ovld ctz(char x);
 uchar __ovld ctz(uchar x);
 char2 __ovld ctz(char2 x);
 uchar2 __ovld ctz(uchar2 x);
 char3 __ovld ctz(char3 x);
 uchar3 __ovld ctz(uchar3 x);
 char4 __ovld ctz(char4 x);
 uchar4 __ovld ctz(uchar4 x);

 long3 __ovld ctz(long3 x);
 ulong3 __ovld ctz(ulong3 x);
 long4 __ovld ctz(long4 x);
 ulong4 __ovld ctz(ulong4 x);
 long8 __ovld ctz(long8 x);
 ulong8 __ovld ctz(ulong8 x);
 long16 __ovld ctz(long16 x);
 ulong16 __ovld ctz(ulong16 x);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Returns mul_hi(a, b) + c.
  */
 char __ovld __cnfn mad_hi(char a, char b, char c);
 uchar __ovld __cnfn mad_hi(uchar a, uchar b, uchar c);
 char2 __ovld __cnfn mad_hi(char2 a, char2 b, char2 c);
 uchar2 __ovld __cnfn mad_hi(uchar2 a, uchar2 b, uchar2 c);

 half __ovld vload(size_t offset, const __constant half *p);
 half2 __ovld vload2(size_t offset, const __constant half *p);
 half3 __ovld vload3(size_t offset, const __constant half *p);
 half4 __ovld vload4(size_t offset, const __constant half *p);
 half8 __ovld vload8(size_t offset, const __constant half *p);
 half16 __ovld vload16(size_t offset, const __constant half *p);
 #endif //cl_khr_fp16
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 char2 __ovld vload2(size_t offset, const char *p);
 uchar2 __ovld vload2(size_t offset, const uchar *p);
 short2 __ovld vload2(size_t offset, const short *p);
 ushort2 __ovld vload2(size_t offset, const ushort *p);
 int2 __ovld vload2(size_t offset, const int *p);
 uint2 __ovld vload2(size_t offset, const uint *p);
 long2 __ovld vload2(size_t offset, const long *p);
 ulong2 __ovld vload2(size_t offset, const ulong *p);

 half16 __ovld vload16(size_t offset, const __local half *p);
 half __ovld vload(size_t offset, const __private half *p);
 half2 __ovld vload2(size_t offset, const __private half *p);
 half3 __ovld vload3(size_t offset, const __private half *p);
 half4 __ovld vload4(size_t offset, const __private half *p);
 half8 __ovld vload8(size_t offset, const __private half *p);
 half16 __ovld vload16(size_t offset, const __private half *p);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld vstore2(char2 data, size_t offset, char *p);
 void __ovld vstore2(uchar2 data, size_t offset, uchar *p);
 void __ovld vstore2(short2 data, size_t offset, short *p);
 void __ovld vstore2(ushort2 data, size_t offset, ushort *p);
 void __ovld vstore2(int2 data, size_t offset, int *p);
 void __ovld vstore2(uint2 data, size_t offset, uint *p);
 void __ovld vstore2(long2 data, size_t offset, long *p);
 void __ovld vstore2(ulong2 data, size_t offset, ulong *p);

 void __ovld vstore16(half16 data, size_t offset, __local half *p);
 void __ovld vstore(half data, size_t offset, __private half *p);
 void __ovld vstore2(half2 data, size_t offset, __private half *p);
 void __ovld vstore3(half3 data, size_t offset, __private half *p);
 void __ovld vstore4(half4 data, size_t offset, __private half *p);
 void __ovld vstore8(half8 data, size_t offset, __private half *p);
 void __ovld vstore16(half16 data, size_t offset, __private half *p);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Read sizeof (half) bytes of data from address
  * (p + offset). The data read is interpreted as a
  * half value. The half value is converted to a
  * float value and the float value is returned.
  * The read address computed as (p + offset)
  * must be 16-bit aligned.
  */
 float __ovld vload_half(size_t offset, const __constant half *p);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld vload_half(size_t offset, const half *p);
 #else
 float __ovld vload_half(size_t offset, const __global half *p);
 float __ovld vload_half(size_t offset, const __local half *p);
 float __ovld vload_half(size_t offset, const __private half *p);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Read sizeof (halfn) bytes of data from address
  * (p + (offset * n)). The data read is interpreted
  * as a halfn value. The halfn value read is
  * converted to a floatn value and the floatn
  * value is returned. The read address computed
  * as (p + (offset * n)) must be 16-bit aligned.
  */
 float2 __ovld vload_half2(size_t offset, const __constant half *p);
 float3 __ovld vload_half3(size_t offset, const __constant half *p);
 float4 __ovld vload_half4(size_t offset, const __constant half *p);
 float8 __ovld vload_half8(size_t offset, const __constant half *p);
 float16 __ovld vload_half16(size_t offset, const __constant half *p);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float2 __ovld vload_half2(size_t offset, const half *p);
 float3 __ovld vload_half3(size_t offset, const half *p);
 float4 __ovld vload_half4(size_t offset, const half *p);
 float8 __ovld vload_half8(size_t offset, const half *p);
 float16 __ovld vload_half16(size_t offset, const half *p);
 #else
 float2 __ovld vload_half2(size_t offset, const __global half *p);
 float3 __ovld vload_half3(size_t offset, const __global half *p);
 float4 __ovld vload_half4(size_t offset, const __global half *p);
 float8 __ovld vload_half8(size_t offset, const __global half *p);
 float16 __ovld vload_half16(size_t offset, const __global half *p);
 float2 __ovld vload_half2(size_t offset, const __local half *p);
 float3 __ovld vload_half3(size_t offset, const __local half *p);
 float4 __ovld vload_half4(size_t offset, const __local half *p);
 float8 __ovld vload_half8(size_t offset, const __local half *p);
 float16 __ovld vload_half16(size_t offset, const __local half *p);
 float2 __ovld vload_half2(size_t offset, const __private half *p);
 float3 __ovld vload_half3(size_t offset, const __private half *p);
 float4 __ovld vload_half4(size_t offset, const __private half *p);
 float8 __ovld vload_half8(size_t offset, const __private half *p);
 float16 __ovld vload_half16(size_t offset, const __private half *p);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * The float value given by data is first
  * converted to a half value using the appropriate
  * rounding mode. The half value is then written
  * to address computed as (p + offset). The
  * address computed as (p + offset) must be 16-
  * bit aligned.
  * vstore_half use the current rounding mode.
  * The default current rounding mode is round to
  * nearest even.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld vstore_half(float data, size_t offset, half *p);
 void __ovld vstore_half_rte(float data, size_t offset, half *p);
 void __ovld vstore_half_rtz(float data, size_t offset, half *p);
 void __ovld vstore_half_rtp(float data, size_t offset, half *p);
 void __ovld vstore_half_rtn(float data, size_t offset, half *p);
 #ifdef cl_khr_fp64
 void __ovld vstore_half(double data, size_t offset, half *p);
 void __ovld vstore_half_rte(double data, size_t offset, half *p);

 void __ovld vstore_half_rtp(double data, size_t offset, __local half *p);
 void __ovld vstore_half_rtn(double data, size_t offset, __local half *p);
 void __ovld vstore_half(double data, size_t offset, __private half *p);
 void __ovld vstore_half_rte(double data, size_t offset, __private half *p);
 void __ovld vstore_half_rtz(double data, size_t offset, __private half *p);
 void __ovld vstore_half_rtp(double data, size_t offset, __private half *p);
 void __ovld vstore_half_rtn(double data, size_t offset, __private half *p);
 #endif //cl_khr_fp64
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * The floatn value given by data is converted to
  * a halfn value using the appropriate rounding
  * mode. The halfn value is then written to
  * address computed as (p + (offset * n)). The
  * address computed as (p + (offset * n)) must be
  * 16-bit aligned.
  * vstore_halfn uses the current rounding mode.
  * The default current rounding mode is round to
  * nearest even.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld vstore_half2(float2 data, size_t offset, half *p);
 void __ovld vstore_half3(float3 data, size_t offset, half *p);
 void __ovld vstore_half4(float4 data, size_t offset, half *p);
 void __ovld vstore_half8(float8 data, size_t offset, half *p);
 void __ovld vstore_half16(float16 data, size_t offset, half *p);
 void __ovld vstore_half2_rte(float2 data, size_t offset, half *p);
 void __ovld vstore_half3_rte(float3 data, size_t offset, half *p);
 void __ovld vstore_half4_rte(float4 data, size_t offset, half *p);

 void __ovld vstore_half8_rtp(double8 data, size_t offset, __private half *p);
 void __ovld vstore_half16_rtp(double16 data, size_t offset, __private half *p);
 void __ovld vstore_half2_rtn(double2 data, size_t offset, __private half *p);
 void __ovld vstore_half3_rtn(double3 data, size_t offset, __private half *p);
 void __ovld vstore_half4_rtn(double4 data, size_t offset, __private half *p);
 void __ovld vstore_half8_rtn(double8 data, size_t offset, __private half *p);
 void __ovld vstore_half16_rtn(double16 data, size_t offset, __private half *p);
 #endif //cl_khr_fp64
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * For n = 1, 2, 4, 8 and 16 read sizeof (halfn)
  * bytes of data from address (p + (offset * n)).
  * The data read is interpreted as a halfn value.
  * The halfn value read is converted to a floatn
  * value and the floatn value is returned.
  * The address computed as (p + (offset * n))
  * must be aligned to sizeof (halfn) bytes.
  * For n = 3, vloada_half3 reads a half3 from
  * address (p + (offset * 4)) and returns a float3.
  * The address computed as (p + (offset * 4))
  * must be aligned to sizeof (half) * 4 bytes.
  */
 float __ovld vloada_half(size_t offset, const __constant half *p);
 float2 __ovld vloada_half2(size_t offset, const __constant half *p);
 float3 __ovld vloada_half3(size_t offset, const __constant half *p);
 float4 __ovld vloada_half4(size_t offset, const __constant half *p);
 float8 __ovld vloada_half8(size_t offset, const __constant half *p);
 float16 __ovld vloada_half16(size_t offset, const __constant half *p);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float __ovld vloada_half(size_t offset, const half *p);
 float2 __ovld vloada_half2(size_t offset, const half *p);
 float3 __ovld vloada_half3(size_t offset, const half *p);
 float4 __ovld vloada_half4(size_t offset, const half *p);
 float8 __ovld vloada_half8(size_t offset, const half *p);
 float16 __ovld vloada_half16(size_t offset, const half *p);
 #else
 float __ovld vloada_half(size_t offset, const __global half *p);

 float8 __ovld vloada_half8(size_t offset, const __local half *p);
 float16 __ovld vloada_half16(size_t offset, const __local half *p);
 float __ovld vloada_half(size_t offset, const __private half *p);
 float2 __ovld vloada_half2(size_t offset, const __private half *p);
 float3 __ovld vloada_half3(size_t offset, const __private half *p);
 float4 __ovld vloada_half4(size_t offset, const __private half *p);
 float8 __ovld vloada_half8(size_t offset, const __private half *p);
 float16 __ovld vloada_half16(size_t offset, const __private half *p);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * The floatn value given by data is converted to
  * a halfn value using the appropriate rounding
  * mode.
  * For n = 1, 2, 4, 8 and 16, the halfn value is
  * written to the address computed as (p + (offset
  * * n)). The address computed as (p + (offset *
  * n)) must be aligned to sizeof (halfn) bytes.
  * For n = 3, the half3 value is written to the
  * address computed as (p + (offset * 4)). The
  * address computed as (p + (offset * 4)) must be
  * aligned to sizeof (half) * 4 bytes.
  * vstorea_halfn uses the current rounding
  * mode. The default current rounding mode is
  * round to nearest even.
  */
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld vstorea_half(float data, size_t offset, half *p);
 void __ovld vstorea_half2(float2 data, size_t offset, half *p);
 void __ovld vstorea_half3(float3 data, size_t offset, half *p);
 void __ovld vstorea_half4(float4 data, size_t offset, half *p);
 void __ovld vstorea_half8(float8 data, size_t offset, half *p);
 void __ovld vstorea_half16(float16 data, size_t offset, half *p);
 
 void __ovld vstorea_half_rte(float data, size_t offset, half *p);

 
 void __ovld vstorea_half_rtn(double data, size_t offset, __private half *p);
 void __ovld vstorea_half2_rtn(double2 data,size_t offset, __private half *p);
 void __ovld vstorea_half3_rtn(double3 data,size_t offset, __private half *p);
 void __ovld vstorea_half4_rtn(double4 data,size_t offset, __private half *p);
 void __ovld vstorea_half8_rtn(double8 data,size_t offset, __private half *p);
 void __ovld vstorea_half16_rtn(double16 data,size_t offset, __private half *p);
 #endif //cl_khr_fp64
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.8, v1.2 s6.12.8, v2.0 s6.13.8 - Synchronization Functions
 
 /**
  * All work-items in a work-group executing the kernel
  * on a processor must execute this function before any
  * are allowed to continue execution beyond the barrier.
  * This function must be encountered by all work-items in

  * will queue a memory fence to ensure correct ordering
  * of memory operations to global memory. This can be
  * useful when work-items, for example, write to buffer or
  * image objects and then want to read the updated data.
  */
 
 void __ovld __conv barrier(cl_mem_fence_flags flags);
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld __conv work_group_barrier(cl_mem_fence_flags flags, memory_scope scope);
 void __ovld __conv work_group_barrier(cl_mem_fence_flags flags);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.9, v1.2 s6.12.9 - Explicit Memory Fence Functions
 
 /**
  * Orders loads and stores of a work-item
  * executing a kernel. This means that loads
  * and stores preceding the mem_fence will
  * be committed to memory before any loads

  * values:
  * CLK_LOCAL_MEM_FENCE
  * CLK_GLOBAL_MEM_FENCE.
  */
 void __ovld write_mem_fence(cl_mem_fence_flags flags);
 
 // OpenCL v2.0 s6.13.9 - Address Space Qualifier Functions
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 cl_mem_fence_flags __ovld get_fence(const void *ptr);
 cl_mem_fence_flags __ovld get_fence(void *ptr);
 
 /**
  * Builtin functions to_global, to_local, and to_private need to be declared as Clang builtin functions
  * and checked in Sema since they should be declared as
  *   addr gentype* to_addr (gentype*);
  * where gentype is builtin type or user defined type.
  */
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.10, v1.2 s6.12.10, v2.0 s6.13.10 - Async Copies from Global to Local Memory, Local to Global Memory, and Prefetch
 
 /**
  * event_t async_work_group_copy (
  * __global gentype *dst,
  * const __local gentype *src,
  * size_t num_elements,

 
 #if defined(cl_khr_int64_base_atomics) && defined(cl_khr_int64_extended_atomics)
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : disable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : disable
 #endif
 
 // OpenCL v2.0 s6.13.11 - Atomics Functions
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // double atomics support requires extensions cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics
 #if defined(cl_khr_int64_base_atomics) && defined(cl_khr_int64_extended_atomics)
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
 #endif
 
 // atomic_init()

 
 bool __ovld atomic_flag_test_and_set(volatile atomic_flag *object);
 bool __ovld atomic_flag_test_and_set_explicit(volatile atomic_flag *object, memory_order order);
 bool __ovld atomic_flag_test_and_set_explicit(volatile atomic_flag *object, memory_order order, memory_scope scope);
 void __ovld atomic_flag_clear(volatile atomic_flag *object);
 void __ovld atomic_flag_clear_explicit(volatile atomic_flag *object, memory_order order);
 void __ovld atomic_flag_clear_explicit(volatile atomic_flag *object, memory_order order, memory_scope scope);
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v1.1 s6.11.12, v1.2 s6.12.12, v2.0 s6.13.12 - Miscellaneous Vector Functions
 
 /**
  * The shuffle and shuffle2 built-in functions construct
  * a permutation of elements from one or two input
  * vectors respectively that are of the same type,
  * returning a vector with the same element type as the

 half8 __ovld __cnfn shuffle2(half16 x, half16 y, ushort8 mask);
 
 half16 __ovld __cnfn shuffle2(half2 x, half2 y, ushort16 mask);
 half16 __ovld __cnfn shuffle2(half4 x, half4 y, ushort16 mask);
 half16 __ovld __cnfn shuffle2(half8 x, half8 y, ushort16 mask);
 half16 __ovld __cnfn shuffle2(half16 x, half16 y, ushort16 mask);
 #endif //cl_khr_fp16
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 // OpenCL v1.2 s6.12.13, v2.0 s6.13.13 - printf
 
 int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
 #endif
 
 // OpenCL v1.1 s6.11.3, v1.2 s6.12.14, v2.0 s6.13.14 - Image Read and Write Functions
 
 #ifdef cl_khr_gl_msaa_sharing

 float4 __purefn __ovld read_imagef(read_only image3d_t image, sampler_t sampler, int4 coord);
 float4 __purefn __ovld read_imagef(read_only image3d_t image, sampler_t sampler, float4 coord);
 
 int4 __purefn __ovld read_imagei(read_only image3d_t image, sampler_t sampler, int4 coord);
 int4 __purefn __ovld read_imagei(read_only image3d_t image, sampler_t sampler, float4 coord);
 uint4 __purefn __ovld read_imageui(read_only image3d_t image, sampler_t sampler, int4 coord);
 uint4 __purefn __ovld read_imageui(read_only image3d_t image, sampler_t sampler, float4 coord);
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 float4 __purefn __ovld read_imagef(read_only image2d_array_t image_array, sampler_t sampler, int4 coord);
 float4 __purefn __ovld read_imagef(read_only image2d_array_t image_array, sampler_t sampler, float4 coord);
 
 int4 __purefn __ovld read_imagei(read_only image2d_array_t image_array, sampler_t sampler, int4 coord);
 int4 __purefn __ovld read_imagei(read_only image2d_array_t image_array, sampler_t sampler, float4 coord);
 uint4 __purefn __ovld read_imageui(read_only image2d_array_t image_array, sampler_t sampler, int4 coord);
 uint4 __purefn __ovld read_imageui(read_only image2d_array_t image_array, sampler_t sampler, float4 coord);
-#endif // __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 
 float4 __purefn __ovld read_imagef(read_only image1d_t image, sampler_t sampler, int coord);
 float4 __purefn __ovld read_imagef(read_only image1d_t image, sampler_t sampler, float coord);
 
 int4 __purefn __ovld read_imagei(read_only image1d_t image, sampler_t sampler, int coord);
 int4 __purefn __ovld read_imagei(read_only image1d_t image, sampler_t sampler, float coord);
 uint4 __purefn __ovld read_imageui(read_only image1d_t image, sampler_t sampler, int coord);
 uint4 __purefn __ovld read_imageui(read_only image1d_t image, sampler_t sampler, float coord);
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 float4 __purefn __ovld read_imagef(read_only image1d_array_t image_array, sampler_t sampler, int2 coord);
 float4 __purefn __ovld read_imagef(read_only image1d_array_t image_array, sampler_t sampler, float2 coord);
 
 int4 __purefn __ovld read_imagei(read_only image1d_array_t image_array, sampler_t sampler, int2 coord);
 int4 __purefn __ovld read_imagei(read_only image1d_array_t image_array, sampler_t sampler, float2 coord);
 uint4 __purefn __ovld read_imageui(read_only image1d_array_t image_array, sampler_t sampler, int2 coord);
 uint4 __purefn __ovld read_imageui(read_only image1d_array_t image_array, sampler_t sampler, float2 coord);
-#endif // __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 
 #ifdef cl_khr_depth_images
 float __purefn __ovld read_imagef(read_only image2d_depth_t image, sampler_t sampler, float2 coord);
 float __purefn __ovld read_imagef(read_only image2d_depth_t image, sampler_t sampler, int2 coord);
 
 float __purefn __ovld read_imagef(read_only image2d_array_depth_t image, sampler_t sampler, float4 coord);
 float __purefn __ovld read_imagef(read_only image2d_array_depth_t image, sampler_t sampler, int4 coord);
 #endif //cl_khr_depth_images
 
 #if defined(cl_khr_gl_msaa_sharing)
 float4 __purefn __ovld read_imagef(read_only image2d_msaa_t image, int2 coord, int sample);
 int4 __purefn __ovld read_imagei(read_only image2d_msaa_t image, int2 coord, int sample);
 uint4 __purefn __ovld read_imageui(read_only image2d_msaa_t image, int2 coord, int sample);
 
 float __purefn __ovld read_imagef(read_only image2d_msaa_depth_t image, int2 coord, int sample);
 
 float4 __purefn __ovld read_imagef(read_only image2d_array_msaa_t image, int4 coord, int sample);
 int4 __purefn __ovld read_imagei(read_only image2d_array_msaa_t image, int4 coord, int sample);
 uint4 __purefn __ovld read_imageui(read_only image2d_array_msaa_t image, int4 coord, int sample);
 
 float __purefn __ovld read_imagef(read_only image2d_array_msaa_depth_t image, int4 coord, int sample);
 #endif //cl_khr_gl_msaa_sharing
 
 // OpenCL Extension v2.0 s9.18 - Mipmaps
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifdef cl_khr_mipmap_image
 
 float4 __purefn __ovld read_imagef(read_only image1d_t image, sampler_t sampler, float coord, float lod);
 int4 __purefn __ovld read_imagei(read_only image1d_t image, sampler_t sampler, float coord, float lod);
 uint4 __purefn __ovld read_imageui(read_only image1d_t image, sampler_t sampler, float coord, float lod);
 
 float4 __purefn __ovld read_imagef(read_only image1d_array_t image_array, sampler_t sampler, float2 coord, float lod);
 int4 __purefn __ovld read_imagei(read_only image1d_array_t image_array, sampler_t sampler, float2 coord, float lod);

 
 float __purefn __ovld read_imagef(read_only image2d_array_depth_t image, sampler_t sampler, float4 coord, float2 gradientX, float2 gradientY);
 
 float4 __purefn __ovld read_imagef(read_only image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 int4 __purefn __ovld read_imagei(read_only image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 uint4 __purefn __ovld read_imageui(read_only image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 
 #endif //cl_khr_mipmap_image
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 
 /**
 * Sampler-less Image Access
 */
 
 float4 __purefn __ovld read_imagef(read_only image1d_t image, int coord);
 int4 __purefn __ovld read_imagei(read_only image1d_t image, int coord);
 uint4 __purefn __ovld read_imageui(read_only image1d_t image, int coord);

 float __purefn __ovld read_imagef(read_only image2d_depth_t image, int2 coord);
 float __purefn __ovld read_imagef(read_only image2d_array_depth_t image, int4 coord);
 #endif //cl_khr_depth_images
 
 float4 __purefn __ovld read_imagef(read_only image3d_t image, int4 coord);
 int4 __purefn __ovld read_imagei(read_only image3d_t image, int4 coord);
 uint4 __purefn __ovld read_imageui(read_only image3d_t image, int4 coord);
 
-#endif // __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 
 // Image read functions returning half4 type
 #ifdef cl_khr_fp16
 half4 __purefn __ovld read_imageh(read_only image1d_t image, sampler_t sampler, int coord);
 half4 __purefn __ovld read_imageh(read_only image1d_t image, sampler_t sampler, float coord);
 half4 __purefn __ovld read_imageh(read_only image2d_t image, sampler_t sampler, int2 coord);
 half4 __purefn __ovld read_imageh(read_only image2d_t image, sampler_t sampler, float2 coord);
 half4 __purefn __ovld read_imageh(read_only image3d_t image, sampler_t sampler, int4 coord);
 half4 __purefn __ovld read_imageh(read_only image3d_t image, sampler_t sampler, float4 coord);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 half4 __purefn __ovld read_imageh(read_only image1d_array_t image, sampler_t sampler, int2 coord);
 half4 __purefn __ovld read_imageh(read_only image1d_array_t image, sampler_t sampler, float2 coord);
 half4 __purefn __ovld read_imageh(read_only image2d_array_t image, sampler_t sampler, int4 coord);
 half4 __purefn __ovld read_imageh(read_only image2d_array_t image, sampler_t sampler, float4 coord);
 /**
  * Sampler-less Image Access
  */
 half4 __purefn __ovld read_imageh(read_only image1d_t image, int coord);
 half4 __purefn __ovld read_imageh(read_only image2d_t image, int2 coord);
 half4 __purefn __ovld read_imageh(read_only image3d_t image, int4 coord);
 half4 __purefn __ovld read_imageh(read_only image1d_array_t image, int2 coord);
 half4 __purefn __ovld read_imageh(read_only image2d_array_t image, int4 coord);
 half4 __purefn __ovld read_imageh(read_only image1d_buffer_t image, int coord);
-#endif // __OPENCL_C_VERSION__ >= CL_VERSION_1_2
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)
 #endif //cl_khr_fp16
 
 // Image read functions for read_write images
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 float4 __purefn __ovld read_imagef(read_write image1d_t image, int coord);
 int4 __purefn __ovld read_imagei(read_write image1d_t image, int coord);
 uint4 __purefn __ovld read_imageui(read_write image1d_t image, int coord);
 
 float4 __purefn __ovld read_imagef(read_write image1d_buffer_t image, int coord);
 int4 __purefn __ovld read_imagei(read_write image1d_buffer_t image, int coord);
 uint4 __purefn __ovld read_imageui(read_write image1d_buffer_t image, int coord);
 

 float4 __purefn __ovld read_imagef(read_write image2d_array_msaa_t image, int4 coord, int sample);
 int4 __purefn __ovld read_imagei(read_write image2d_array_msaa_t image, int4 coord, int sample);
 uint4 __purefn __ovld read_imageui(read_write image2d_array_msaa_t image, int4 coord, int sample);
 
 float __purefn __ovld read_imagef(read_write image2d_msaa_depth_t image, int2 coord, int sample);
 float __purefn __ovld read_imagef(read_write image2d_array_msaa_depth_t image, int4 coord, int sample);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifdef cl_khr_mipmap_image
 float4 __purefn __ovld read_imagef(read_write image1d_t image, sampler_t sampler, float coord, float lod);
 int4 __purefn __ovld read_imagei(read_write image1d_t image, sampler_t sampler, float coord, float lod);
 uint4 __purefn __ovld read_imageui(read_write image1d_t image, sampler_t sampler, float coord, float lod);
 
 float4 __purefn __ovld read_imagef(read_write image1d_array_t image_array, sampler_t sampler, float2 coord, float lod);
 int4 __purefn __ovld read_imagei(read_write image1d_array_t image_array, sampler_t sampler, float2 coord, float lod);
 uint4 __purefn __ovld read_imageui(read_write image1d_array_t image_array, sampler_t sampler, float2 coord, float lod);

 
 float __purefn __ovld read_imagef(read_write image2d_array_depth_t image, sampler_t sampler, float4 coord, float2 gradientX, float2 gradientY);
 
 float4 __purefn __ovld read_imagef(read_write image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 int4 __purefn __ovld read_imagei(read_write image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 uint4 __purefn __ovld read_imageui(read_write image3d_t image, sampler_t sampler, float4 coord, float4 gradientX, float4 gradientY);
 
 #endif //cl_khr_mipmap_image
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // Image read functions returning half4 type
 #ifdef cl_khr_fp16
 half4 __purefn __ovld read_imageh(read_write image1d_t image, int coord);
 half4 __purefn __ovld read_imageh(read_write image2d_t image, int2 coord);
 half4 __purefn __ovld read_imageh(read_write image3d_t image, int4 coord);
 half4 __purefn __ovld read_imageh(read_write image1d_array_t image, int2 coord);
 half4 __purefn __ovld read_imageh(read_write image2d_array_t image, int4 coord);
 half4 __purefn __ovld read_imageh(read_write image1d_buffer_t image, int coord);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Write color value to location specified by coordinate
  * (coord.x, coord.y) in the 2D image object specified by image.
  * (coord.x, coord.y) are considered to be unnormalized coordinates
  * and must be in the range 0 ... image width - 1, and 0
  * ... image height - 1.
 

 #endif
 
 #ifdef cl_khr_depth_images
 void __ovld write_imagef(write_only image2d_depth_t image, int2 coord, float color);
 void __ovld write_imagef(write_only image2d_array_depth_t image, int4 coord, float color);
 #endif //cl_khr_depth_images
 
 // OpenCL Extension v2.0 s9.18 - Mipmaps
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifdef cl_khr_mipmap_image
 void __ovld write_imagef(write_only image1d_t image, int coord, int lod, float4 color);
 void __ovld write_imagei(write_only image1d_t image, int coord, int lod, int4 color);
 void __ovld write_imageui(write_only image1d_t image, int coord, int lod, uint4 color);
 
 void __ovld write_imagef(write_only image1d_array_t image_array, int2 coord, int lod, float4 color);
 void __ovld write_imagei(write_only image1d_array_t image_array, int2 coord, int lod, int4 color);
 void __ovld write_imageui(write_only image1d_array_t image_array, int2 coord, int lod, uint4 color);

 void __ovld write_imagef(write_only image2d_array_depth_t image, int4 coord, int lod, float color);
 
 #ifdef cl_khr_3d_image_writes
 void __ovld write_imagef(write_only image3d_t image, int4 coord, int lod, float4 color);
 void __ovld write_imagei(write_only image3d_t image, int4 coord, int lod, int4 color);
 void __ovld write_imageui(write_only image3d_t image, int4 coord, int lod, uint4 color);
 #endif
 #endif //cl_khr_mipmap_image
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // Image write functions for half4 type
 #ifdef cl_khr_fp16
 void __ovld write_imageh(write_only image1d_t image, int coord, half4 color);
 void __ovld write_imageh(write_only image2d_t image, int2 coord, half4 color);
 #ifdef cl_khr_3d_image_writes
 void __ovld write_imageh(write_only image3d_t image, int4 coord, half4 color);
 #endif
 void __ovld write_imageh(write_only image1d_array_t image, int2 coord, half4 color);
 void __ovld write_imageh(write_only image2d_array_t image, int4 coord, half4 color);
 void __ovld write_imageh(write_only image1d_buffer_t image, int coord, half4 color);
 #endif //cl_khr_fp16
 
 // Image write functions for read_write images
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void __ovld write_imagef(read_write image2d_t image, int2 coord, float4 color);
 void __ovld write_imagei(read_write image2d_t image, int2 coord, int4 color);
 void __ovld write_imageui(read_write image2d_t image, int2 coord, uint4 color);
 
 void __ovld write_imagef(read_write image2d_array_t image_array, int4 coord, float4 color);
 void __ovld write_imagei(read_write image2d_array_t image_array, int4 coord, int4 color);
 void __ovld write_imageui(read_write image2d_array_t image_array, int4 coord, uint4 color);
 

 void __ovld write_imageui(read_write image3d_t image, int4 coord, uint4 color);
 #endif
 
 #ifdef cl_khr_depth_images
 void __ovld write_imagef(read_write image2d_depth_t image, int2 coord, float color);
 void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, float color);
 #endif //cl_khr_depth_images
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifdef cl_khr_mipmap_image
 void __ovld write_imagef(read_write image1d_t image, int coord, int lod, float4 color);
 void __ovld write_imagei(read_write image1d_t image, int coord, int lod, int4 color);
 void __ovld write_imageui(read_write image1d_t image, int coord, int lod, uint4 color);
 
 void __ovld write_imagef(read_write image1d_array_t image_array, int2 coord, int lod, float4 color);
 void __ovld write_imagei(read_write image1d_array_t image_array, int2 coord, int lod, int4 color);
 void __ovld write_imageui(read_write image1d_array_t image_array, int2 coord, int lod, uint4 color);

 void __ovld write_imagef(read_write image2d_array_depth_t image, int4 coord, int lod, float color);
 
 #ifdef cl_khr_3d_image_writes
 void __ovld write_imagef(read_write image3d_t image, int4 coord, int lod, float4 color);
 void __ovld write_imagei(read_write image3d_t image, int4 coord, int lod, int4 color);
 void __ovld write_imageui(read_write image3d_t image, int4 coord, int lod, uint4 color);
 #endif
 #endif //cl_khr_mipmap_image
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // Image write functions for half4 type
 #ifdef cl_khr_fp16
 void __ovld write_imageh(read_write image1d_t image, int coord, half4 color);
 void __ovld write_imageh(read_write image2d_t image, int2 coord, half4 color);
 #ifdef cl_khr_3d_image_writes
 void __ovld write_imageh(read_write image3d_t image, int4 coord, half4 color);
 #endif
 void __ovld write_imageh(read_write image1d_array_t image, int2 coord, half4 color);
 void __ovld write_imageh(read_write image2d_array_t image, int4 coord, half4 color);
 void __ovld write_imageh(read_write image1d_buffer_t image, int coord, half4 color);
 #endif //cl_khr_fp16
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // Note: In OpenCL v1.0/1.1/1.2, image argument of image query builtin functions does not have
 // access qualifier, which by default assume read_only access qualifier. Image query builtin
 // functions with write_only image argument should also be declared.
 
 /**
  * Return the image width in pixels.
  *

 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_width(write_only image2d_msaa_t image);
 int __ovld __cnfn get_image_width(write_only image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_width(write_only image2d_array_msaa_t image);
 int __ovld __cnfn get_image_width(write_only image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __cnfn get_image_width(read_write image1d_t image);
 int __ovld __cnfn get_image_width(read_write image1d_buffer_t image);
 int __ovld __cnfn get_image_width(read_write image2d_t image);
 int __ovld __cnfn get_image_width(read_write image3d_t image);
 int __ovld __cnfn get_image_width(read_write image1d_array_t image);
 int __ovld __cnfn get_image_width(read_write image2d_array_t image);
 #ifdef cl_khr_depth_images
 int __ovld __cnfn get_image_width(read_write image2d_depth_t image);
 int __ovld __cnfn get_image_width(read_write image2d_array_depth_t image);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_width(read_write image2d_msaa_t image);
 int __ovld __cnfn get_image_width(read_write image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_width(read_write image2d_array_msaa_t image);
 int __ovld __cnfn get_image_width(read_write image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the image height in pixels.
  */
 int __ovld __cnfn get_image_height(read_only image2d_t image);
 int __ovld __cnfn get_image_height(read_only image3d_t image);
 int __ovld __cnfn get_image_height(read_only image2d_array_t image);
 #ifdef cl_khr_depth_images

 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_height(write_only image2d_msaa_t image);
 int __ovld __cnfn get_image_height(write_only image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_height(write_only image2d_array_msaa_t image);
 int __ovld __cnfn get_image_height(write_only image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __cnfn get_image_height(read_write image2d_t image);
 int __ovld __cnfn get_image_height(read_write image3d_t image);
 int __ovld __cnfn get_image_height(read_write image2d_array_t image);
 #ifdef cl_khr_depth_images
 int __ovld __cnfn get_image_height(read_write image2d_depth_t image);
 int __ovld __cnfn get_image_height(read_write image2d_array_depth_t image);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_height(read_write image2d_msaa_t image);
 int __ovld __cnfn get_image_height(read_write image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_height(read_write image2d_array_msaa_t image);
 int __ovld __cnfn get_image_height(read_write image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the image depth in pixels.
  */
 int __ovld __cnfn get_image_depth(read_only image3d_t image);
 
 #ifdef cl_khr_3d_image_writes
 int __ovld __cnfn get_image_depth(write_only image3d_t image);
 #endif
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __cnfn get_image_depth(read_write image3d_t image);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL Extension v2.0 s9.18 - Mipmaps
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #ifdef cl_khr_mipmap_image
 /**
  * Return the image miplevels.
  */
 
 int __ovld get_image_num_mip_levels(read_only image1d_t image);
 int __ovld get_image_num_mip_levels(read_only image2d_t image);
 int __ovld get_image_num_mip_levels(read_only image3d_t image);

 int __ovld get_image_num_mip_levels(write_only image2d_depth_t image);
 
 int __ovld get_image_num_mip_levels(read_write image1d_array_t image);
 int __ovld get_image_num_mip_levels(read_write image2d_array_t image);
 int __ovld get_image_num_mip_levels(read_write image2d_array_depth_t image);
 int __ovld get_image_num_mip_levels(read_write image2d_depth_t image);
 
 #endif //cl_khr_mipmap_image
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the channel data type. Valid values are:
  * CLK_SNORM_INT8
  * CLK_SNORM_INT16
  * CLK_UNORM_INT8
  * CLK_UNORM_INT16
  * CLK_UNORM_SHORT_565

 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_channel_data_type(write_only image2d_msaa_t image);
 int __ovld __cnfn get_image_channel_data_type(write_only image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_channel_data_type(write_only image2d_array_msaa_t image);
 int __ovld __cnfn get_image_channel_data_type(write_only image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __cnfn get_image_channel_data_type(read_write image1d_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image1d_buffer_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image3d_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image1d_array_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_array_t image);
 #ifdef cl_khr_depth_images
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_depth_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_array_depth_t image);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_msaa_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_array_msaa_t image);
 int __ovld __cnfn get_image_channel_data_type(read_write image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the image channel order. Valid values are:
  * CLK_A
  * CLK_R
  * CLK_Rx
  * CLK_RG
  * CLK_RGx

 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_channel_order(write_only image2d_msaa_t image);
 int __ovld __cnfn get_image_channel_order(write_only image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_channel_order(write_only image2d_array_msaa_t image);
 int __ovld __cnfn get_image_channel_order(write_only image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __cnfn get_image_channel_order(read_write image1d_t image);
 int __ovld __cnfn get_image_channel_order(read_write image1d_buffer_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_t image);
 int __ovld __cnfn get_image_channel_order(read_write image3d_t image);
 int __ovld __cnfn get_image_channel_order(read_write image1d_array_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_array_t image);
 #ifdef cl_khr_depth_images
 int __ovld __cnfn get_image_channel_order(read_write image2d_depth_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_array_depth_t image);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld __cnfn get_image_channel_order(read_write image2d_msaa_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_msaa_depth_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_array_msaa_t image);
 int __ovld __cnfn get_image_channel_order(read_write image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the 2D image width and height as an int2
  * type. The width is returned in the x component, and
  * the height in the y component.
  */
 int2 __ovld __cnfn get_image_dim(read_only image2d_t image);
 int2 __ovld __cnfn get_image_dim(read_only image2d_array_t image);

 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int2 __ovld __cnfn get_image_dim(write_only image2d_msaa_t image);
 int2 __ovld __cnfn get_image_dim(write_only image2d_msaa_depth_t image);
 int2 __ovld __cnfn get_image_dim(write_only image2d_array_msaa_t image);
 int2 __ovld __cnfn get_image_dim(write_only image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int2 __ovld __cnfn get_image_dim(read_write image2d_t image);
 int2 __ovld __cnfn get_image_dim(read_write image2d_array_t image);
 #ifdef cl_khr_depth_images
 int2 __ovld __cnfn get_image_dim(read_write image2d_array_depth_t image);
 int2 __ovld __cnfn get_image_dim(read_write image2d_depth_t image);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 int2 __ovld __cnfn get_image_dim(read_write image2d_msaa_t image);
 int2 __ovld __cnfn get_image_dim(read_write image2d_msaa_depth_t image);
 int2 __ovld __cnfn get_image_dim(read_write image2d_array_msaa_t image);
 int2 __ovld __cnfn get_image_dim(read_write image2d_array_msaa_depth_t image);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the 3D image width, height, and depth as an
  * int4 type. The width is returned in the x
  * component, height in the y component, depth in the z
  * component and the w component is 0.
  */
 int4 __ovld __cnfn get_image_dim(read_only image3d_t image);
 #ifdef cl_khr_3d_image_writes
 int4 __ovld __cnfn get_image_dim(write_only image3d_t image);
 #endif
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int4 __ovld __cnfn get_image_dim(read_write image3d_t image);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
  * Return the image array size.
  */
 
 size_t __ovld __cnfn get_image_array_size(read_only image1d_array_t image_array);
 size_t __ovld __cnfn get_image_array_size(read_only image2d_array_t image_array);
 #ifdef cl_khr_depth_images

 #ifdef cl_khr_depth_images
 size_t __ovld __cnfn get_image_array_size(write_only image2d_array_depth_t image_array);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 size_t __ovld __cnfn get_image_array_size(write_only image2d_array_msaa_t image_array);
 size_t __ovld __cnfn get_image_array_size(write_only image2d_array_msaa_depth_t image_array);
 #endif //cl_khr_gl_msaa_sharing
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 size_t __ovld __cnfn get_image_array_size(read_write image1d_array_t image_array);
 size_t __ovld __cnfn get_image_array_size(read_write image2d_array_t image_array);
 #ifdef cl_khr_depth_images
 size_t __ovld __cnfn get_image_array_size(read_write image2d_array_depth_t image_array);
 #endif //cl_khr_depth_images
 #if defined(cl_khr_gl_msaa_sharing)
 size_t __ovld __cnfn get_image_array_size(read_write image2d_array_msaa_t image_array);
 size_t __ovld __cnfn get_image_array_size(read_write image2d_array_msaa_depth_t image_array);
 #endif //cl_khr_gl_msaa_sharing
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 /**
 * Return the number of samples associated with image
 */
 #if defined(cl_khr_gl_msaa_sharing)
 int __ovld get_image_num_samples(read_only image2d_msaa_t image);
 int __ovld get_image_num_samples(read_only image2d_msaa_depth_t image);
 int __ovld get_image_num_samples(read_only image2d_array_msaa_t image);
 int __ovld get_image_num_samples(read_only image2d_array_msaa_depth_t image);
 
 int __ovld get_image_num_samples(write_only image2d_msaa_t image);
 int __ovld get_image_num_samples(write_only image2d_msaa_depth_t image);
 int __ovld get_image_num_samples(write_only image2d_array_msaa_t image);
 int __ovld get_image_num_samples(write_only image2d_array_msaa_depth_t image);
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld get_image_num_samples(read_write image2d_msaa_t image);
 int __ovld get_image_num_samples(read_write image2d_msaa_depth_t image);
 int __ovld get_image_num_samples(read_write image2d_array_msaa_t image);
 int __ovld get_image_num_samples(read_write image2d_array_msaa_depth_t image);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 #endif
 
 // OpenCL v2.0 s6.13.15 - Work-group Functions
 
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 int __ovld __conv work_group_all(int predicate);
 int __ovld __conv work_group_any(int predicate);
 
 #ifdef cl_khr_fp16
 half __ovld __conv work_group_broadcast(half a, size_t local_id);
 half __ovld __conv work_group_broadcast(half a, size_t x, size_t y);
 half __ovld __conv work_group_broadcast(half a, size_t x, size_t y, size_t z);
 #endif

 double __ovld __conv work_group_scan_exclusive_add(double x);
 double __ovld __conv work_group_scan_exclusive_min(double x);
 double __ovld __conv work_group_scan_exclusive_max(double x);
 double __ovld __conv work_group_scan_inclusive_add(double x);
 double __ovld __conv work_group_scan_inclusive_min(double x);
 double __ovld __conv work_group_scan_inclusive_max(double x);
 #endif //cl_khr_fp64
 
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL v2.0 s6.13.16 - Pipe Functions
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 bool __ovld is_valid_reserve_id(reserve_id_t reserve_id);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 
 // OpenCL v2.0 s6.13.17 - Enqueue Kernels
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 ndrange_t __ovld ndrange_1D(size_t);
 ndrange_t __ovld ndrange_1D(size_t, size_t);
 ndrange_t __ovld ndrange_1D(size_t, size_t, size_t);
 
 ndrange_t __ovld ndrange_2D(const size_t[2]);
 ndrange_t __ovld ndrange_2D(const size_t[2], const size_t[2]);
 ndrange_t __ovld ndrange_2D(const size_t[2], const size_t[2], const size_t[2]);

 
 void __ovld set_user_event_status(clk_event_t e, int state);
 
 bool __ovld is_valid_event (clk_event_t event);
 
 void __ovld capture_event_profiling_info(clk_event_t, clk_profiling_info, __global void* value);
 
 queue_t __ovld get_default_queue(void);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 // OpenCL Extension v2.0 s9.17 - Sub-groups
 
 #if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
 // Shared Sub Group Functions
 uint    __ovld get_sub_group_size(void);
 uint    __ovld get_max_sub_group_size(void);
 uint    __ovld get_num_sub_groups(void);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 uint    __ovld get_enqueued_num_sub_groups(void);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 uint    __ovld get_sub_group_id(void);
 uint    __ovld get_sub_group_local_id(void);
 
 void    __ovld __conv sub_group_barrier(cl_mem_fence_flags flags);
-#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void    __ovld __conv sub_group_barrier(cl_mem_fence_flags flags, memory_scope scope);
-#endif //__OPENCL_C_VERSION__ >= CL_VERSION_2_0
+#endif //defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 int     __ovld __conv sub_group_all(int predicate);
 int     __ovld __conv sub_group_any(int predicate);
 
 int     __ovld __conv sub_group_broadcast(int   x, uint sub_group_local_id);
 uint    __ovld __conv sub_group_broadcast(uint  x, uint sub_group_local_id);
 long    __ovld __conv sub_group_broadcast(long  x, uint sub_group_local_id);
 ulong   __ovld __conv sub_group_broadcast(ulong x, uint sub_group_local_id);

 long    __ovld __conv intel_sub_group_shuffle_xor( long x, uint c );
 ulong   __ovld __conv intel_sub_group_shuffle_xor( ulong x, uint c );
 
 uint    __ovld __conv intel_sub_group_block_read( read_only image2d_t image, int2 coord );
 uint2   __ovld __conv intel_sub_group_block_read2( read_only image2d_t image, int2 coord );
 uint4   __ovld __conv intel_sub_group_block_read4( read_only image2d_t image, int2 coord );
 uint8   __ovld __conv intel_sub_group_block_read8( read_only image2d_t image, int2 coord );
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 uint    __ovld __conv intel_sub_group_block_read(read_write image2d_t image, int2 coord);
 uint2   __ovld __conv intel_sub_group_block_read2(read_write image2d_t image, int2 coord);
 uint4   __ovld __conv intel_sub_group_block_read4(read_write image2d_t image, int2 coord);
 uint8   __ovld __conv intel_sub_group_block_read8(read_write image2d_t image, int2 coord);
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 uint    __ovld __conv intel_sub_group_block_read( const __global uint* p );
 uint2   __ovld __conv intel_sub_group_block_read2( const __global uint* p );
 uint4   __ovld __conv intel_sub_group_block_read4( const __global uint* p );
 uint8   __ovld __conv intel_sub_group_block_read8( const __global uint* p );
 
 void    __ovld __conv intel_sub_group_block_write(write_only image2d_t image, int2 coord, uint data);
 void    __ovld __conv intel_sub_group_block_write2(write_only image2d_t image, int2 coord, uint2 data);
 void    __ovld __conv intel_sub_group_block_write4(write_only image2d_t image, int2 coord, uint4 data);
 void    __ovld __conv intel_sub_group_block_write8(write_only image2d_t image, int2 coord, uint8 data);
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void    __ovld __conv intel_sub_group_block_write(read_write image2d_t image, int2 coord, uint data);
 void    __ovld __conv intel_sub_group_block_write2(read_write image2d_t image, int2 coord, uint2 data);
 void    __ovld __conv intel_sub_group_block_write4(read_write image2d_t image, int2 coord, uint4 data);
 void    __ovld __conv intel_sub_group_block_write8(read_write image2d_t image, int2 coord, uint8 data);
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 void    __ovld __conv intel_sub_group_block_write( __global uint* p, uint data );
 void    __ovld __conv intel_sub_group_block_write2( __global uint* p, uint2 data );
 void    __ovld __conv intel_sub_group_block_write4( __global uint* p, uint4 data );
 void    __ovld __conv intel_sub_group_block_write8( __global uint* p, uint8 data );
 
 #ifdef cl_khr_fp16
 half    __ovld __conv intel_sub_group_shuffle( half x, uint c );

 short       __ovld __conv intel_sub_group_scan_inclusive_max( short   x );
 ushort      __ovld __conv intel_sub_group_scan_inclusive_max( ushort  x );
 
 uint       __ovld __conv intel_sub_group_block_read_ui( read_only image2d_t image, int2 byte_coord );
 uint2      __ovld __conv intel_sub_group_block_read_ui2( read_only image2d_t image, int2 byte_coord );
 uint4      __ovld __conv intel_sub_group_block_read_ui4( read_only image2d_t image, int2 byte_coord );
 uint8      __ovld __conv intel_sub_group_block_read_ui8( read_only image2d_t image, int2 byte_coord );
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 uint       __ovld __conv intel_sub_group_block_read_ui( read_write image2d_t image, int2 byte_coord );
 uint2      __ovld __conv intel_sub_group_block_read_ui2( read_write image2d_t image, int2 byte_coord );
 uint4      __ovld __conv intel_sub_group_block_read_ui4( read_write image2d_t image, int2 byte_coord );
 uint8      __ovld __conv intel_sub_group_block_read_ui8( read_write image2d_t image, int2 byte_coord );
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 uint       __ovld __conv intel_sub_group_block_read_ui( const __global uint* p );
 uint2      __ovld __conv intel_sub_group_block_read_ui2( const __global uint* p );
 uint4      __ovld __conv intel_sub_group_block_read_ui4( const __global uint* p );
 uint8      __ovld __conv intel_sub_group_block_read_ui8( const __global uint* p );
 
 void       __ovld __conv intel_sub_group_block_write_ui( read_only image2d_t image, int2 byte_coord, uint data );
 void       __ovld __conv intel_sub_group_block_write_ui2( read_only image2d_t image, int2 byte_coord, uint2 data );
 void       __ovld __conv intel_sub_group_block_write_ui4( read_only image2d_t image, int2 byte_coord, uint4 data );
 void       __ovld __conv intel_sub_group_block_write_ui8( read_only image2d_t image, int2 byte_coord, uint8 data );
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void       __ovld __conv intel_sub_group_block_write_ui( read_write image2d_t image, int2 byte_coord, uint data );
 void       __ovld __conv intel_sub_group_block_write_ui2( read_write image2d_t image, int2 byte_coord, uint2 data );
 void       __ovld __conv intel_sub_group_block_write_ui4( read_write image2d_t image, int2 byte_coord, uint4 data );
 void       __ovld __conv intel_sub_group_block_write_ui8( read_write image2d_t image, int2 byte_coord, uint8 data );
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 void       __ovld __conv intel_sub_group_block_write_ui( __global uint* p, uint data );
 void       __ovld __conv intel_sub_group_block_write_ui2( __global uint* p, uint2 data );
 void       __ovld __conv intel_sub_group_block_write_ui4( __global uint* p, uint4 data );
 void       __ovld __conv intel_sub_group_block_write_ui8( __global uint* p, uint8 data );
 
 ushort      __ovld __conv intel_sub_group_block_read_us( read_only image2d_t image, int2 coord );
 ushort2     __ovld __conv intel_sub_group_block_read_us2( read_only image2d_t image, int2 coord );
 ushort4     __ovld __conv intel_sub_group_block_read_us4( read_only image2d_t image, int2 coord );
 ushort8     __ovld __conv intel_sub_group_block_read_us8( read_only image2d_t image, int2 coord );
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 ushort      __ovld __conv intel_sub_group_block_read_us(read_write image2d_t image, int2 coord);
 ushort2     __ovld __conv intel_sub_group_block_read_us2(read_write image2d_t image, int2 coord);
 ushort4     __ovld __conv intel_sub_group_block_read_us4(read_write image2d_t image, int2 coord);
 ushort8     __ovld __conv intel_sub_group_block_read_us8(read_write image2d_t image, int2 coord);
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 ushort      __ovld __conv intel_sub_group_block_read_us(  const __global ushort* p );
 ushort2     __ovld __conv intel_sub_group_block_read_us2( const __global ushort* p );
 ushort4     __ovld __conv intel_sub_group_block_read_us4( const __global ushort* p );
 ushort8     __ovld __conv intel_sub_group_block_read_us8( const __global ushort* p );
 
 void        __ovld __conv intel_sub_group_block_write_us(write_only image2d_t image, int2 coord, ushort  data);
 void        __ovld __conv intel_sub_group_block_write_us2(write_only image2d_t image, int2 coord, ushort2 data);
 void        __ovld __conv intel_sub_group_block_write_us4(write_only image2d_t image, int2 coord, ushort4 data);
 void        __ovld __conv intel_sub_group_block_write_us8(write_only image2d_t image, int2 coord, ushort8 data);
 
-#if (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 void        __ovld __conv intel_sub_group_block_write_us(read_write image2d_t image, int2 coord, ushort  data);
 void        __ovld __conv intel_sub_group_block_write_us2(read_write image2d_t image, int2 coord, ushort2 data);
 void        __ovld __conv intel_sub_group_block_write_us4(read_write image2d_t image, int2 coord, ushort4 data);
 void        __ovld __conv intel_sub_group_block_write_us8(read_write image2d_t image, int2 coord, ushort8 data);
-#endif // (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
+#endif // defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_2_0)
 
 void        __ovld __conv intel_sub_group_block_write_us(  __global ushort* p, ushort  data );
 void        __ovld __conv intel_sub_group_block_write_us2( __global ushort* p, ushort2 data );
 void        __ovld __conv intel_sub_group_block_write_us4( __global ushort* p, ushort4 data );
 void        __ovld __conv intel_sub_group_block_write_us8( __global ushort* p, ushort8 data );
 #endif // cl_intel_subgroups_short
 
 #ifdef cl_intel_device_side_avc_motion_estimation