diff --git a/CMakeLists.txt b/CMakeLists.txt
index 35fb3c7db78cb9962a314f36da8933ce2900205d..775c07d2ca33794855f7bbede8760df23aeab860 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1308,7 +1308,7 @@ set(HOST_DEVICE_CL_VERSION_MINOR 0)
 set(HOST_DEVICE_EXTENSIONS "cl_khr_byte_addressable_store cl_khr_global_int32_base_atomics \
 cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics \
 cl_khr_local_int32_extended_atomics cl_khr_3d_image_writes \
-cl_khr_command_buffer cl_khr_command_buffer_multi_device cl_khr_subgroups \
+cl_khr_command_buffer cl_khr_command_buffer_multi_device \
 cl_intel_unified_shared_memory cl_ext_buffer_device_address")
 
 # Host CPU device: list of OpenCL 3.0 features that are always enabled
@@ -1332,7 +1332,7 @@ if(NOT ENABLE_CONFORMANCE)
     set(HOST_DEVICE_EXTENSIONS "${HOST_DEVICE_EXTENSIONS} \
       cl_exp_tensor cl_exp_defined_builtin_kernels")
   endif()
-  set(HOST_DEVICE_EXTENSIONS "${HOST_DEVICE_EXTENSIONS} cl_khr_subgroup_ballot \
+  set(HOST_DEVICE_EXTENSIONS "${HOST_DEVICE_EXTENSIONS} cl_khr_subgroups cl_khr_subgroup_ballot \
 cl_khr_subgroup_shuffle cl_intel_subgroups cl_intel_subgroups_short cl_intel_subgroups_char \
 cl_ext_float_atomics cl_intel_required_subgroup_size")
   # read-write images are still partially broken
@@ -1365,6 +1365,9 @@ endif()
 #   Those with either 1) expression implementation, or 2) Clang builtin
 #   implementation.
 #
+# * cl_khr_subgroup: still need some patches to OpenCL-CTS, which are
+#   not upstreamed. Removing this extension from conformance-enabled build
+#   is required for PoCL CPU device pass the official upstream OpenCL-CTS.
 
 if(HOST_DEVICE_EXTENSIONS MATCHES "cl_khr_subgroup")
   set(HOST_DEVICE_FEATURES_30 "${HOST_DEVICE_FEATURES_30} __opencl_c_subgroups")
diff --git a/examples/conformance/CMakeLists.txt b/examples/conformance/CMakeLists.txt
index 13cfb68bee7b29e47c5b636c124536959570c4e1..d7f1f7aa46189367e77e23b865fba18cc8a63d44 100644
--- a/examples/conformance/CMakeLists.txt
+++ b/examples/conformance/CMakeLists.txt
@@ -85,7 +85,7 @@ if(CUSTOM_CTS_GIT_TAG)
   set(CTS_GIT_TAG "${CUSTOM_CTS_GIT_TAG}")
 else()
   # Use PoCL's fork which has pending fixes.
-  set(CTS_GIT_TAG "v2024.08.28")
+  set(CTS_GIT_TAG "v2024.10.09")
 endif()
 
 set(TS_BUILDDIR "${TS_BUILDDIR}/test_conformance")
@@ -1045,13 +1045,16 @@ set(CTS_LEVEL0_LIST
   conformance_main_mem_host_flags
   conformance_main_multiples
   conformance_main_non_uniform_work_group
-  conformance_main_printf
+  # disabled. Fails since CTS 2024.10.09; this CTS added
+  # testcase printf("%s", "") which causes crash in L0
+  #conformance_main_printf
   conformance_main_profiling
   conformance_main_subgroups_micro
   conformance_main_cl_khr_command_buffer
   conformance_main_cl_khr_command_buffer_mutable
   conformance_main_api_micro_ocl_30
   conformance_main_basic_micro_other_ocl_30
+  # disabled. Fails because of a bug in the L0 driver
   #conformance_main_basic_micro_progvar
   conformance_main_SVM_micro_1
   conformance_main_SVM_micro_2
@@ -1106,8 +1109,10 @@ set(CTS_LEVEL0_LIST
   conformance_main_math_micro_cosh
   conformance_main_math_micro_cospi
 
-  conformance_main_math_micro_divide
-  conformance_main_math_micro_divide_cr
+  # disabled. Fails since CTS 2024.10.09;
+  # this CTS checks FP16 ULP precision
+  # conformance_main_math_micro_divide
+  # conformance_main_math_micro_divide_cr
 
   conformance_main_math_micro_exp
   conformance_main_math_micro_exp2
@@ -1121,7 +1126,9 @@ set(CTS_LEVEL0_LIST
   conformance_main_math_micro_fmax
   conformance_main_math_micro_fmin
   conformance_main_math_micro_fmod
-  conformance_main_math_micro_fract
+  # disabled. Fails since CTS 2024.10.09;
+  # this CTS checks FP16 ULP precision
+  # conformance_main_math_micro_fract
   conformance_main_math_micro_frexp
 
   conformance_main_math_micro_hypot
diff --git a/include/CL/cl.h b/include/CL/cl.h
index afeeb4ee5e587b9771af1efcd3ed950849360113..792e20cc8a1d6c405d29329dc0a737f658827094 100644
--- a/include/CL/cl.h
+++ b/include/CL/cl.h
@@ -112,9 +112,9 @@ typedef cl_uint             cl_kernel_exec_info;
 typedef cl_bitfield         cl_device_atomic_capabilities;
 typedef cl_bitfield         cl_device_device_enqueue_capabilities;
 typedef cl_uint             cl_khronos_vendor_id;
-typedef cl_properties       cl_mem_properties;
-typedef cl_uint             cl_version;
+typedef cl_properties cl_mem_properties;
 #endif
+typedef cl_uint cl_version;
 
 typedef struct _cl_image_format {
     cl_channel_order        image_channel_order;
@@ -914,8 +914,6 @@ typedef struct _cl_name_version {
 /* cl_khronos_vendor_id */
 #define CL_KHRONOS_VENDOR_ID_CODEPLAY               0x10004
 
-#ifdef CL_VERSION_3_0
-
 /* cl_version */
 #define CL_VERSION_MAJOR_BITS (10)
 #define CL_VERSION_MINOR_BITS (10)
@@ -939,8 +937,6 @@ typedef struct _cl_name_version {
    (((minor) & CL_VERSION_MINOR_MASK) << CL_VERSION_PATCH_BITS) | \
    ((patch) & CL_VERSION_PATCH_MASK))
 
-#endif
-
 /********************************************************************************************************/
 
 /* CL_NO_PROTOTYPES implies CL_NO_CORE_PROTOTYPES: */
diff --git a/include/CL/cl_d3d10.h b/include/CL/cl_d3d10.h
index 8404644a471ff7f2dea7792bd4a63425752a0aa8..6b56c775bf8a53bbabb7cf6c6afe4b1de81fe3ae 100644
--- a/include/CL/cl_d3d10.h
+++ b/include/CL/cl_d3d10.h
@@ -65,6 +65,9 @@ extern "C" {
 #define CL_KHR_D3D10_SHARING_EXTENSION_NAME \
     "cl_khr_d3d10_sharing"
 
+
+#define CL_KHR_D3D10_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_uint             cl_d3d10_device_source_khr;
 typedef cl_uint             cl_d3d10_device_set_khr;
 
@@ -228,6 +231,9 @@ clEnqueueReleaseD3D10ObjectsKHR(
 #define CL_INTEL_SHARING_FORMAT_QUERY_D3D10_EXTENSION_NAME \
     "cl_intel_sharing_format_query_d3d10"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_D3D10_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* when cl_khr_d3d10_sharing is supported */
 
 typedef cl_int CL_API_CALL
diff --git a/include/CL/cl_d3d11.h b/include/CL/cl_d3d11.h
index ade879509cf1b41126b986c93a7686126e23a739..384c8f428fc5b8560fc9a3dfa89e9b85dc80420d 100644
--- a/include/CL/cl_d3d11.h
+++ b/include/CL/cl_d3d11.h
@@ -65,6 +65,9 @@ extern "C" {
 #define CL_KHR_D3D11_SHARING_EXTENSION_NAME \
     "cl_khr_d3d11_sharing"
 
+
+#define CL_KHR_D3D11_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_uint             cl_d3d11_device_source_khr;
 typedef cl_uint             cl_d3d11_device_set_khr;
 
@@ -228,6 +231,9 @@ clEnqueueReleaseD3D11ObjectsKHR(
 #define CL_INTEL_SHARING_FORMAT_QUERY_D3D11_EXTENSION_NAME \
     "cl_intel_sharing_format_query_d3d11"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_D3D11_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* when cl_khr_d3d11_sharing is supported */
 
 typedef cl_int CL_API_CALL
diff --git a/include/CL/cl_dx9_media_sharing.h b/include/CL/cl_dx9_media_sharing.h
index c0df5c919894b4bac70da73101177e10d3196f68..b079379d0a726970ffc4cdb5f857a029b110b783 100644
--- a/include/CL/cl_dx9_media_sharing.h
+++ b/include/CL/cl_dx9_media_sharing.h
@@ -67,6 +67,9 @@ extern "C" {
 #define CL_KHR_DX9_MEDIA_SHARING_EXTENSION_NAME \
     "cl_khr_dx9_media_sharing"
 
+
+#define CL_KHR_DX9_MEDIA_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_uint             cl_dx9_media_adapter_type_khr;
 typedef cl_uint             cl_dx9_media_adapter_set_khr;
 
@@ -209,6 +212,9 @@ clEnqueueReleaseDX9MediaSurfacesKHR(
 #define CL_INTEL_DX9_MEDIA_SHARING_EXTENSION_NAME \
     "cl_intel_dx9_media_sharing"
 
+
+#define CL_INTEL_DX9_MEDIA_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_uint             cl_dx9_device_source_intel;
 typedef cl_uint             cl_dx9_device_set_intel;
 
@@ -341,6 +347,9 @@ clEnqueueReleaseDX9ObjectsINTEL(
 #define CL_INTEL_SHARING_FORMAT_QUERY_DX9_EXTENSION_NAME \
     "cl_intel_sharing_format_query_dx9"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_DX9_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* when cl_khr_dx9_media_sharing or cl_intel_dx9_media_sharing is supported */
 
 typedef cl_int CL_API_CALL
diff --git a/include/CL/cl_egl.h b/include/CL/cl_egl.h
index 25cd5e0cb565e8cf06999295a162a98778b96113..68aefec7619fb42928e2afbf313f0f4b9d324c38 100644
--- a/include/CL/cl_egl.h
+++ b/include/CL/cl_egl.h
@@ -51,6 +51,9 @@ extern "C" {
 #define CL_KHR_EGL_IMAGE_EXTENSION_NAME \
     "cl_khr_egl_image"
 
+
+#define CL_KHR_EGL_IMAGE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* Command type for events created with clEnqueueAcquireEGLObjectsKHR */
 #define CL_COMMAND_EGL_FENCE_SYNC_OBJECT_KHR                0x202F
 #define CL_COMMAND_ACQUIRE_EGL_OBJECTS_KHR                  0x202D
@@ -144,6 +147,9 @@ clEnqueueReleaseEGLObjectsKHR(
 #define CL_KHR_EGL_EVENT_EXTENSION_NAME \
     "cl_khr_egl_event"
 
+
+#define CL_KHR_EGL_EVENT_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* CLeglDisplayKHR is an opaque handle to an EGLDisplay */
 /* type CLeglDisplayKHR */
 
diff --git a/include/CL/cl_ext.h b/include/CL/cl_ext.h
index b2db99270b3bded53116ae643b69dd58fc2f8c42..d6ca6982c90a3cc572ef00272f7986baf7d88197 100644
--- a/include/CL/cl_ext.h
+++ b/include/CL/cl_ext.h
@@ -51,6 +51,9 @@ extern "C" {
 #define CL_KHR_COMMAND_BUFFER_EXTENSION_NAME \
     "cl_khr_command_buffer"
 
+
+#define CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION CL_MAKE_VERSION(0, 9, 5)
+
 typedef cl_bitfield         cl_device_command_buffer_capabilities_khr;
 typedef struct _cl_command_buffer_khr* cl_command_buffer_khr;
 typedef cl_uint             cl_sync_point_khr;
@@ -58,7 +61,7 @@ typedef cl_uint             cl_command_buffer_info_khr;
 typedef cl_uint             cl_command_buffer_state_khr;
 typedef cl_properties       cl_command_buffer_properties_khr;
 typedef cl_bitfield         cl_command_buffer_flags_khr;
-typedef cl_properties       cl_ndrange_kernel_command_properties_khr;
+typedef cl_properties       cl_command_properties_khr;
 typedef struct _cl_mutable_command_khr* cl_mutable_command_khr;
 
 /* cl_device_info */
@@ -146,6 +149,7 @@ typedef cl_int CL_API_CALL
 clCommandBarrierWithWaitListKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr* sync_point_wait_list,
     cl_sync_point_khr* sync_point,
@@ -158,6 +162,7 @@ typedef cl_int CL_API_CALL
 clCommandCopyBufferKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_buffer,
     size_t src_offset,
@@ -175,6 +180,7 @@ typedef cl_int CL_API_CALL
 clCommandCopyBufferRectKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_buffer,
     const size_t* src_origin,
@@ -196,6 +202,7 @@ typedef cl_int CL_API_CALL
 clCommandCopyBufferToImageKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_image,
     size_t src_offset,
@@ -213,6 +220,7 @@ typedef cl_int CL_API_CALL
 clCommandCopyImageKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image,
     cl_mem dst_image,
     const size_t* src_origin,
@@ -230,6 +238,7 @@ typedef cl_int CL_API_CALL
 clCommandCopyImageToBufferKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image,
     cl_mem dst_buffer,
     const size_t* src_origin,
@@ -247,6 +256,7 @@ typedef cl_int CL_API_CALL
 clCommandFillBufferKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem buffer,
     const void* pattern,
     size_t pattern_size,
@@ -264,6 +274,7 @@ typedef cl_int CL_API_CALL
 clCommandFillImageKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem image,
     const void* fill_color,
     const size_t* origin,
@@ -280,7 +291,7 @@ typedef cl_int CL_API_CALL
 clCommandNDRangeKernelKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
-    const cl_ndrange_kernel_command_properties_khr* properties,
+    const cl_command_properties_khr* properties,
     cl_kernel kernel,
     cl_uint work_dim,
     const size_t* global_work_offset,
@@ -339,6 +350,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandBarrierWithWaitListKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr* sync_point_wait_list,
     cl_sync_point_khr* sync_point,
@@ -348,6 +360,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandCopyBufferKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_buffer,
     size_t src_offset,
@@ -362,6 +375,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandCopyBufferRectKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_buffer,
     const size_t* src_origin,
@@ -380,6 +394,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandCopyBufferToImageKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer,
     cl_mem dst_image,
     size_t src_offset,
@@ -394,6 +409,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandCopyImageKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image,
     cl_mem dst_image,
     const size_t* src_origin,
@@ -408,6 +424,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandCopyImageToBufferKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image,
     cl_mem dst_buffer,
     const size_t* src_origin,
@@ -422,6 +439,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandFillBufferKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem buffer,
     const void* pattern,
     size_t pattern_size,
@@ -436,6 +454,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandFillImageKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem image,
     const void* fill_color,
     const size_t* origin,
@@ -449,7 +468,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandNDRangeKernelKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
-    const cl_ndrange_kernel_command_properties_khr* properties,
+    const cl_command_properties_khr* properties,
     cl_kernel kernel,
     cl_uint work_dim,
     const size_t* global_work_offset,
@@ -476,6 +495,7 @@ typedef cl_int CL_API_CALL
 clCommandSVMMemcpyKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void* dst_ptr,
     const void* src_ptr,
     size_t size,
@@ -491,6 +511,7 @@ typedef cl_int CL_API_CALL
 clCommandSVMMemFillKHR_t(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void* svm_ptr,
     const void* pattern,
     size_t pattern_size,
@@ -509,6 +530,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandSVMMemcpyKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void* dst_ptr,
     const void* src_ptr,
     size_t size,
@@ -521,6 +543,7 @@ extern CL_API_ENTRY cl_int CL_API_CALL
 clCommandSVMMemFillKHR(
     cl_command_buffer_khr command_buffer,
     cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void* svm_ptr,
     const void* pattern,
     size_t pattern_size,
@@ -539,6 +562,9 @@ clCommandSVMMemFillKHR(
 #define CL_KHR_COMMAND_BUFFER_MULTI_DEVICE_EXTENSION_NAME \
     "cl_khr_command_buffer_multi_device"
 
+
+#define CL_KHR_COMMAND_BUFFER_MULTI_DEVICE_EXTENSION_VERSION CL_MAKE_VERSION(0, 9, 1)
+
 typedef cl_bitfield         cl_platform_command_buffer_capabilities_khr;
 
 /* cl_platform_info */
@@ -596,7 +622,10 @@ clRemapCommandBufferKHR(
 #define CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_EXTENSION_NAME \
     "cl_khr_command_buffer_mutable_dispatch"
 
-typedef cl_uint             cl_command_buffer_structure_type_khr;
+
+#define CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_EXTENSION_VERSION CL_MAKE_VERSION(0, 9, 3)
+
+typedef cl_uint             cl_command_buffer_update_type_khr;
 typedef cl_bitfield         cl_mutable_dispatch_fields_khr;
 typedef cl_uint             cl_mutable_command_info_khr;
 typedef struct _cl_mutable_dispatch_arg_khr {
@@ -610,8 +639,6 @@ typedef struct _cl_mutable_dispatch_exec_info_khr {
     const void* param_value;
 } cl_mutable_dispatch_exec_info_khr;
 typedef struct _cl_mutable_dispatch_config_khr {
-    cl_command_buffer_structure_type_khr type;
-    const void* next;
     cl_mutable_command_khr command;
     cl_uint num_args;
     cl_uint num_svm_args;
@@ -624,12 +651,6 @@ typedef struct _cl_mutable_dispatch_config_khr {
     const size_t* global_work_size;
     const size_t* local_work_size;
 } cl_mutable_dispatch_config_khr;
-typedef struct _cl_mutable_base_config_khr {
-    cl_command_buffer_structure_type_khr type;
-    const void* next;
-    cl_uint num_mutable_dispatch;
-    const cl_mutable_dispatch_config_khr* mutable_dispatch_list;
-} cl_mutable_base_config_khr;
 typedef cl_bitfield         cl_mutable_dispatch_asserts_khr;
 
 /* cl_command_buffer_flags_khr - bitfield */
@@ -641,7 +662,7 @@ typedef cl_bitfield         cl_mutable_dispatch_asserts_khr;
 /* cl_device_info */
 #define CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR         0x12B0
 
-/* cl_ndrange_kernel_command_properties_khr */
+/* cl_command_properties_khr */
 #define CL_MUTABLE_DISPATCH_UPDATABLE_FIELDS_KHR            0x12B1
 
 /* cl_mutable_dispatch_fields_khr - bitfield */
@@ -655,21 +676,20 @@ typedef cl_bitfield         cl_mutable_dispatch_asserts_khr;
 #define CL_MUTABLE_COMMAND_COMMAND_QUEUE_KHR                0x12A0
 #define CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR               0x12A1
 #define CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR                 0x12AD
-#define CL_MUTABLE_DISPATCH_PROPERTIES_ARRAY_KHR            0x12A2
+#define CL_MUTABLE_COMMAND_PROPERTIES_ARRAY_KHR             0x12A2
 #define CL_MUTABLE_DISPATCH_KERNEL_KHR                      0x12A3
 #define CL_MUTABLE_DISPATCH_DIMENSIONS_KHR                  0x12A4
 #define CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR          0x12A5
 #define CL_MUTABLE_DISPATCH_GLOBAL_WORK_SIZE_KHR            0x12A6
 #define CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR             0x12A7
 
-/* cl_command_buffer_structure_type_khr */
-#define CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR           0
-#define CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR       1
+/* cl_command_buffer_update_type_khr */
+#define CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR       0
 
 /* cl_command_buffer_properties_khr */
 #define CL_COMMAND_BUFFER_MUTABLE_DISPATCH_ASSERTS_KHR      0x12B7
 
-/* cl_ndrange_kernel_command_properties_khr */
+/* cl_command_properties_khr */
 #define CL_MUTABLE_DISPATCH_ASSERTS_KHR                     0x12B8
 
 /* cl_mutable_dispatch_asserts_khr - bitfield */
@@ -679,7 +699,9 @@ typedef cl_bitfield         cl_mutable_dispatch_asserts_khr;
 typedef cl_int CL_API_CALL
 clUpdateMutableCommandsKHR_t(
     cl_command_buffer_khr command_buffer,
-    const cl_mutable_base_config_khr* mutable_config);
+    cl_uint num_configs,
+    const cl_command_buffer_update_type_khr* config_types,
+    const void** configs);
 
 typedef clUpdateMutableCommandsKHR_t *
 clUpdateMutableCommandsKHR_fn ;
@@ -700,7 +722,9 @@ clGetMutableCommandInfoKHR_fn ;
 extern CL_API_ENTRY cl_int CL_API_CALL
 clUpdateMutableCommandsKHR(
     cl_command_buffer_khr command_buffer,
-    const cl_mutable_base_config_khr* mutable_config) ;
+    cl_uint num_configs,
+    const cl_command_buffer_update_type_khr* config_types,
+    const void** configs) ;
 
 extern CL_API_ENTRY cl_int CL_API_CALL
 clGetMutableCommandInfoKHR(
@@ -719,6 +743,9 @@ clGetMutableCommandInfoKHR(
 #define CL_KHR_FP64_EXTENSION_NAME \
     "cl_khr_fp64"
 
+
+#define CL_KHR_FP64_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 #if !defined(CL_VERSION_1_2)
 /* cl_device_info - defined in CL.h for OpenCL 1.2 and newer */
 #define CL_DEVICE_DOUBLE_FP_CONFIG                          0x1032
@@ -732,6 +759,9 @@ clGetMutableCommandInfoKHR(
 #define CL_KHR_FP16_EXTENSION_NAME \
     "cl_khr_fp16"
 
+
+#define CL_KHR_FP16_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_HALF_FP_CONFIG                            0x1033
 
@@ -743,6 +773,9 @@ clGetMutableCommandInfoKHR(
     "cl_APPLE_SetMemObjectDestructor"
 
 
+#define CL_APPLE_SETMEMOBJECTDESTRUCTOR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
+
 typedef cl_int CL_API_CALL
 clSetMemObjectDestructorAPPLE_t(
     cl_mem memobj,
@@ -770,6 +803,9 @@ clSetMemObjectDestructorAPPLE(
     "cl_APPLE_ContextLoggingFunctions"
 
 
+#define CL_APPLE_CONTEXTLOGGINGFUNCTIONS_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
+
 typedef void CL_API_CALL
 clLogMessagesToSystemLogAPPLE_t(
     const char* errstr,
@@ -832,6 +868,9 @@ clLogMessagesToStderrAPPLE(
 #define CL_KHR_ICD_EXTENSION_NAME \
     "cl_khr_icd"
 
+
+#define CL_KHR_ICD_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_platform_info */
 #define CL_PLATFORM_ICD_SUFFIX_KHR                          0x0920
 
@@ -865,6 +904,9 @@ clIcdGetPlatformIDsKHR(
 #define CL_KHR_IL_PROGRAM_EXTENSION_NAME \
     "cl_khr_il_program"
 
+
+#define CL_KHR_IL_PROGRAM_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_IL_VERSION_KHR                            0x105B
 
@@ -900,6 +942,9 @@ clCreateProgramWithILKHR(
 #define CL_KHR_IMAGE2D_FROM_BUFFER_EXTENSION_NAME \
     "cl_khr_image2d_from_buffer"
 
+
+#define CL_KHR_IMAGE2D_FROM_BUFFER_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_IMAGE_PITCH_ALIGNMENT_KHR                 0x104A
 #define CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT_KHR          0x104B
@@ -911,6 +956,9 @@ clCreateProgramWithILKHR(
 #define CL_KHR_INITIALIZE_MEMORY_EXTENSION_NAME \
     "cl_khr_initialize_memory"
 
+
+#define CL_KHR_INITIALIZE_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_bitfield         cl_context_memory_initialize_khr;
 
 /* cl_context_properties */
@@ -927,6 +975,9 @@ typedef cl_bitfield         cl_context_memory_initialize_khr;
 #define CL_KHR_TERMINATE_CONTEXT_EXTENSION_NAME \
     "cl_khr_terminate_context"
 
+
+#define CL_KHR_TERMINATE_CONTEXT_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_bitfield         cl_device_terminate_capability_khr;
 
 /* cl_device_info */
@@ -964,6 +1015,9 @@ clTerminateContextKHR(
 #define CL_KHR_SPIR_EXTENSION_NAME \
     "cl_khr_spir"
 
+
+#define CL_KHR_SPIR_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_SPIR_VERSIONS                             0x40E0
 
@@ -977,6 +1031,9 @@ clTerminateContextKHR(
 #define CL_KHR_CREATE_COMMAND_QUEUE_EXTENSION_NAME \
     "cl_khr_create_command_queue"
 
+
+#define CL_KHR_CREATE_COMMAND_QUEUE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_properties       cl_queue_properties_khr;
 
 
@@ -1008,6 +1065,9 @@ clCreateCommandQueueWithPropertiesKHR(
 #define CL_NV_DEVICE_ATTRIBUTE_QUERY_EXTENSION_NAME \
     "cl_nv_device_attribute_query"
 
+
+#define CL_NV_DEVICE_ATTRIBUTE_QUERY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV               0x4000
 #define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV               0x4001
@@ -1030,6 +1090,9 @@ clCreateCommandQueueWithPropertiesKHR(
 #define CL_AMD_DEVICE_ATTRIBUTE_QUERY_EXTENSION_NAME \
     "cl_amd_device_attribute_query"
 
+
+#define CL_AMD_DEVICE_ATTRIBUTE_QUERY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_PROFILING_TIMER_OFFSET_AMD                0x4036
 #define CL_DEVICE_TOPOLOGY_AMD                              0x4037
@@ -1060,6 +1123,9 @@ clCreateCommandQueueWithPropertiesKHR(
 #define CL_ARM_PRINTF_EXTENSION_NAME \
     "cl_arm_printf"
 
+
+#define CL_ARM_PRINTF_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_context_properties */
 #define CL_PRINTF_CALLBACK_ARM                              0x40B0
 #define CL_PRINTF_BUFFERSIZE_ARM                            0x40B1
@@ -1071,6 +1137,9 @@ clCreateCommandQueueWithPropertiesKHR(
 #define CL_EXT_DEVICE_FISSION_EXTENSION_NAME \
     "cl_ext_device_fission"
 
+
+#define CL_EXT_DEVICE_FISSION_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_ulong            cl_device_partition_property_ext;
 
 /* Error codes */
@@ -1157,6 +1226,9 @@ clCreateSubDevicesEXT(
 #define CL_EXT_MIGRATE_MEMOBJECT_EXTENSION_NAME \
     "cl_ext_migrate_memobject"
 
+
+#define CL_EXT_MIGRATE_MEMOBJECT_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_bitfield         cl_mem_migration_flags_ext;
 
 /* cl_mem_migration_flags_ext */
@@ -1200,6 +1272,9 @@ clEnqueueMigrateMemObjectEXT(
 #define CL_EXT_CXX_FOR_OPENCL_EXTENSION_NAME \
     "cl_ext_cxx_for_opencl"
 
+
+#define CL_EXT_CXX_FOR_OPENCL_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_CXX_FOR_OPENCL_NUMERIC_VERSION_EXT        0x4230
 
@@ -1210,6 +1285,9 @@ clEnqueueMigrateMemObjectEXT(
 #define CL_QCOM_EXT_HOST_PTR_EXTENSION_NAME \
     "cl_qcom_ext_host_ptr"
 
+
+#define CL_QCOM_EXT_HOST_PTR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_uint             cl_image_pitch_info_qcom;
 typedef struct _cl_mem_ext_host_ptr {
     cl_uint allocation_type;
@@ -1270,6 +1348,9 @@ clGetDeviceImageInfoQCOM(
 #define CL_QCOM_EXT_HOST_PTR_IOCOHERENT_EXTENSION_NAME \
     "cl_qcom_ext_host_ptr_iocoherent"
 
+
+#define CL_QCOM_EXT_HOST_PTR_IOCOHERENT_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_uint host_cache_policy */
 #define CL_MEM_HOST_IOCOHERENT_QCOM                         0x40A9
 
@@ -1280,6 +1361,9 @@ clGetDeviceImageInfoQCOM(
 #define CL_QCOM_ION_HOST_PTR_EXTENSION_NAME \
     "cl_qcom_ion_host_ptr"
 
+
+#define CL_QCOM_ION_HOST_PTR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* type cl_mem_ext_host_ptr */
 typedef struct _cl_mem_ion_host_ptr {
     cl_mem_ext_host_ptr ext_host_ptr;
@@ -1297,6 +1381,9 @@ typedef struct _cl_mem_ion_host_ptr {
 #define CL_QCOM_ANDROID_NATIVE_BUFFER_HOST_PTR_EXTENSION_NAME \
     "cl_qcom_android_native_buffer_host_ptr"
 
+
+#define CL_QCOM_ANDROID_NATIVE_BUFFER_HOST_PTR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* type cl_mem_ext_host_ptr */
 typedef struct _cl_mem_android_native_buffer_host_ptr {
     cl_mem_ext_host_ptr ext_host_ptr;
@@ -1313,6 +1400,9 @@ typedef struct _cl_mem_android_native_buffer_host_ptr {
 #define CL_IMG_YUV_IMAGE_EXTENSION_NAME \
     "cl_img_yuv_image"
 
+
+#define CL_IMG_YUV_IMAGE_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_channel_order */
 #define CL_NV21_IMG                                         0x40D0
 #define CL_YV12_IMG                                         0x40D1
@@ -1324,6 +1414,9 @@ typedef struct _cl_mem_android_native_buffer_host_ptr {
 #define CL_IMG_CACHED_ALLOCATIONS_EXTENSION_NAME \
     "cl_img_cached_allocations"
 
+
+#define CL_IMG_CACHED_ALLOCATIONS_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_flags */
 #define CL_MEM_USE_UNCACHED_CPU_MEMORY_IMG                  (1 << 26)
 #define CL_MEM_USE_CACHED_CPU_MEMORY_IMG                    (1 << 27)
@@ -1335,6 +1428,9 @@ typedef struct _cl_mem_android_native_buffer_host_ptr {
 #define CL_IMG_USE_GRALLOC_PTR_EXTENSION_NAME \
     "cl_img_use_gralloc_ptr"
 
+
+#define CL_IMG_USE_GRALLOC_PTR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* Error codes */
 #define CL_GRALLOC_RESOURCE_NOT_ACQUIRED_IMG                0x40D4
 #define CL_INVALID_GRALLOC_OBJECT_IMG                       0x40D5
@@ -1400,6 +1496,9 @@ clEnqueueReleaseGrallocObjectsIMG(
 #define CL_IMG_GENERATE_MIPMAP_EXTENSION_NAME \
     "cl_img_generate_mipmap"
 
+
+#define CL_IMG_GENERATE_MIPMAP_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_uint             cl_mipmap_filter_mode_img;
 
 /* cl_mipmap_filter_mode_img */
@@ -1448,6 +1547,9 @@ clEnqueueGenerateMipmapIMG(
 #define CL_IMG_MEM_PROPERTIES_EXTENSION_NAME \
     "cl_img_mem_properties"
 
+
+#define CL_IMG_MEM_PROPERTIES_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_properties */
 #define CL_MEM_ALLOC_FLAGS_IMG                              0x40D7
 
@@ -1469,6 +1571,9 @@ clEnqueueGenerateMipmapIMG(
 #define CL_KHR_SUBGROUPS_EXTENSION_NAME \
     "cl_khr_subgroups"
 
+
+#define CL_KHR_SUBGROUPS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 #if !defined(CL_VERSION_2_1)
 /* defined in CL.h for OpenCL 2.1 and newer */
 typedef cl_uint             cl_kernel_sub_group_info;
@@ -1516,6 +1621,9 @@ clGetKernelSubGroupInfoKHR(
 #define CL_KHR_MIPMAP_IMAGE_EXTENSION_NAME \
     "cl_khr_mipmap_image"
 
+
+#define CL_KHR_MIPMAP_IMAGE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_sampler_properties */
 #define CL_SAMPLER_MIP_FILTER_MODE_KHR                      0x1155
 #define CL_SAMPLER_LOD_MIN_KHR                              0x1156
@@ -1528,6 +1636,9 @@ clGetKernelSubGroupInfoKHR(
 #define CL_KHR_PRIORITY_HINTS_EXTENSION_NAME \
     "cl_khr_priority_hints"
 
+
+#define CL_KHR_PRIORITY_HINTS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* To be used by clGetEventInfo */
 typedef cl_uint             cl_queue_priority_khr;
 
@@ -1546,6 +1657,9 @@ typedef cl_uint             cl_queue_priority_khr;
 #define CL_KHR_THROTTLE_HINTS_EXTENSION_NAME \
     "cl_khr_throttle_hints"
 
+
+#define CL_KHR_THROTTLE_HINTS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* To be used by clGetEventInfo */
 typedef cl_uint             cl_queue_throttle_khr;
 
@@ -1564,6 +1678,9 @@ typedef cl_uint             cl_queue_throttle_khr;
 #define CL_KHR_SUBGROUP_NAMED_BARRIER_EXTENSION_NAME \
     "cl_khr_subgroup_named_barrier"
 
+
+#define CL_KHR_SUBGROUP_NAMED_BARRIER_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_MAX_NAMED_BARRIER_COUNT_KHR               0x2035
 
@@ -1574,6 +1691,9 @@ typedef cl_uint             cl_queue_throttle_khr;
 #define CL_KHR_EXTENDED_VERSIONING_EXTENSION_NAME \
     "cl_khr_extended_versioning"
 
+
+#define CL_KHR_EXTENDED_VERSIONING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 #define CL_VERSION_MAJOR_BITS_KHR                           10
 #define CL_VERSION_MINOR_BITS_KHR                           10
 #define CL_VERSION_PATCH_BITS_KHR                           12
@@ -1617,6 +1737,9 @@ typedef struct _cl_name_version_khr {
 #define CL_KHR_DEVICE_UUID_EXTENSION_NAME \
     "cl_khr_device_uuid"
 
+
+#define CL_KHR_DEVICE_UUID_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* Size Constants */
 #define CL_UUID_SIZE_KHR                                    16
 #define CL_LUID_SIZE_KHR                                    8
@@ -1635,6 +1758,9 @@ typedef struct _cl_name_version_khr {
 #define CL_KHR_PCI_BUS_INFO_EXTENSION_NAME \
     "cl_khr_pci_bus_info"
 
+
+#define CL_KHR_PCI_BUS_INFO_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef struct _cl_device_pci_bus_info_khr {
     cl_uint pci_domain;
     cl_uint pci_bus;
@@ -1653,6 +1779,9 @@ typedef struct _cl_device_pci_bus_info_khr {
     "cl_khr_suggested_local_work_size"
 
 
+#define CL_KHR_SUGGESTED_LOCAL_WORK_SIZE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
+
 typedef cl_int CL_API_CALL
 clGetKernelSuggestedLocalWorkSizeKHR_t(
     cl_command_queue command_queue,
@@ -1685,6 +1814,9 @@ clGetKernelSuggestedLocalWorkSizeKHR(
 #define CL_KHR_INTEGER_DOT_PRODUCT_EXTENSION_NAME \
     "cl_khr_integer_dot_product"
 
+
+#define CL_KHR_INTEGER_DOT_PRODUCT_EXTENSION_VERSION CL_MAKE_VERSION(2, 0, 0)
+
 typedef cl_bitfield         cl_device_integer_dot_product_capabilities_khr;
 typedef struct _cl_device_integer_dot_product_acceleration_properties_khr {
     cl_bool signed_accelerated;
@@ -1711,6 +1843,9 @@ typedef struct _cl_device_integer_dot_product_acceleration_properties_khr {
 #define CL_KHR_EXTERNAL_MEMORY_EXTENSION_NAME \
     "cl_khr_external_memory"
 
+
+#define CL_KHR_EXTERNAL_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 1)
+
 typedef cl_uint             cl_external_memory_handle_type_khr;
 
 /* cl_platform_info */
@@ -1782,21 +1917,11 @@ clEnqueueReleaseExternalMemObjectsKHR(
 #define CL_KHR_EXTERNAL_MEMORY_DMA_BUF_EXTENSION_NAME \
     "cl_khr_external_memory_dma_buf"
 
-/* cl_external_memory_handle_type_khr */
-#define CL_EXTERNAL_MEMORY_HANDLE_DMA_BUF_KHR               0x2067
 
-/***************************************************************
-* cl_khr_external_memory_dx
-***************************************************************/
-#define cl_khr_external_memory_dx 1
-#define CL_KHR_EXTERNAL_MEMORY_DX_EXTENSION_NAME \
-    "cl_khr_external_memory_dx"
+#define CL_KHR_EXTERNAL_MEMORY_DMA_BUF_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
 
 /* cl_external_memory_handle_type_khr */
-#define CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KHR         0x2063
-#define CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KMT_KHR     0x2064
-#define CL_EXTERNAL_MEMORY_HANDLE_D3D12_HEAP_KHR            0x2065
-#define CL_EXTERNAL_MEMORY_HANDLE_D3D12_RESOURCE_KHR        0x2066
+#define CL_EXTERNAL_MEMORY_HANDLE_DMA_BUF_KHR               0x2067
 
 /***************************************************************
 * cl_khr_external_memory_opaque_fd
@@ -1805,6 +1930,9 @@ clEnqueueReleaseExternalMemObjectsKHR(
 #define CL_KHR_EXTERNAL_MEMORY_OPAQUE_FD_EXTENSION_NAME \
     "cl_khr_external_memory_opaque_fd"
 
+
+#define CL_KHR_EXTERNAL_MEMORY_OPAQUE_FD_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_external_memory_handle_type_khr */
 #define CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR             0x2060
 
@@ -1815,9 +1943,13 @@ clEnqueueReleaseExternalMemObjectsKHR(
 #define CL_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME \
     "cl_khr_external_memory_win32"
 
+
+#define CL_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_VERSION CL_MAKE_VERSION(1, 1, 0)
+
 /* cl_external_memory_handle_type_khr */
 #define CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR          0x2061
 #define CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR      0x2062
+#define CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_NAME_KHR     0x2069
 
 /***************************************************************
 * cl_khr_external_semaphore
@@ -1826,6 +1958,9 @@ clEnqueueReleaseExternalMemObjectsKHR(
 #define CL_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME \
     "cl_khr_external_semaphore"
 
+
+#define CL_KHR_EXTERNAL_SEMAPHORE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 1)
+
 typedef struct _cl_semaphore_khr * cl_semaphore_khr;
 typedef cl_uint             cl_external_semaphore_handle_type_khr;
 
@@ -1870,16 +2005,6 @@ clGetSemaphoreHandleForTypeKHR(
 
 #endif /* !defined(CL_NO_NON_ICD_DISPATCH_EXTENSION_PROTOTYPES) */
 
-/***************************************************************
-* cl_khr_external_semaphore_dx_fence
-***************************************************************/
-#define cl_khr_external_semaphore_dx_fence 1
-#define CL_KHR_EXTERNAL_SEMAPHORE_DX_FENCE_EXTENSION_NAME \
-    "cl_khr_external_semaphore_dx_fence"
-
-/* cl_external_semaphore_handle_type_khr */
-#define CL_SEMAPHORE_HANDLE_D3D12_FENCE_KHR                 0x2059
-
 /***************************************************************
 * cl_khr_external_semaphore_opaque_fd
 ***************************************************************/
@@ -1887,6 +2012,9 @@ clGetSemaphoreHandleForTypeKHR(
 #define CL_KHR_EXTERNAL_SEMAPHORE_OPAQUE_FD_EXTENSION_NAME \
     "cl_khr_external_semaphore_opaque_fd"
 
+
+#define CL_KHR_EXTERNAL_SEMAPHORE_OPAQUE_FD_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_external_semaphore_handle_type_khr */
 #define CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR                   0x2055
 
@@ -1897,6 +2025,9 @@ clGetSemaphoreHandleForTypeKHR(
 #define CL_KHR_EXTERNAL_SEMAPHORE_SYNC_FD_EXTENSION_NAME \
     "cl_khr_external_semaphore_sync_fd"
 
+
+#define CL_KHR_EXTERNAL_SEMAPHORE_SYNC_FD_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_properties       cl_semaphore_reimport_properties_khr;
 
 /* cl_external_semaphore_handle_type_khr */
@@ -1929,9 +2060,13 @@ clReImportSemaphoreSyncFdKHR(
 #define CL_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME \
     "cl_khr_external_semaphore_win32"
 
+
+#define CL_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_VERSION CL_MAKE_VERSION(0, 9, 1)
+
 /* cl_external_semaphore_handle_type_khr */
 #define CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR                0x2056
 #define CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR            0x2057
+#define CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_NAME_KHR           0x2068
 
 /***************************************************************
 * cl_khr_semaphore
@@ -1940,6 +2075,9 @@ clReImportSemaphoreSyncFdKHR(
 #define CL_KHR_SEMAPHORE_EXTENSION_NAME \
     "cl_khr_semaphore"
 
+
+#define CL_KHR_SEMAPHORE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* type cl_semaphore_khr */
 typedef cl_properties       cl_semaphore_properties_khr;
 typedef cl_uint             cl_semaphore_info_khr;
@@ -2087,6 +2225,9 @@ clRetainSemaphoreKHR(
 #define CL_ARM_IMPORT_MEMORY_EXTENSION_NAME \
     "cl_arm_import_memory"
 
+
+#define CL_ARM_IMPORT_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef intptr_t            cl_import_properties_arm;
 
 /* cl_import_properties_arm */
@@ -2133,6 +2274,9 @@ clImportMemoryARM(
 #define CL_ARM_SHARED_VIRTUAL_MEMORY_EXTENSION_NAME \
     "cl_arm_shared_virtual_memory"
 
+
+#define CL_ARM_SHARED_VIRTUAL_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_bitfield         cl_svm_mem_flags_arm;
 typedef cl_uint             cl_kernel_exec_info_arm;
 typedef cl_bitfield         cl_device_svm_capabilities_arm;
@@ -2359,6 +2503,9 @@ clSetKernelExecInfoARM(
 #define CL_ARM_GET_CORE_ID_EXTENSION_NAME \
     "cl_arm_get_core_id"
 
+
+#define CL_ARM_GET_CORE_ID_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_COMPUTE_UNITS_BITFIELD_ARM                0x40BF
 
@@ -2371,6 +2518,9 @@ clSetKernelExecInfoARM(
 #define CL_ARM_JOB_SLOT_SELECTION_EXTENSION_NAME \
     "cl_arm_job_slot_selection"
 
+
+#define CL_ARM_JOB_SLOT_SELECTION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_JOB_SLOTS_ARM                             0x41E0
 
@@ -2384,6 +2534,9 @@ clSetKernelExecInfoARM(
 #define CL_ARM_SCHEDULING_CONTROLS_EXTENSION_NAME \
     "cl_arm_scheduling_controls"
 
+
+#define CL_ARM_SCHEDULING_CONTROLS_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* Types */
 typedef cl_bitfield         cl_device_scheduling_controls_capabilities_arm;
 
@@ -2423,6 +2576,9 @@ typedef cl_bitfield         cl_device_scheduling_controls_capabilities_arm;
 #define CL_ARM_CONTROLLED_KERNEL_TERMINATION_EXTENSION_NAME \
     "cl_arm_controlled_kernel_termination"
 
+
+#define CL_ARM_CONTROLLED_KERNEL_TERMINATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* Types */
 typedef cl_bitfield         cl_device_controlled_termination_capabilities_arm;
 
@@ -2453,6 +2609,9 @@ typedef cl_bitfield         cl_device_controlled_termination_capabilities_arm;
 #define CL_ARM_PROTECTED_MEMORY_ALLOCATION_EXTENSION_NAME \
     "cl_arm_protected_memory_allocation"
 
+
+#define CL_ARM_PROTECTED_MEMORY_ALLOCATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 #define CL_MEM_PROTECTED_ALLOC_ARM                          ((cl_bitfield)1 << 36)
 
 /***************************************************************
@@ -2462,6 +2621,9 @@ typedef cl_bitfield         cl_device_controlled_termination_capabilities_arm;
 #define CL_INTEL_EXEC_BY_LOCAL_THREAD_EXTENSION_NAME \
     "cl_intel_exec_by_local_thread"
 
+
+#define CL_INTEL_EXEC_BY_LOCAL_THREAD_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_command_queue_properties - bitfield */
 #define CL_QUEUE_THREAD_LOCAL_EXEC_ENABLE_INTEL             ((cl_bitfield)1 << 31)
 
@@ -2472,6 +2634,9 @@ typedef cl_bitfield         cl_device_controlled_termination_capabilities_arm;
 #define CL_INTEL_DEVICE_ATTRIBUTE_QUERY_EXTENSION_NAME \
     "cl_intel_device_attribute_query"
 
+
+#define CL_INTEL_DEVICE_ATTRIBUTE_QUERY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_bitfield         cl_device_feature_capabilities_intel;
 
 /* cl_device_feature_capabilities_intel */
@@ -2494,6 +2659,9 @@ typedef cl_bitfield         cl_device_feature_capabilities_intel;
 #define CL_INTEL_DEVICE_PARTITION_BY_NAMES_EXTENSION_NAME \
     "cl_intel_device_partition_by_names"
 
+
+#define CL_INTEL_DEVICE_PARTITION_BY_NAMES_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 #define CL_DEVICE_PARTITION_BY_NAMES_INTEL                  0x4052
 #define CL_PARTITION_BY_NAMES_LIST_END_INTEL                -1
 
@@ -2504,6 +2672,9 @@ typedef cl_bitfield         cl_device_feature_capabilities_intel;
 #define CL_INTEL_ACCELERATOR_EXTENSION_NAME \
     "cl_intel_accelerator"
 
+
+#define CL_INTEL_ACCELERATOR_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef struct _cl_accelerator_intel* cl_accelerator_intel;
 typedef cl_uint             cl_accelerator_type_intel;
 typedef cl_uint             cl_accelerator_info_intel;
@@ -2592,6 +2763,9 @@ clReleaseAcceleratorINTEL(
 #define CL_INTEL_MOTION_ESTIMATION_EXTENSION_NAME \
     "cl_intel_motion_estimation"
 
+
+#define CL_INTEL_MOTION_ESTIMATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef struct _cl_motion_estimation_desc_intel {
     cl_uint mb_block_type;
     cl_uint subpixel_mode;
@@ -2628,6 +2802,9 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_ADVANCED_MOTION_ESTIMATION_EXTENSION_NAME \
     "cl_intel_advanced_motion_estimation"
 
+
+#define CL_INTEL_ADVANCED_MOTION_ESTIMATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_ME_VERSION_INTEL                          0x407E
 
@@ -2684,6 +2861,9 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_SIMULTANEOUS_SHARING_EXTENSION_NAME \
     "cl_intel_simultaneous_sharing"
 
+
+#define CL_INTEL_SIMULTANEOUS_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_SIMULTANEOUS_INTEROPS_INTEL               0x4104
 #define CL_DEVICE_NUM_SIMULTANEOUS_INTEROPS_INTEL           0x4105
@@ -2695,6 +2875,9 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_EGL_IMAGE_YUV_EXTENSION_NAME \
     "cl_intel_egl_image_yuv"
 
+
+#define CL_INTEL_EGL_IMAGE_YUV_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_egl_image_properties_khr */
 #define CL_EGL_YUV_PLANE_INTEL                              0x4107
 
@@ -2705,6 +2888,9 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_PACKED_YUV_EXTENSION_NAME \
     "cl_intel_packed_yuv"
 
+
+#define CL_INTEL_PACKED_YUV_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_channel_order */
 #define CL_YUYV_INTEL                                       0x4076
 #define CL_UYVY_INTEL                                       0x4077
@@ -2718,6 +2904,9 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_REQUIRED_SUBGROUP_SIZE_EXTENSION_NAME \
     "cl_intel_required_subgroup_size"
 
+
+#define CL_INTEL_REQUIRED_SUBGROUP_SIZE_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_SUB_GROUP_SIZES_INTEL                     0x4108
 
@@ -2734,10 +2923,15 @@ typedef struct _cl_motion_estimation_desc_intel {
 #define CL_INTEL_DRIVER_DIAGNOSTICS_EXTENSION_NAME \
     "cl_intel_driver_diagnostics"
 
-typedef cl_uint             cl_diagnostics_verbose_level;
+
+#define CL_INTEL_DRIVER_DIAGNOSTICS_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
+typedef cl_bitfield         cl_diagnostic_verbose_level_intel;
 
 /* cl_context_properties */
 #define CL_CONTEXT_SHOW_DIAGNOSTICS_INTEL                   0x4106
+
+/* cl_diagnostic_verbose_level_intel */
 #define CL_CONTEXT_DIAGNOSTICS_LEVEL_ALL_INTEL              0xff
 #define CL_CONTEXT_DIAGNOSTICS_LEVEL_GOOD_INTEL             (1 << 0)
 #define CL_CONTEXT_DIAGNOSTICS_LEVEL_BAD_INTEL              (1 << 1)
@@ -2750,6 +2944,9 @@ typedef cl_uint             cl_diagnostics_verbose_level;
 #define CL_INTEL_PLANAR_YUV_EXTENSION_NAME \
     "cl_intel_planar_yuv"
 
+
+#define CL_INTEL_PLANAR_YUV_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_channel_order */
 #define CL_NV12_INTEL                                       0x410E
 
@@ -2768,6 +2965,9 @@ typedef cl_uint             cl_diagnostics_verbose_level;
 #define CL_INTEL_DEVICE_SIDE_AVC_MOTION_ESTIMATION_EXTENSION_NAME \
     "cl_intel_device_side_avc_motion_estimation"
 
+
+#define CL_INTEL_DEVICE_SIDE_AVC_MOTION_ESTIMATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_AVC_ME_VERSION_INTEL                      0x410B
 #define CL_DEVICE_AVC_ME_SUPPORTS_TEXTURE_SAMPLER_USE_INTEL 0x410C
@@ -2925,6 +3125,9 @@ typedef cl_uint             cl_diagnostics_verbose_level;
 #define CL_INTEL_UNIFIED_SHARED_MEMORY_EXTENSION_NAME \
     "cl_intel_unified_shared_memory"
 
+
+#define CL_INTEL_UNIFIED_SHARED_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_bitfield         cl_device_unified_shared_memory_capabilities_intel;
 typedef cl_properties       cl_mem_properties_intel;
 typedef cl_bitfield         cl_mem_alloc_flags_intel;
@@ -3246,6 +3449,9 @@ clEnqueueMemsetINTEL(
 #define CL_INTEL_MEM_ALLOC_BUFFER_LOCATION_EXTENSION_NAME \
     "cl_intel_mem_alloc_buffer_location"
 
+
+#define CL_INTEL_MEM_ALLOC_BUFFER_LOCATION_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_properties_intel */
 #define CL_MEM_ALLOC_BUFFER_LOCATION_INTEL                  0x419E
 
@@ -3259,6 +3465,9 @@ clEnqueueMemsetINTEL(
 #define CL_INTEL_CREATE_BUFFER_WITH_PROPERTIES_EXTENSION_NAME \
     "cl_intel_create_buffer_with_properties"
 
+
+#define CL_INTEL_CREATE_BUFFER_WITH_PROPERTIES_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* type cl_mem_properties_intel */
 
 
@@ -3294,6 +3503,9 @@ clCreateBufferWithPropertiesINTEL(
 #define CL_INTEL_PROGRAM_SCOPE_HOST_PIPE_EXTENSION_NAME \
     "cl_intel_program_scope_host_pipe"
 
+
+#define CL_INTEL_PROGRAM_SCOPE_HOST_PIPE_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* clGetEventInfo response when param_name is CL_EVENT_COMMAND_TYPE */
 #define CL_COMMAND_READ_HOST_PIPE_INTEL                     0x4214
 #define CL_COMMAND_WRITE_HOST_PIPE_INTEL                    0x4215
@@ -3368,6 +3580,9 @@ clEnqueueWriteHostPipeINTEL(
 #define CL_INTEL_MEM_CHANNEL_PROPERTY_EXTENSION_NAME \
     "cl_intel_mem_channel_property"
 
+
+#define CL_INTEL_MEM_CHANNEL_PROPERTY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_properties_intel */
 #define CL_MEM_CHANNEL_INTEL                                0x4213
 
@@ -3378,6 +3593,9 @@ clEnqueueWriteHostPipeINTEL(
 #define CL_INTEL_MEM_FORCE_HOST_MEMORY_EXTENSION_NAME \
     "cl_intel_mem_force_host_memory"
 
+
+#define CL_INTEL_MEM_FORCE_HOST_MEMORY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_flags */
 #define CL_MEM_FORCE_HOST_MEMORY_INTEL                      (1 << 20)
 
@@ -3388,6 +3606,9 @@ clEnqueueWriteHostPipeINTEL(
 #define CL_INTEL_COMMAND_QUEUE_FAMILIES_EXTENSION_NAME \
     "cl_intel_command_queue_families"
 
+
+#define CL_INTEL_COMMAND_QUEUE_FAMILIES_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_bitfield         cl_command_queue_capabilities_intel;
 
 #define CL_QUEUE_FAMILY_MAX_NAME_SIZE_INTEL                 64
@@ -3432,6 +3653,9 @@ typedef struct _cl_queue_family_properties_intel {
 #define CL_INTEL_QUEUE_NO_SYNC_OPERATIONS_EXTENSION_NAME \
     "cl_intel_queue_no_sync_operations"
 
+
+#define CL_INTEL_QUEUE_NO_SYNC_OPERATIONS_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_command_queue_properties */
 #define CL_QUEUE_NO_SYNC_OPERATIONS_INTEL                   (1 << 29)
 
@@ -3442,6 +3666,9 @@ typedef struct _cl_queue_family_properties_intel {
 #define CL_INTEL_SHARING_FORMAT_QUERY_EXTENSION_NAME \
     "cl_intel_sharing_format_query"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /***************************************************************
 * cl_ext_image_requirements_info
 ***************************************************************/
@@ -3451,6 +3678,9 @@ typedef struct _cl_queue_family_properties_intel {
 #define CL_EXT_IMAGE_REQUIREMENTS_INFO_EXTENSION_NAME \
     "cl_ext_image_requirements_info"
 
+
+#define CL_EXT_IMAGE_REQUIREMENTS_INFO_EXTENSION_VERSION CL_MAKE_VERSION(0, 5, 0)
+
 /* Types */
 typedef cl_uint             cl_image_requirements_info_ext;
 
@@ -3507,6 +3737,9 @@ clGetImageRequirementsInfoEXT(
 #define CL_EXT_IMAGE_FROM_BUFFER_EXTENSION_NAME \
     "cl_ext_image_from_buffer"
 
+
+#define CL_EXT_IMAGE_FROM_BUFFER_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_image_requirements_info_ext */
 #define CL_IMAGE_REQUIREMENTS_SLICE_PITCH_ALIGNMENT_EXT     0x1291
 
@@ -3519,6 +3752,9 @@ clGetImageRequirementsInfoEXT(
 #define CL_LOADER_INFO_EXTENSION_NAME \
     "cl_loader_info"
 
+
+#define CL_LOADER_INFO_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_uint             cl_icdl_info;
 
 /* cl_icdl_info */
@@ -3556,6 +3792,9 @@ clGetICDLoaderInfoOCLICD(
 #define CL_KHR_DEPTH_IMAGES_EXTENSION_NAME \
     "cl_khr_depth_images"
 
+
+#define CL_KHR_DEPTH_IMAGES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 #if !defined(CL_VERSION_2_0)
 /* cl_channel_order - defined in CL.h for OpenCL 2.0 and newer */
 #define CL_DEPTH                                            0x10BD
@@ -3569,6 +3808,9 @@ clGetICDLoaderInfoOCLICD(
 #define CL_EXT_FLOAT_ATOMICS_EXTENSION_NAME \
     "cl_ext_float_atomics"
 
+
+#define CL_EXT_FLOAT_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_bitfield         cl_device_fp_atomic_capabilities_ext;
 
 /* cl_device_fp_atomic_capabilities_ext */
@@ -3591,6 +3833,9 @@ typedef cl_bitfield         cl_device_fp_atomic_capabilities_ext;
 #define CL_INTEL_CREATE_MEM_OBJECT_PROPERTIES_EXTENSION_NAME \
     "cl_intel_create_mem_object_properties"
 
+
+#define CL_INTEL_CREATE_MEM_OBJECT_PROPERTIES_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* cl_mem_properties */
 #define CL_MEM_LOCALLY_UNCACHED_RESOURCE_INTEL              0x4218
 #define CL_MEM_DEVICE_ID_INTEL                              0x4219
@@ -3603,6 +3848,9 @@ typedef cl_bitfield         cl_device_fp_atomic_capabilities_ext;
     "cl_pocl_content_size"
 
 
+#define CL_POCL_CONTENT_SIZE_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
+
 typedef cl_int CL_API_CALL
 clSetContentSizeBufferPoCL_t(
     cl_mem buffer,
@@ -3627,6 +3875,9 @@ clSetContentSizeBufferPoCL(
 #define CL_EXT_IMAGE_RAW10_RAW12_EXTENSION_NAME \
     "cl_ext_image_raw10_raw12"
 
+
+#define CL_EXT_IMAGE_RAW10_RAW12_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_channel_type */
 #define CL_UNSIGNED_INT_RAW10_EXT                           0x10E3
 #define CL_UNSIGNED_INT_RAW12_EXT                           0x10E4
@@ -3638,6 +3889,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_3D_IMAGE_WRITES_EXTENSION_NAME \
     "cl_khr_3d_image_writes"
 
+
+#define CL_KHR_3D_IMAGE_WRITES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_async_work_group_copy_fence
 ***************************************************************/
@@ -3645,6 +3899,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_ASYNC_WORK_GROUP_COPY_FENCE_EXTENSION_NAME \
     "cl_khr_async_work_group_copy_fence"
 
+
+#define CL_KHR_ASYNC_WORK_GROUP_COPY_FENCE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_byte_addressable_store
 ***************************************************************/
@@ -3652,6 +3909,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_BYTE_ADDRESSABLE_STORE_EXTENSION_NAME \
     "cl_khr_byte_addressable_store"
 
+
+#define CL_KHR_BYTE_ADDRESSABLE_STORE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_device_enqueue_local_arg_types
 ***************************************************************/
@@ -3659,6 +3919,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_DEVICE_ENQUEUE_LOCAL_ARG_TYPES_EXTENSION_NAME \
     "cl_khr_device_enqueue_local_arg_types"
 
+
+#define CL_KHR_DEVICE_ENQUEUE_LOCAL_ARG_TYPES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_expect_assume
 ***************************************************************/
@@ -3666,6 +3929,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_EXPECT_ASSUME_EXTENSION_NAME \
     "cl_khr_expect_assume"
 
+
+#define CL_KHR_EXPECT_ASSUME_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_extended_async_copies
 ***************************************************************/
@@ -3673,6 +3939,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_EXTENDED_ASYNC_COPIES_EXTENSION_NAME \
     "cl_khr_extended_async_copies"
 
+
+#define CL_KHR_EXTENDED_ASYNC_COPIES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_extended_bit_ops
 ***************************************************************/
@@ -3680,6 +3949,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_EXTENDED_BIT_OPS_EXTENSION_NAME \
     "cl_khr_extended_bit_ops"
 
+
+#define CL_KHR_EXTENDED_BIT_OPS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_global_int32_base_atomics
 ***************************************************************/
@@ -3687,6 +3959,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_GLOBAL_INT32_BASE_ATOMICS_EXTENSION_NAME \
     "cl_khr_global_int32_base_atomics"
 
+
+#define CL_KHR_GLOBAL_INT32_BASE_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_global_int32_extended_atomics
 ***************************************************************/
@@ -3694,6 +3969,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_GLOBAL_INT32_EXTENDED_ATOMICS_EXTENSION_NAME \
     "cl_khr_global_int32_extended_atomics"
 
+
+#define CL_KHR_GLOBAL_INT32_EXTENDED_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_int64_base_atomics
 ***************************************************************/
@@ -3701,6 +3979,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_INT64_BASE_ATOMICS_EXTENSION_NAME \
     "cl_khr_int64_base_atomics"
 
+
+#define CL_KHR_INT64_BASE_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_int64_extended_atomics
 ***************************************************************/
@@ -3708,6 +3989,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_INT64_EXTENDED_ATOMICS_EXTENSION_NAME \
     "cl_khr_int64_extended_atomics"
 
+
+#define CL_KHR_INT64_EXTENDED_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_kernel_clock
 ***************************************************************/
@@ -3715,6 +3999,9 @@ clSetContentSizeBufferPoCL(
 #define CL_KHR_KERNEL_CLOCK_EXTENSION_NAME \
     "cl_khr_kernel_clock"
 
+
+#define CL_KHR_KERNEL_CLOCK_EXTENSION_VERSION CL_MAKE_VERSION(0, 9, 0)
+
 /* cl_device_info */
 #define CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR             0x1076
 
@@ -3732,6 +4019,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_LOCAL_INT32_BASE_ATOMICS_EXTENSION_NAME \
     "cl_khr_local_int32_base_atomics"
 
+
+#define CL_KHR_LOCAL_INT32_BASE_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_local_int32_extended_atomics
 ***************************************************************/
@@ -3739,6 +4029,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_LOCAL_INT32_EXTENDED_ATOMICS_EXTENSION_NAME \
     "cl_khr_local_int32_extended_atomics"
 
+
+#define CL_KHR_LOCAL_INT32_EXTENDED_ATOMICS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_mipmap_image_writes
 ***************************************************************/
@@ -3746,6 +4039,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_MIPMAP_IMAGE_WRITES_EXTENSION_NAME \
     "cl_khr_mipmap_image_writes"
 
+
+#define CL_KHR_MIPMAP_IMAGE_WRITES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_select_fprounding_mode
 ***************************************************************/
@@ -3753,6 +4049,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SELECT_FPROUNDING_MODE_EXTENSION_NAME \
     "cl_khr_select_fprounding_mode"
 
+
+#define CL_KHR_SELECT_FPROUNDING_MODE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_spirv_extended_debug_info
 ***************************************************************/
@@ -3760,6 +4059,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SPIRV_EXTENDED_DEBUG_INFO_EXTENSION_NAME \
     "cl_khr_spirv_extended_debug_info"
 
+
+#define CL_KHR_SPIRV_EXTENDED_DEBUG_INFO_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_spirv_linkonce_odr
 ***************************************************************/
@@ -3767,6 +4069,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SPIRV_LINKONCE_ODR_EXTENSION_NAME \
     "cl_khr_spirv_linkonce_odr"
 
+
+#define CL_KHR_SPIRV_LINKONCE_ODR_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_spirv_no_integer_wrap_decoration
 ***************************************************************/
@@ -3774,6 +4079,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SPIRV_NO_INTEGER_WRAP_DECORATION_EXTENSION_NAME \
     "cl_khr_spirv_no_integer_wrap_decoration"
 
+
+#define CL_KHR_SPIRV_NO_INTEGER_WRAP_DECORATION_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_srgb_image_writes
 ***************************************************************/
@@ -3781,6 +4089,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SRGB_IMAGE_WRITES_EXTENSION_NAME \
     "cl_khr_srgb_image_writes"
 
+
+#define CL_KHR_SRGB_IMAGE_WRITES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_ballot
 ***************************************************************/
@@ -3788,6 +4099,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_BALLOT_EXTENSION_NAME \
     "cl_khr_subgroup_ballot"
 
+
+#define CL_KHR_SUBGROUP_BALLOT_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_clustered_reduce
 ***************************************************************/
@@ -3795,6 +4109,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_CLUSTERED_REDUCE_EXTENSION_NAME \
     "cl_khr_subgroup_clustered_reduce"
 
+
+#define CL_KHR_SUBGROUP_CLUSTERED_REDUCE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_extended_types
 ***************************************************************/
@@ -3802,6 +4119,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_EXTENDED_TYPES_EXTENSION_NAME \
     "cl_khr_subgroup_extended_types"
 
+
+#define CL_KHR_SUBGROUP_EXTENDED_TYPES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_non_uniform_arithmetic
 ***************************************************************/
@@ -3809,6 +4129,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_NON_UNIFORM_ARITHMETIC_EXTENSION_NAME \
     "cl_khr_subgroup_non_uniform_arithmetic"
 
+
+#define CL_KHR_SUBGROUP_NON_UNIFORM_ARITHMETIC_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_non_uniform_vote
 ***************************************************************/
@@ -3816,6 +4139,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_NON_UNIFORM_VOTE_EXTENSION_NAME \
     "cl_khr_subgroup_non_uniform_vote"
 
+
+#define CL_KHR_SUBGROUP_NON_UNIFORM_VOTE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_rotate
 ***************************************************************/
@@ -3823,6 +4149,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_ROTATE_EXTENSION_NAME \
     "cl_khr_subgroup_rotate"
 
+
+#define CL_KHR_SUBGROUP_ROTATE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_shuffle
 ***************************************************************/
@@ -3830,6 +4159,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_SHUFFLE_EXTENSION_NAME \
     "cl_khr_subgroup_shuffle"
 
+
+#define CL_KHR_SUBGROUP_SHUFFLE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_subgroup_shuffle_relative
 ***************************************************************/
@@ -3837,6 +4169,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_SUBGROUP_SHUFFLE_RELATIVE_EXTENSION_NAME \
     "cl_khr_subgroup_shuffle_relative"
 
+
+#define CL_KHR_SUBGROUP_SHUFFLE_RELATIVE_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /***************************************************************
 * cl_khr_work_group_uniform_arithmetic
 ***************************************************************/
@@ -3844,6 +4179,22 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_KHR_WORK_GROUP_UNIFORM_ARITHMETIC_EXTENSION_NAME \
     "cl_khr_work_group_uniform_arithmetic"
 
+
+#define CL_KHR_WORK_GROUP_UNIFORM_ARITHMETIC_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
+/***************************************************************
+* cl_ext_image_unorm_int_2_101010
+***************************************************************/
+#define cl_ext_image_unorm_int_2_101010 1
+#define CL_EXT_IMAGE_UNORM_INT_2_101010_EXTENSION_NAME \
+    "cl_ext_image_unorm_int_2_101010"
+
+
+#define CL_EXT_IMAGE_UNORM_INT_2_101010_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
+/* cl_channel_type */
+#define CL_UNORM_INT_2_101010_EXT                           0x10E5
+
 /***************************************************************
 * cl_img_cancel_command
 ***************************************************************/
@@ -3851,6 +4202,9 @@ typedef cl_bitfield         cl_device_kernel_clock_capabilities_khr;
 #define CL_IMG_CANCEL_COMMAND_EXTENSION_NAME \
     "cl_img_cancel_command"
 
+
+#define CL_IMG_CANCEL_COMMAND_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* Error codes */
 #define CL_CANCELLED_IMG                                    -1126
 
diff --git a/include/CL/cl_gl.h b/include/CL/cl_gl.h
index f5b1e37bbe9ef69ce9326904b6dc6fe429873cd9..552560f71c4e4f107c546131306c20bb0479ed85 100644
--- a/include/CL/cl_gl.h
+++ b/include/CL/cl_gl.h
@@ -51,6 +51,13 @@ extern "C" {
 #define CL_KHR_GL_SHARING_EXTENSION_NAME \
     "cl_khr_gl_sharing"
 
+
+#define CL_KHR_GL_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
+typedef int                 cl_GLint;
+typedef unsigned int        cl_GLenum;
+typedef unsigned int        cl_GLuint;
+
 typedef cl_uint             cl_gl_context_info;
 
 /* Error codes */
@@ -313,6 +320,9 @@ clCreateFromGLTexture3D(
 #define CL_KHR_GL_EVENT_EXTENSION_NAME \
     "cl_khr_gl_event"
 
+
+#define CL_KHR_GL_EVENT_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef struct __GLsync *   cl_GLsync;
 
 /* cl_command_type */
@@ -345,6 +355,9 @@ clCreateEventFromGLsyncKHR(
 #define CL_KHR_GL_DEPTH_IMAGES_EXTENSION_NAME \
     "cl_khr_gl_depth_images"
 
+
+#define CL_KHR_GL_DEPTH_IMAGES_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_channel_order */
 #define CL_DEPTH_STENCIL                                    0x10BE
 
@@ -358,6 +371,9 @@ clCreateEventFromGLsyncKHR(
 #define CL_KHR_GL_MSAA_SHARING_EXTENSION_NAME \
     "cl_khr_gl_msaa_sharing"
 
+
+#define CL_KHR_GL_MSAA_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 /* cl_gl_texture_info */
 #define CL_GL_NUM_SAMPLES                                   0x2012
 
@@ -368,6 +384,9 @@ clCreateEventFromGLsyncKHR(
 #define CL_INTEL_SHARING_FORMAT_QUERY_GL_EXTENSION_NAME \
     "cl_intel_sharing_format_query_gl"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_GL_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* when cl_khr_gl_sharing is supported */
 
 typedef cl_int CL_API_CALL
diff --git a/include/CL/cl_layer.h b/include/CL/cl_layer.h
index a43b89783dbc5e57488070f62d71bb685d590d00..245f7b53295a285c25d566a218aecbb0a224ad01 100644
--- a/include/CL/cl_layer.h
+++ b/include/CL/cl_layer.h
@@ -53,6 +53,9 @@ extern "C" {
 #define CL_LOADER_LAYERS_EXTENSION_NAME \
     "cl_loader_layers"
 
+
+#define CL_LOADER_LAYERS_EXTENSION_VERSION CL_MAKE_VERSION(1, 0, 0)
+
 typedef cl_uint             cl_layer_info;
 typedef cl_uint             cl_layer_api_version;
 
diff --git a/include/CL/cl_platform.h b/include/CL/cl_platform.h
index 2d8907697a399e61b316c5c580f58decb004f8cb..5f92d6faad4e68b66c75514320e245fe6b74ebd6 100644
--- a/include/CL/cl_platform.h
+++ b/include/CL/cl_platform.h
@@ -361,11 +361,6 @@ typedef double          cl_double;
 
 #include <stddef.h>
 
-/* Mirror types to GL types. Mirror types allow us to avoid deciding which 87s to load based on whether we are using GL or GLES here. */
-typedef unsigned int cl_GLuint;
-typedef int          cl_GLint;
-typedef unsigned int cl_GLenum;
-
 /*
  * Vector types
  *
diff --git a/include/CL/cl_va_api_media_sharing_intel.h b/include/CL/cl_va_api_media_sharing_intel.h
index 93f5d8bbcfd818eec7e544c0b239b9466fd14a7e..9fb8863f24d9f3ad179bd839af70be1e084c749c 100644
--- a/include/CL/cl_va_api_media_sharing_intel.h
+++ b/include/CL/cl_va_api_media_sharing_intel.h
@@ -53,6 +53,9 @@ extern "C" {
 #define CL_INTEL_SHARING_FORMAT_QUERY_VA_API_EXTENSION_NAME \
     "cl_intel_sharing_format_query_va_api"
 
+
+#define CL_INTEL_SHARING_FORMAT_QUERY_VA_API_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 /* when cl_intel_va_api_media_sharing is supported */
 
 typedef cl_int CL_API_CALL
@@ -89,6 +92,9 @@ clGetSupportedVA_APIMediaSurfaceFormatsINTEL(
 #define CL_INTEL_VA_API_MEDIA_SHARING_EXTENSION_NAME \
     "cl_intel_va_api_media_sharing"
 
+
+#define CL_INTEL_VA_API_MEDIA_SHARING_EXTENSION_VERSION CL_MAKE_VERSION(0, 0, 0)
+
 typedef cl_uint             cl_va_api_device_source_intel;
 typedef cl_uint             cl_va_api_device_set_intel;
 
diff --git a/include/CL/opencl.h b/include/CL/opencl.h
index bc0cde34d3f7328cc75950bbecb723c4f21b11d1..84641d7571040074fb94d0ca42ab2770dde8854a 100644
--- a/include/CL/opencl.h
+++ b/include/CL/opencl.h
@@ -22,11 +22,11 @@ extern "C" {
 #endif
 
 #include <CL/cl.h>
+#include <CL/cl_gl.h>
 #include <CL/cl_exp_tensor.h>
 #include <CL/cl_exp_defined_builtin_kernels.h>
 #include <CL/cl_ext.h>
 #include <CL/cl_ext_pocl.h>
-#include <CL/cl_gl.h>
 
 #ifdef __cplusplus
 }
diff --git a/include/hpp/CL/opencl.hpp b/include/hpp/CL/opencl.hpp
index d8d227e1b2d6c86343651ca54dfebe4c133424b6..970e67b1e2f7f70ff7983b800159936005c2a595 100644
--- a/include/hpp/CL/opencl.hpp
+++ b/include/hpp/CL/opencl.hpp
@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2008-2023 The Khronos Group Inc.
+// Copyright (c) 2008-2024 The Khronos Group Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -336,7 +336,7 @@
         // Traditional cl_mem allocations
 
         std::vector<int> output(numElements, 0xdeadbeef);
-        cl::Buffer outputBuffer(begin(output), end(output), false);
+        cl::Buffer outputBuffer(output.begin(), output.end(), false);
         cl::Pipe aPipe(sizeof(cl_int), numElements / 2);
 
         // Default command queue, also passed in as a parameter
@@ -373,7 +373,7 @@
             error
             );
 
-        cl::copy(outputBuffer, begin(output), end(output));
+        cl::copy(outputBuffer, output.begin(), output.end());
 
         cl::Device d = cl::Device::getDefault();
 
@@ -525,12 +525,6 @@
 #include <CL/opencl.h>
 #endif // !__APPLE__
 
-#if (__cplusplus >= 201103L || _MSVC_LANG >= 201103L )
-#define CL_HPP_NOEXCEPT_ noexcept
-#else
-#define CL_HPP_NOEXCEPT_
-#endif
-
 #if __cplusplus >= 201703L
 # define CL_HPP_DEFINE_STATIC_MEMBER_ inline
 #elif defined(_MSC_VER)
@@ -725,6 +719,7 @@ namespace cl {
 #endif
 
     class Memory;
+    class Platform;
     class Program;
     class Device;
     class Context;
@@ -764,13 +759,11 @@ namespace cl {
         Error(cl_int err, const char * errStr = nullptr) : err_(err), errStr_(errStr)
         {}
 
-        ~Error() throw() {}
-
         /*! \brief Get error string associated with exception
          *
          * \return A memory pointer to the error message string.
          */
-        virtual const char * what() const throw ()
+        const char * what() const noexcept override
         {
             if (errStr_ == nullptr) {
                 return "empty";
@@ -904,6 +897,10 @@ static inline cl_int errHandler (cl_int err, const char * errStr = nullptr)
 #define __ENQUEUE_COPY_IMAGE_TO_BUFFER_ERR  CL_HPP_ERR_STR_(clEnqueueCopyImageToBuffer)
 #define __ENQUEUE_COPY_BUFFER_TO_IMAGE_ERR  CL_HPP_ERR_STR_(clEnqueueCopyBufferToImage)
 #define __ENQUEUE_MAP_BUFFER_ERR            CL_HPP_ERR_STR_(clEnqueueMapBuffer)
+#define __ENQUEUE_MAP_SVM_ERR               CL_HPP_ERR_STR_(clEnqueueSVMMap)
+#define __ENQUEUE_FILL_SVM_ERR              CL_HPP_ERR_STR_(clEnqueueSVMMemFill)
+#define __ENQUEUE_COPY_SVM_ERR              CL_HPP_ERR_STR_(clEnqueueSVMMemcpy)
+#define __ENQUEUE_UNMAP_SVM_ERR             CL_HPP_ERR_STR_(clEnqueueSVMUnmap)
 #define __ENQUEUE_MAP_IMAGE_ERR             CL_HPP_ERR_STR_(clEnqueueMapImage)
 #define __ENQUEUE_UNMAP_MEM_OBJECT_ERR      CL_HPP_ERR_STR_(clEnqueueUnMapMemObject)
 #define __ENQUEUE_NDRANGE_KERNEL_ERR        CL_HPP_ERR_STR_(clEnqueueNDRangeKernel)
@@ -951,6 +948,11 @@ static inline cl_int errHandler (cl_int err, const char * errStr = nullptr)
 #define __RETAIN_SEMAPHORE_KHR_ERR                  CL_HPP_ERR_STR_(clRetainSemaphoreKHR)
 #define __RELEASE_SEMAPHORE_KHR_ERR                 CL_HPP_ERR_STR_(clReleaseSemaphoreKHR)
 #endif
+
+#ifdef cl_khr_external_semaphore
+#define __GET_SEMAPHORE_HANDLE_FOR_TYPE_KHR_ERR         CL_HPP_ERR_STR_(clGetSemaphoreHandleForTypeKHR)
+#endif // cl_khr_external_semaphore
+
 #if defined(cl_khr_command_buffer)
 #define __CREATE_COMMAND_BUFFER_KHR_ERR             CL_HPP_ERR_STR_(clCreateCommandBufferKHR)
 #define __GET_COMMAND_BUFFER_INFO_KHR_ERR           CL_HPP_ERR_STR_(clGetCommandBufferInfoKHR)
@@ -971,6 +973,10 @@ static inline cl_int errHandler (cl_int err, const char * errStr = nullptr)
 #define __RELEASE_COMMAND_BUFFER_KHR_ERR            CL_HPP_ERR_STR_(clReleaseCommandBufferKHR)
 #endif // cl_khr_command_buffer
 
+#if defined(cl_ext_image_requirements_info)
+#define __GET_IMAGE_REQUIREMENT_INFO_EXT_ERR            CL_HPP_ERR_STR_(clGetImageRequirementsInfoEXT)
+#endif //cl_ext_image_requirements_info
+
 /**
  * CL 1.2 version that uses device fission.
  */
@@ -1042,6 +1048,11 @@ CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clEnqueueSignalSemaphoresKHR       pfn_clEnqueu
 CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clGetSemaphoreInfoKHR              pfn_clGetSemaphoreInfoKHR               = nullptr;
 #endif // cl_khr_semaphore
 
+#ifdef cl_khr_external_semaphore
+CL_HPP_CREATE_CL_EXT_FCN_PTR_ALIAS_(clGetSemaphoreHandleForTypeKHR);
+CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clGetSemaphoreHandleForTypeKHR     pfn_clGetSemaphoreHandleForTypeKHR      = nullptr;
+#endif // cl_khr_external_semaphore
+
 #if defined(cl_khr_command_buffer)
 CL_HPP_CREATE_CL_EXT_FCN_PTR_ALIAS_(clCreateCommandBufferKHR);
 CL_HPP_CREATE_CL_EXT_FCN_PTR_ALIAS_(clFinalizeCommandBufferKHR);
@@ -1084,6 +1095,17 @@ CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clUpdateMutableCommandsKHR pfn_clUpdateMutableC
 CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clGetMutableCommandInfoKHR pfn_clGetMutableCommandInfoKHR           = nullptr;
 #endif /* cl_khr_command_buffer_mutable_dispatch */
 
+#if defined(cl_ext_image_requirements_info)
+CL_HPP_CREATE_CL_EXT_FCN_PTR_ALIAS_(clGetImageRequirementsInfoEXT);
+CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clGetImageRequirementsInfoEXT pfn_clGetImageRequirementsInfoEXT  = nullptr;
+#endif
+
+#if defined(cl_ext_device_fission)
+CL_HPP_CREATE_CL_EXT_FCN_PTR_ALIAS_(clCreateSubDevicesEXT);
+CL_HPP_DEFINE_STATIC_MEMBER_ PFN_clCreateSubDevicesEXT
+    pfn_clCreateSubDevicesEXT = nullptr;
+#endif
+
 namespace detail {
 
 // Generic getInfoHelper. The final parameter is used to guide overload
@@ -1121,7 +1143,6 @@ inline cl_int getInfoHelper(Func f, cl_uint name, vector<vector<unsigned char>>*
         }
     }
 
-
     return CL_SUCCESS;
 }
 
@@ -1205,7 +1226,7 @@ inline cl_int getInfoHelper(Func f, cl_uint name, string* param, long)
             return err;
         }
         if (param) {
-            param->assign(begin(value), prev(end(value)));
+            param->assign(value.begin(), value.end() - 1);
         }
     }
     else if (param) {
@@ -1322,7 +1343,7 @@ inline cl_int getInfoHelper(Func f, cl_uint name, T* param, int, typename T::cl_
     F(cl_device_info, CL_DEVICE_AVAILABLE, cl_bool) \
     F(cl_device_info, CL_DEVICE_COMPILER_AVAILABLE, cl_bool) \
     F(cl_device_info, CL_DEVICE_EXECUTION_CAPABILITIES, cl_device_exec_capabilities) \
-    F(cl_device_info, CL_DEVICE_PLATFORM, cl_platform_id) \
+    F(cl_device_info, CL_DEVICE_PLATFORM, cl::Platform) \
     F(cl_device_info, CL_DEVICE_NAME, string) \
     F(cl_device_info, CL_DEVICE_VENDOR, string) \
     F(cl_device_info, CL_DRIVER_VERSION, string) \
@@ -1519,8 +1540,13 @@ inline cl_int getInfoHelper(Func f, cl_uint name, T* param, int, typename T::cl_
 #define CL_HPP_PARAM_NAME_CL_KHR_EXTENDED_VERSIONING_KHRONLY_(F) \
     F(cl_device_info, CL_DEVICE_OPENCL_C_NUMERIC_VERSION_KHR, cl_version_khr)
 
+// Note: the query for CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR is handled specially!
 #define CL_HPP_PARAM_NAME_CL_KHR_SEMAPHORE_(F) \
+    F(cl_semaphore_info_khr, CL_SEMAPHORE_CONTEXT_KHR, cl::Context) \
+    F(cl_semaphore_info_khr, CL_SEMAPHORE_REFERENCE_COUNT_KHR, cl_uint) \
     F(cl_semaphore_info_khr, CL_SEMAPHORE_PROPERTIES_KHR, cl::vector<cl_semaphore_properties_khr>) \
+    F(cl_semaphore_info_khr, CL_SEMAPHORE_TYPE_KHR, cl_semaphore_type_khr) \
+    F(cl_semaphore_info_khr, CL_SEMAPHORE_PAYLOAD_KHR, cl_semaphore_payload_khr) \
     F(cl_platform_info, CL_PLATFORM_SEMAPHORE_TYPES_KHR,  cl::vector<cl_semaphore_type_khr>) \
     F(cl_device_info, CL_DEVICE_SEMAPHORE_TYPES_KHR,      cl::vector<cl_semaphore_type_khr>) \
 
@@ -1528,6 +1554,23 @@ inline cl_int getInfoHelper(Func f, cl_uint name, T* param, int, typename T::cl_
     F(cl_device_info, CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, cl::vector<cl::ExternalMemoryType>) \
     F(cl_platform_info, CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, cl::vector<cl::ExternalMemoryType>)
 
+#define CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_(F) \
+    F(cl_platform_info, CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR,  cl::vector<cl_external_semaphore_handle_type_khr>) \
+    F(cl_platform_info, CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR,  cl::vector<cl_external_semaphore_handle_type_khr>) \
+    F(cl_device_info, CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR,      cl::vector<cl_external_semaphore_handle_type_khr>) \
+    F(cl_device_info, CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR,      cl::vector<cl_external_semaphore_handle_type_khr>) \
+    F(cl_semaphore_info_khr, CL_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR,      cl::vector<cl_external_semaphore_handle_type_khr>) \
+
+#define CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_OPAQUE_FD_EXT(F) \
+    F(cl_external_semaphore_handle_type_khr, CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR, int) \
+
+#define CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_SYNC_FD_EXT(F) \
+    F(cl_external_semaphore_handle_type_khr, CL_SEMAPHORE_HANDLE_SYNC_FD_KHR, int) \
+
+#define CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_WIN32_EXT(F) \
+    F(cl_external_semaphore_handle_type_khr, CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR, void*) \
+    F(cl_external_semaphore_handle_type_khr, CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR, void*) \
+
 #define CL_HPP_PARAM_NAME_INFO_3_0_(F) \
     F(cl_platform_info, CL_PLATFORM_NUMERIC_VERSION, cl_version) \
     F(cl_platform_info, CL_PLATFORM_EXTENSIONS_WITH_VERSION, cl::vector<cl_name_version>) \
@@ -1553,6 +1596,18 @@ inline cl_int getInfoHelper(Func f, cl_uint name, T* param, int, typename T::cl_
     F(cl_pipe_info, CL_PIPE_PROPERTIES, cl::vector<cl_pipe_properties>) \
     F(cl_sampler_info, CL_SAMPLER_PROPERTIES, cl::vector<cl_sampler_properties>) \
 
+#define CL_HPP_PARAM_NAME_CL_IMAGE_REQUIREMENTS_EXT(F) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_ROW_PITCH_ALIGNMENT_EXT, size_type) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_BASE_ADDRESS_ALIGNMENT_EXT, size_type) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_SIZE_EXT, size_type) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_MAX_WIDTH_EXT, cl_uint) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_MAX_HEIGHT_EXT, cl_uint) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_MAX_DEPTH_EXT, cl_uint) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_MAX_ARRAY_SIZE_EXT, cl_uint) \
+
+#define CL_HPP_PARAM_NAME_CL_IMAGE_REQUIREMENTS_SLICE_PITCH_ALIGNMENT_EXT(F) \
+    F(cl_image_requirements_info_ext, CL_IMAGE_REQUIREMENTS_SLICE_PITCH_ALIGNMENT_EXT, size_type) \
+
 template <typename enum_type, cl_int Name>
 struct param_traits {};
 
@@ -1629,12 +1684,29 @@ CL_HPP_PARAM_NAME_CL_KHR_EXTENDED_VERSIONING_KHRONLY_(CL_HPP_DECLARE_PARAM_TRAIT
 
 #if defined(cl_khr_semaphore)
 CL_HPP_PARAM_NAME_CL_KHR_SEMAPHORE_(CL_HPP_DECLARE_PARAM_TRAITS_)
-#endif // cl_khr_semaphore
+#if defined(CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR)
+CL_HPP_DECLARE_PARAM_TRAITS_(cl_semaphore_info_khr, CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR, cl::vector<cl::Device>)
+#endif // defined(CL_SEMAPHORE_DEVICE_HANDLE_LIST_KHR)
+#endif // defined(cl_khr_semaphore)
 
 #ifdef cl_khr_external_memory
 CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_MEMORY_(CL_HPP_DECLARE_PARAM_TRAITS_)
 #endif // cl_khr_external_memory
 
+#if defined(cl_khr_external_semaphore)
+CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_khr_external_semaphore
+
+#if defined(cl_khr_external_semaphore_opaque_fd)
+CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_OPAQUE_FD_EXT(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_khr_external_semaphore_opaque_fd
+#if defined(cl_khr_external_semaphore_sync_fd)
+CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_SYNC_FD_EXT(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_khr_external_semaphore_sync_fd
+#if defined(cl_khr_external_semaphore_win32)
+CL_HPP_PARAM_NAME_CL_KHR_EXTERNAL_SEMAPHORE_WIN32_EXT(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_khr_external_semaphore_win32
+
 #if defined(cl_khr_device_uuid)
 using uuid_array = array<cl_uchar, CL_UUID_SIZE_KHR>;
 using luid_array = array<cl_uchar, CL_LUID_SIZE_KHR>;
@@ -1667,6 +1739,14 @@ CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_INTEGER_DOT_PRODUCT_ACCEL
 #endif // defined(CL_DEVICE_INTEGER_DOT_PRODUCT_ACCELERATION_PROPERTIES_8BIT_KHR)
 #endif // defined(cl_khr_integer_dot_product)
 
+#if defined(cl_ext_image_requirements_info)
+CL_HPP_PARAM_NAME_CL_IMAGE_REQUIREMENTS_EXT(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_ext_image_requirements_info
+
+#if defined(cl_ext_image_from_buffer)
+CL_HPP_PARAM_NAME_CL_IMAGE_REQUIREMENTS_SLICE_PITCH_ALIGNMENT_EXT(CL_HPP_DECLARE_PARAM_TRAITS_)
+#endif // cl_ext_image_from_buffer
+
 #ifdef CL_PLATFORM_ICD_SUFFIX_KHR
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_platform_info, CL_PLATFORM_ICD_SUFFIX_KHR, string)
 #endif
@@ -1763,7 +1843,7 @@ CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_INTEGRATED_MEMORY_NV, cl_
 
 #if defined(cl_khr_command_buffer)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR, cl_device_command_buffer_capabilities_khr)
-CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR, cl_command_buffer_properties_khr)
+CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR, cl_command_queue_properties)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_command_buffer_info_khr, CL_COMMAND_BUFFER_QUEUES_KHR, cl::vector<CommandQueue>)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_command_buffer_info_khr, CL_COMMAND_BUFFER_NUM_QUEUES_KHR, cl_uint)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_command_buffer_info_khr, CL_COMMAND_BUFFER_REFERENCE_COUNT_KHR, cl_uint)
@@ -1775,7 +1855,12 @@ CL_HPP_DECLARE_PARAM_TRAITS_(cl_command_buffer_info_khr, CL_COMMAND_BUFFER_PROPE
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_COMMAND_COMMAND_QUEUE_KHR, CommandQueue)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_COMMAND_COMMAND_BUFFER_KHR, CommandBufferKhr)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR, cl_command_type)
+
+#if CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 2)
+CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_COMMAND_PROPERTIES_ARRAY_KHR, cl::vector<cl_command_properties_khr>)
+#else
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_PROPERTIES_ARRAY_KHR, cl::vector<cl_ndrange_kernel_command_properties_khr>)
+#endif
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_KERNEL_KHR, cl_kernel)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_DIMENSIONS_KHR, cl_uint)
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_GLOBAL_WORK_OFFSET_KHR, cl::vector<size_type>)
@@ -1783,6 +1868,10 @@ CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_GL
 CL_HPP_DECLARE_PARAM_TRAITS_(cl_mutable_command_info_khr, CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR, cl::vector<size_type>)
 #endif /* cl_khr_command_buffer_mutable_dispatch */
 
+#if defined(cl_khr_kernel_clock)
+CL_HPP_DECLARE_PARAM_TRAITS_(cl_device_info, CL_DEVICE_KERNEL_CLOCK_CAPABILITIES_KHR, cl_device_kernel_clock_capabilities_khr)
+#endif /* cl_khr_kernel_clock */
+
 // Convenience functions
 
 template <typename Func, typename T>
@@ -2008,7 +2097,8 @@ struct ReferenceHandler<cl_mutable_command_khr>
 #endif // cl_khr_command_buffer
 
 
-#if CL_HPP_TARGET_OPENCL_VERSION >= 120 && CL_HPP_MINIMUM_OPENCL_VERSION < 120
+#if (CL_HPP_TARGET_OPENCL_VERSION >= 120 && CL_HPP_MINIMUM_OPENCL_VERSION < 120) || \
+    (CL_HPP_TARGET_OPENCL_VERSION >= 200 && CL_HPP_MINIMUM_OPENCL_VERSION < 200)
 // Extracts version number with major in the upper 16 bits, minor in the lower 16
 static cl_uint getVersion(const vector<char> &versionInfo)
 {
@@ -2058,7 +2148,7 @@ static cl_uint getContextPlatformVersion(cl_context context)
     clGetContextInfo(context, CL_CONTEXT_DEVICES, size, devices.data(), nullptr);
     return getDevicePlatformVersion(devices[0]);
 }
-#endif // CL_HPP_TARGET_OPENCL_VERSION >= 120 && CL_HPP_MINIMUM_OPENCL_VERSION < 120
+#endif // CL_HPP_TARGET_OPENCL_VERSION && CL_HPP_MINIMUM_OPENCL_VERSION
 
 template <typename T>
 class Wrapper
@@ -2090,7 +2180,7 @@ public:
         detail::errHandler(retain(), __RETAIN_ERR);
     }
 
-    Wrapper(Wrapper<cl_type>&& rhs) CL_HPP_NOEXCEPT_
+    Wrapper(Wrapper<cl_type>&& rhs) noexcept
     {
         object_ = rhs.object_;
         rhs.object_ = nullptr;
@@ -2167,18 +2257,16 @@ protected:
     static bool isReferenceCountable(cl_device_id device)
     {
         bool retVal = false;
-#if CL_HPP_TARGET_OPENCL_VERSION >= 120
-#if CL_HPP_MINIMUM_OPENCL_VERSION < 120
+#if CL_HPP_TARGET_OPENCL_VERSION >= 120 && CL_HPP_MINIMUM_OPENCL_VERSION < 120
         if (device != nullptr) {
             int version = getDevicePlatformVersion(device);
             if(version > ((1 << 16) + 1)) {
                 retVal = true;
             }
         }
-#else // CL_HPP_MINIMUM_OPENCL_VERSION < 120
+#elif CL_HPP_TARGET_OPENCL_VERSION >= 120
         retVal = true;
-#endif // CL_HPP_MINIMUM_OPENCL_VERSION < 120
-#endif // CL_HPP_TARGET_OPENCL_VERSION >= 120
+#endif // CL_HPP_TARGET_OPENCL_VERSION
         (void)device;
         return retVal;
     }
@@ -2211,7 +2299,7 @@ public:
         detail::errHandler(retain(), __RETAIN_ERR);
     }
 
-    Wrapper(Wrapper<cl_type>&& rhs) CL_HPP_NOEXCEPT_
+    Wrapper(Wrapper<cl_type>&& rhs) noexcept
     {
         object_ = rhs.object_;
         referenceCountable_ = rhs.referenceCountable_;
@@ -2498,91 +2586,14 @@ public:
 
 #if CL_HPP_TARGET_OPENCL_VERSION >= 120
     //! \brief Wrapper for clCreateSubDevices().
-    cl_int createSubDevices(
-        const cl_device_partition_property * properties,
-        vector<Device>* devices)
-    {
-        cl_uint n = 0;
-        cl_int err = clCreateSubDevices(object_, properties, 0, nullptr, &n);
-        if (err != CL_SUCCESS) {
-            return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
-        }
-
-        vector<cl_device_id> ids(n);
-        err = clCreateSubDevices(object_, properties, n, ids.data(), nullptr);
-        if (err != CL_SUCCESS) {
-            return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
-        }
-
-        // Cannot trivially assign because we need to capture intermediates 
-        // with safe construction
-        if (devices) {
-            devices->resize(ids.size());
-
-            // Assign to param, constructing with retain behaviour
-            // to correctly capture each underlying CL object
-            for (size_type i = 0; i < ids.size(); i++) {
-                // We do not need to retain because this device is being created 
-                // by the runtime
-                (*devices)[i] = Device(ids[i], false);
-            }
-        }
-
-        return CL_SUCCESS;
-    }
-#endif
+    cl_int createSubDevices(const cl_device_partition_property* properties,
+                            vector<Device>* devices);
+#endif // defined (CL_HPP_TARGET_OPENCL_VERSION >= 120)
 
 #if defined(cl_ext_device_fission)
     //! \brief Wrapper for clCreateSubDevices().
-    cl_int createSubDevices(
-        const cl_device_partition_property_ext * properties,
-        vector<Device>* devices)
-    {
-        typedef CL_API_ENTRY cl_int 
-            ( CL_API_CALL * PFN_clCreateSubDevicesEXT)(
-                cl_device_id /*in_device*/,
-                const cl_device_partition_property_ext * /* properties */,
-                cl_uint /*num_entries*/,
-                cl_device_id * /*out_devices*/,
-                cl_uint * /*num_devices*/ ) CL_API_SUFFIX__VERSION_1_1;
-
-        static PFN_clCreateSubDevicesEXT pfn_clCreateSubDevicesEXT = nullptr;
-#if CL_HPP_TARGET_OPENCL_VERSION >= 120
-        cl::Device device(object_);
-        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>();
-        CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clCreateSubDevicesEXT);
-#endif
-#if CL_HPP_MINIMUM_OPENCL_VERSION < 120
-        CL_HPP_INIT_CL_EXT_FCN_PTR_(clCreateSubDevicesEXT);
-#endif
-
-        cl_uint n = 0;
-        cl_int err = pfn_clCreateSubDevicesEXT(object_, properties, 0, nullptr, &n);
-        if (err != CL_SUCCESS) {
-            return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
-        }
-
-        vector<cl_device_id> ids(n);
-        err = pfn_clCreateSubDevicesEXT(object_, properties, n, ids.data(), nullptr);
-        if (err != CL_SUCCESS) {
-            return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
-        }
-        // Cannot trivially assign because we need to capture intermediates 
-        // with safe construction
-        if (devices) {
-            devices->resize(ids.size());
-
-            // Assign to param, constructing with retain behaviour
-            // to correctly capture each underlying CL object
-            for (size_type i = 0; i < ids.size(); i++) {
-                // We do not need to retain because this device is being created 
-                // by the runtime
-                (*devices)[i] = Device(ids[i], false);
-            }
-        }
-
-        return CL_SUCCESS;
-    }
+    cl_int createSubDevices(const cl_device_partition_property_ext* properties,
+                            vector<Device>* devices);
 #endif // defined(cl_ext_device_fission)
 };
 
@@ -3000,6 +3011,93 @@ public:
 #endif // CL_HPP_TARGET_OPENCL_VERSION >= 120
 }; // class Platform
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 120
+   //! \brief Wrapper for clCreateSubDevices().
+inline cl_int Device::createSubDevices(const cl_device_partition_property* properties,
+                         vector<Device>* devices)
+{
+    cl_uint n = 0;
+    cl_int err = clCreateSubDevices(object_, properties, 0, nullptr, &n);
+    if (err != CL_SUCCESS)
+    {
+        return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
+    }
+
+    vector<cl_device_id> ids(n);
+    err = clCreateSubDevices(object_, properties, n, ids.data(), nullptr);
+    if (err != CL_SUCCESS)
+    {
+        return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
+    }
+
+    // Cannot trivially assign because we need to capture intermediates
+    // with safe construction
+    if (devices)
+    {
+        devices->resize(ids.size());
+
+        // Assign to param, constructing with retain behaviour
+        // to correctly capture each underlying CL object
+        for (size_type i = 0; i < ids.size(); i++)
+        {
+            // We do not need to retain because this device is being created
+            // by the runtime
+            (*devices)[i] = Device(ids[i], false);
+        }
+    }
+
+    return CL_SUCCESS;
+}
+#endif // defined (CL_HPP_TARGET_OPENCL_VERSION >= 120)
+
+#if defined(cl_ext_device_fission)
+   //! \brief Wrapper for clCreateSubDevices().
+inline cl_int Device::createSubDevices(const cl_device_partition_property_ext* properties,
+                        vector<Device>* devices)
+{
+#if CL_HPP_TARGET_OPENCL_VERSION >= 120
+    cl::Device device(object_);
+    cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>()();
+    CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clCreateSubDevicesEXT);
+#endif
+#if CL_HPP_MINIMUM_OPENCL_VERSION < 120
+    CL_HPP_INIT_CL_EXT_FCN_PTR_(clCreateSubDevicesEXT);
+#endif
+
+    cl_uint n = 0;
+    cl_int err = pfn_clCreateSubDevicesEXT(object_, properties, 0, nullptr, &n);
+    if (err != CL_SUCCESS)
+    {
+        return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
+    }
+
+    vector<cl_device_id> ids(n);
+    err =
+        pfn_clCreateSubDevicesEXT(object_, properties, n, ids.data(), nullptr);
+    if (err != CL_SUCCESS)
+    {
+        return detail::errHandler(err, __CREATE_SUB_DEVICES_ERR);
+    }
+    // Cannot trivially assign because we need to capture intermediates
+    // with safe construction
+    if (devices)
+    {
+        devices->resize(ids.size());
+
+        // Assign to param, constructing with retain behaviour
+        // to correctly capture each underlying CL object
+        for (size_type i = 0; i < ids.size(); i++)
+        {
+            // We do not need to retain because this device is being created
+            // by the runtime
+            (*devices)[i] = Device(ids[i], false);
+        }
+    }
+
+    return CL_SUCCESS;
+}
+#endif // defined(cl_ext_device_fission)
+
 CL_HPP_DEFINE_STATIC_MEMBER_ std::once_flag Platform::default_initialized_;
 CL_HPP_DEFINE_STATIC_MEMBER_ Platform Platform::default_;
 CL_HPP_DEFINE_STATIC_MEMBER_ cl_int Platform::default_error_ = CL_SUCCESS;
@@ -3022,6 +3120,25 @@ UnloadCompiler()
 }
 #endif // #if defined(CL_USE_DEPRECATED_OPENCL_1_1_APIS)
 
+
+#if defined(cl_ext_image_requirements_info)
+enum ImageRequirementsInfoExt : cl_image_requirements_info_ext
+{
+    RowPitchAlign = CL_IMAGE_REQUIREMENTS_ROW_PITCH_ALIGNMENT_EXT,
+    BaseAddAlign = CL_IMAGE_REQUIREMENTS_BASE_ADDRESS_ALIGNMENT_EXT,
+    Size = CL_IMAGE_REQUIREMENTS_SIZE_EXT,
+    MaxWidth = CL_IMAGE_REQUIREMENTS_MAX_WIDTH_EXT,
+    MaxHeight = CL_IMAGE_REQUIREMENTS_MAX_HEIGHT_EXT,
+    MaxDepth = CL_IMAGE_REQUIREMENTS_MAX_DEPTH_EXT,
+    MaxArraySize = CL_IMAGE_REQUIREMENTS_MAX_ARRAY_SIZE_EXT,
+#if defined(cl_ext_image_from_buffer)
+    SlicePitchAlign = CL_IMAGE_REQUIREMENTS_SLICE_PITCH_ALIGNMENT_EXT,
+#endif
+};
+
+#endif // cl_ext_image_requirements_info
+
+
 /*! \brief Class interface for cl_context.
  *
  *  \note Copies of these objects are shallow, meaning that the copy will refer
@@ -3084,6 +3201,51 @@ private:
     static void makeDefaultProvided(const Context &c) {
         default_ = c;
     }
+
+#if defined(cl_ext_image_requirements_info)
+    struct ImageRequirementsInfo {
+
+        ImageRequirementsInfo(cl_mem_flags f, const cl_mem_properties* mem_properties, const ImageFormat* format, const cl_image_desc* desc)
+        {
+            flags = f;
+            properties = mem_properties;
+            image_format = format;
+            image_desc = desc;
+        }
+
+        cl_mem_flags flags = 0;
+        const cl_mem_properties* properties;
+        const ImageFormat* image_format;
+        const cl_image_desc* image_desc;
+    };
+
+    static cl_int getImageRequirementsInfoExtHelper(const Context &context,
+        const ImageRequirementsInfo &info,
+        cl_image_requirements_info_ext param_name,
+        size_type param_value_size,
+        void* param_value,
+        size_type* param_value_size_ret)
+    {
+
+#if CL_HPP_TARGET_OPENCL_VERSION >= 120
+        Device device = context.getInfo<CL_CONTEXT_DEVICES>().at(0);
+        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>()();
+        CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clGetImageRequirementsInfoEXT);
+#else
+        CL_HPP_INIT_CL_EXT_FCN_PTR_(clGetImageRequirementsInfoEXT);
+#endif
+
+        if (pfn_clGetImageRequirementsInfoEXT == nullptr) {
+            return detail::errHandler(CL_INVALID_OPERATION, __GET_IMAGE_REQUIREMENT_INFO_EXT_ERR);
+        }
+
+        return detail::errHandler(
+            pfn_clGetImageRequirementsInfoEXT(context(), info.properties,
+                info.flags, info.image_format, info.image_desc, param_name,
+                param_value_size, param_value, param_value_size_ret),
+            __GET_IMAGE_REQUIREMENT_INFO_EXT_ERR);
+    }
+#endif // cl_ext_image_requirements_info
     
 public:
 #ifdef CL_HPP_UNIT_TEST_ENABLE
@@ -3363,7 +3525,7 @@ public:
                 return detail::errHandler(err, __GET_SUPPORTED_IMAGE_FORMATS_ERR);
             }
 
-            formats->assign(begin(value), end(value));
+            formats->assign(value.begin(), value.end());
         }
         else {
             // If no values are being returned, ensure an empty vector comes back
@@ -3373,6 +3535,41 @@ public:
         return CL_SUCCESS;
     }
 
+#if defined(cl_ext_image_requirements_info)
+    template <typename T>
+    cl_int getImageRequirementsInfoExt(cl_image_requirements_info_ext name,
+        T* param,
+        cl_mem_flags flags = 0,
+        const cl_mem_properties* properties = nullptr,
+        const ImageFormat* image_format = nullptr,
+        const cl_image_desc* image_desc = nullptr) const
+    {
+        ImageRequirementsInfo imageInfo = {flags, properties, image_format, image_desc};
+
+        return detail::errHandler(
+            detail::getInfo(
+                Context::getImageRequirementsInfoExtHelper, *this, imageInfo, name, param),
+                __GET_IMAGE_REQUIREMENT_INFO_EXT_ERR);
+    }
+
+    template <cl_image_requirements_info_ext type> typename
+    detail::param_traits<detail::cl_image_requirements_info_ext, type>::param_type
+        getImageRequirementsInfoExt(cl_mem_flags flags = 0,
+            const cl_mem_properties* properties = nullptr,
+            const ImageFormat* image_format = nullptr,
+            const cl_image_desc* image_desc = nullptr,
+            cl_int* err = nullptr) const
+    {
+        typename detail::param_traits<
+        detail::cl_image_requirements_info_ext, type>::param_type param;
+        cl_int result = getImageRequirementsInfoExt(type, &param, flags, properties, image_format, image_desc);
+        if (err != nullptr) {
+            *err = result;
+        }
+        return param;
+    }
+#endif // cl_ext_image_requirements_info
+
 #if CL_HPP_TARGET_OPENCL_VERSION >= 300
     /*! \brief  Registers a destructor callback function with a context.
      *
@@ -3870,12 +4067,12 @@ public:
     {
     }
 
-    pointer address(reference r) CL_HPP_NOEXCEPT_
+    pointer address(reference r) noexcept
     {
         return std::addressof(r);
     }
 
-    const_pointer address(const_reference r) CL_HPP_NOEXCEPT_
+    const_pointer address(const_reference r) noexcept
     {
         return std::addressof(r);
     }
@@ -3888,7 +4085,8 @@ public:
      */
     pointer allocate(
         size_type size,
-        typename cl::SVMAllocator<void, SVMTrait>::const_pointer = 0)
+        typename cl::SVMAllocator<void, SVMTrait>::const_pointer = 0,
+        bool map = true)
     {
         // Allocate memory with default alignment matching the size of the type
         void* voidPointer =
@@ -3907,11 +4105,15 @@ public:
 #endif // #if defined(CL_HPP_ENABLE_EXCEPTIONS)
 
         // If allocation was coarse-grained then map it
-        if (!(SVMTrait::getSVMMemFlags() & CL_MEM_SVM_FINE_GRAIN_BUFFER)) {
+        if (map && !(SVMTrait::getSVMMemFlags() & CL_MEM_SVM_FINE_GRAIN_BUFFER)) {
             cl_int err = enqueueMapSVM(retValue, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, size*sizeof(T));
             if (err != CL_SUCCESS) {
+                clSVMFree(context_(), retValue);
+                retValue = nullptr;
+#if defined(CL_HPP_ENABLE_EXCEPTIONS)
                 std::bad_alloc excep;
                 throw excep;
+#endif
             }
         }
 
@@ -3928,7 +4130,7 @@ public:
      * Return the maximum possible allocation size.
      * This is the minimum of the maximum sizes of all devices in the context.
      */
-    size_type max_size() const CL_HPP_NOEXCEPT_
+    size_type max_size() const noexcept
     {
         size_type maxSize = std::numeric_limits<size_type>::max() / sizeof(T);
 
@@ -4026,10 +4228,18 @@ cl::pointer<T, detail::Deleter<Alloc>> allocate_pointer(const Alloc &alloc_, Arg
 
     T* tmp = std::allocator_traits<Alloc>::allocate(alloc, copies);
     if (!tmp) {
+#if defined(CL_HPP_ENABLE_EXCEPTIONS)
         std::bad_alloc excep;
         throw excep;
+#else
+        return nullptr;
+#endif
     }
-    try {
+
+#if defined(CL_HPP_ENABLE_EXCEPTIONS)
+    try
+#endif
+    {
         std::allocator_traits<Alloc>::construct(
             alloc,
             std::addressof(*tmp),
@@ -4037,11 +4247,13 @@ cl::pointer<T, detail::Deleter<Alloc>> allocate_pointer(const Alloc &alloc_, Arg
 
         return cl::pointer<T, detail::Deleter<Alloc>>(tmp, detail::Deleter<Alloc>{alloc, copies});
     }
+#if defined(CL_HPP_ENABLE_EXCEPTIONS)
     catch (std::bad_alloc&)
     {
         std::allocator_traits<Alloc>::deallocate(alloc, tmp, copies);
         throw;
     }
+#endif
 }
 
 template< class T, class SVMTrait, class... Args >
@@ -4633,6 +4845,42 @@ public:
     //! \brief Default constructor - initializes to nullptr.
     Image1D() { }
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image1D with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param host_ptr Storage to be used if the CL_MEM_USE_HOST_PTR flag was
+     *                  specified. Note alignment & exclusivity requirements.
+     */
+    Image1D(const Context &context, const vector<cl_mem_properties> &properties,
+            cl_mem_flags flags, ImageFormat format, size_type width,
+            void *host_ptr = nullptr, cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE1D;
+      desc.image_width = width;
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, host_ptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, host_ptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     /*! \brief Constructor from cl_mem - takes ownership.
      *
      * \param retainObject will cause the constructor to retain its cl object.
@@ -4693,6 +4941,43 @@ public:
 
     Image1DBuffer() { }
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image1DBuffer with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param buffer Refer to a valid buffer or image memory object.
+     */
+    Image1DBuffer(const Context &context,
+                  const vector<cl_mem_properties> &properties,
+                  cl_mem_flags flags, ImageFormat format, size_type width,
+                  const Buffer &buffer, cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+      desc.image_width = width;
+      desc.buffer = buffer();
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, nullptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, nullptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     /*! \brief Constructor from cl_mem - takes ownership.
      *
      * \param retainObject will cause the constructor to retain its cl object.
@@ -4708,9 +4993,6 @@ public:
         Image::operator=(rhs);
         return *this;
     }
-
-
-
 };
 
 /*! \class Image1DArray
@@ -4752,7 +5034,47 @@ public:
     }
 
     Image1DArray() { }
-  
+
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image1DArray with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param host_ptr Storage to be used if the CL_MEM_USE_HOST_PTR flag was
+     *                  specified. Note alignment & exclusivity requirements.
+     */
+    Image1DArray(const Context &context,
+                 const vector<cl_mem_properties> &properties,
+                 cl_mem_flags flags, ImageFormat format, size_type arraySize,
+                 size_type width, size_type rowPitch = 0,
+                 void *host_ptr = nullptr, cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
+      desc.image_width = width;
+      desc.image_array_size = arraySize;
+      desc.image_row_pitch = rowPitch;
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, host_ptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, host_ptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     /*! \brief Constructor from cl_mem - takes ownership.
      *
      * \param retainObject will cause the constructor to retain its cl object.
@@ -4955,6 +5277,83 @@ public:
     }
 #endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 200
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image2D with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param host_ptr Storage to be used if the CL_MEM_USE_HOST_PTR flag was
+     *                  specified. Note alignment & exclusivity requirements.
+     */
+    Image2D(const Context &context, const vector<cl_mem_properties> &properties,
+            cl_mem_flags flags, ImageFormat format, size_type width,
+            size_type height, size_type row_pitch = 0, void *host_ptr = nullptr,
+            cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+      desc.image_width = width;
+      desc.image_height = height;
+      desc.image_row_pitch = row_pitch;
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, host_ptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, host_ptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+
+    /*! \brief Constructs a Image2D with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param buffer Refer to a valid buffer or image memory object.
+     */
+    Image2D(const Context &context, const vector<cl_mem_properties> &properties,
+            cl_mem_flags flags, ImageFormat format, const Buffer &buffer,
+            size_type width, size_type height, size_type row_pitch = 0,
+            cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+      desc.image_width = width;
+      desc.image_height = height;
+      desc.image_row_pitch = row_pitch;
+      desc.buffer = buffer();
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, nullptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, nullptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     //! \brief Default constructor - initializes to nullptr.
     Image2D() { }
 
@@ -4977,10 +5376,6 @@ public:
         Image::operator=(rhs);
         return *this;
     }
-
-
-
-
 };
 
 
@@ -5097,6 +5492,49 @@ public:
         }
     }
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image2DArray with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param host_ptr Storage to be used if the CL_MEM_USE_HOST_PTR flag was
+     *                  specified. Note alignment & exclusivity requirements.
+     */
+    Image2DArray(const Context &context,
+                 const vector<cl_mem_properties> &properties,
+                 cl_mem_flags flags, ImageFormat format, size_type arraySize,
+                 size_type width, size_type height, size_type rowPitch = 0,
+                 size_type slicePitch = 0, void *host_ptr = nullptr,
+                 cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY;
+      desc.image_width = width;
+      desc.image_height = height;
+      desc.image_array_size = arraySize;
+      desc.image_row_pitch = rowPitch;
+      desc.image_slice_pitch = slicePitch;
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, host_ptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, host_ptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     Image2DArray() { }
     
     /*! \brief Constructor from cl_mem - takes ownership.
@@ -5197,6 +5635,48 @@ public:
 #endif // CL_HPP_MINIMUM_OPENCL_VERSION < 120
     }
 
+#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+    /*! \brief Constructs a Image3D with specified properties.
+     *
+     *  Wraps clCreateImageWithProperties().
+     *
+     *  \param properties Optional list of properties for the image object and
+     *                    their corresponding values. The non-empty list must
+     *                    end with 0.
+     *  \param host_ptr Storage to be used if the CL_MEM_USE_HOST_PTR flag was
+     *                  specified. Note alignment & exclusivity requirements.
+     */
+    Image3D(const Context &context, const vector<cl_mem_properties> &properties,
+            cl_mem_flags flags, ImageFormat format, size_type width,
+            size_type height, size_type depth, size_type row_pitch = 0,
+            size_type slice_pitch = 0, void *host_ptr = nullptr,
+            cl_int *err = nullptr) {
+      cl_int error;
+
+      cl_image_desc desc = {};
+      desc.image_type = CL_MEM_OBJECT_IMAGE3D;
+      desc.image_width = width;
+      desc.image_height = height;
+      desc.image_depth = depth;
+      desc.image_row_pitch = row_pitch;
+      desc.image_slice_pitch = slice_pitch;
+
+      if (properties.empty()) {
+        object_ = ::clCreateImageWithProperties(
+            context(), nullptr, flags, &format, &desc, host_ptr, &error);
+      } else {
+        object_ =
+            ::clCreateImageWithProperties(context(), properties.data(), flags,
+                                          &format, &desc, host_ptr, &error);
+      }
+
+      detail::errHandler(error, __CREATE_IMAGE_ERR);
+      if (err != nullptr) {
+        *err = error;
+      }
+    }
+#endif //#if CL_HPP_TARGET_OPENCL_VERSION >= 300
+
     //! \brief Default constructor - initializes to nullptr.
     Image3D() : Image() { }
 
@@ -5718,6 +6198,7 @@ Local(size_type size)
 class Kernel : public detail::Wrapper<cl_kernel>
 {
 public:
+    inline Kernel(const Program& program, const string& name, cl_int* err = nullptr);
     inline Kernel(const Program& program, const char* name, cl_int* err = nullptr);
 
     //! \brief Default constructor - initializes to nullptr.
@@ -6196,7 +6677,6 @@ public:
         }
     }
 
-
 #if defined(CL_HPP_USE_IL_KHR) || CL_HPP_TARGET_OPENCL_VERSION >= 210
     /**
      * Program constructor to allow construction of program from SPIR-V or another IL.
@@ -6346,7 +6826,6 @@ public:
             return;
         }
 
-
         vector<size_type> lengths(numDevices);
         vector<const unsigned char*> images(numDevices);
 #if !defined(CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY)
@@ -6360,7 +6839,7 @@ public:
             lengths[i] = binaries[(int)i].second;
         }
 #endif // #if !defined(CL_HPP_ENABLE_PROGRAM_CONSTRUCTION_FROM_ARRAY_COMPATIBILITY)
-        
+
         vector<cl_device_id> deviceIDs(numDevices);
         for( size_type deviceIndex = 0; deviceIndex < numDevices; ++deviceIndex ) {
             deviceIDs[deviceIndex] = (devices[deviceIndex])();
@@ -6369,7 +6848,7 @@ public:
         if(binaryStatus) {
             binaryStatus->resize(numDevices);
         }
-        
+
         object_ = ::clCreateProgramWithBinary(
             context(), (cl_uint) devices.size(),
             deviceIDs.data(),
@@ -6436,6 +6915,14 @@ public:
         return *this;
     }
 
+    cl_int build(
+        const vector<Device>& devices,
+        const string& options,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return build(devices, options.c_str(), notifyFptr, data);
+    }
 
     cl_int build(
         const vector<Device>& devices,
@@ -6445,7 +6932,7 @@ public:
     {
         size_type numDevices = devices.size();
         vector<cl_device_id> deviceIDs(numDevices);
-        
+
         for( size_type deviceIndex = 0; deviceIndex < numDevices; ++deviceIndex ) {
             deviceIDs[deviceIndex] = (devices[deviceIndex])();
         }
@@ -6462,6 +6949,15 @@ public:
         return detail::buildErrHandler(buildError, __BUILD_PROGRAM_ERR, getBuildInfo<CL_PROGRAM_BUILD_LOG>());
     }
 
+    cl_int build(
+        const Device& device,
+        const string& options,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return build(device, options.c_str(), notifyFptr, data);
+    }
+
     cl_int build(
         const Device& device,
         const char* options = nullptr,
@@ -6483,6 +6979,14 @@ public:
         return detail::buildErrHandler(buildError, __BUILD_PROGRAM_ERR, buildLog);
     }
 
+    cl_int build(
+        const string& options,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return build(options.c_str(), notifyFptr, data);
+    }
+
     cl_int build(
         const char* options = nullptr,
         void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
@@ -6500,6 +7004,14 @@ public:
     }
 
 #if CL_HPP_TARGET_OPENCL_VERSION >= 120
+    cl_int compile(
+        const string& options,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return compile(options.c_str(), notifyFptr, data);
+    }
+
     cl_int compile(
         const char* options = nullptr,
         void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
@@ -6517,6 +7029,84 @@ public:
             data);
         return detail::buildErrHandler(error, __COMPILE_PROGRAM_ERR, getBuildInfo<CL_PROGRAM_BUILD_LOG>());
     }
+
+    cl_int compile(
+        const string& options,
+        const vector<Program>& inputHeaders,
+        const vector<string>& headerIncludeNames,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return compile(options.c_str(), inputHeaders, headerIncludeNames, notifyFptr, data);
+    }
+
+    cl_int compile(
+        const char* options,
+        const vector<Program>& inputHeaders,
+        const vector<string>& headerIncludeNames,
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        static_assert(sizeof(cl::Program) == sizeof(cl_program),
+            "Size of cl::Program must be equal to size of cl_program");
+        vector<const char*> headerIncludeNamesCStr;
+        for(const string& name: headerIncludeNames) {
+            headerIncludeNamesCStr.push_back(name.c_str());
+        }
+        cl_int error = ::clCompileProgram(
+            object_,
+            0,
+            nullptr,
+            options,
+            static_cast<cl_uint>(inputHeaders.size()),
+            reinterpret_cast<const cl_program*>(inputHeaders.data()),
+            reinterpret_cast<const char**>(headerIncludeNamesCStr.data()),
+            notifyFptr,
+            data);
+        return detail::buildErrHandler(error, __COMPILE_PROGRAM_ERR, getBuildInfo<CL_PROGRAM_BUILD_LOG>());
+    }
+
+    cl_int compile(
+        const string& options,
+        const vector<Device>& deviceList,
+        const vector<Program>& inputHeaders = vector<Program>(),
+        const vector<string>& headerIncludeNames = vector<string>(),
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        return compile(options.c_str(), deviceList, inputHeaders, headerIncludeNames, notifyFptr, data);
+    }
+
+    cl_int compile(
+        const char* options,
+        const vector<Device>& deviceList,
+        const vector<Program>& inputHeaders = vector<Program>(),
+        const vector<string>& headerIncludeNames = vector<string>(),
+        void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+        void* data = nullptr) const
+    {
+        static_assert(sizeof(cl::Program) == sizeof(cl_program),
+            "Size of cl::Program must be equal to size of cl_program");
+        vector<const char*> headerIncludeNamesCStr;
+        for(const string& name: headerIncludeNames) {
+            headerIncludeNamesCStr.push_back(name.c_str());
+        }
+        vector<cl_device_id> deviceIDList;
+        for(const Device& device: deviceList) {
+            deviceIDList.push_back(device());
+        }
+        cl_int error = ::clCompileProgram(
+            object_,
+            static_cast<cl_uint>(deviceList.size()),
+            reinterpret_cast<const cl_device_id*>(deviceIDList.data()),
+            options,
+            static_cast<cl_uint>(inputHeaders.size()),
+            reinterpret_cast<const cl_program*>(inputHeaders.data()),
+            reinterpret_cast<const char**>(headerIncludeNamesCStr.data()),
+            notifyFptr,
+            data);
+        return detail::buildErrHandler(error, __COMPILE_PROGRAM_ERR, getBuildInfo<CL_PROGRAM_BUILD_LOG>());
+    }
 #endif // CL_HPP_TARGET_OPENCL_VERSION >= 120
 
     template <typename T>
@@ -6732,6 +7322,17 @@ inline Program linkProgram(
     return Program(prog);
 }
 
+inline Program linkProgram(
+    const Program& input1,
+    const Program& input2,
+    const string& options,
+    void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+    void* data = nullptr,
+    cl_int* err = nullptr)
+{
+    return linkProgram(input1, input2, options.c_str(), notifyFptr, data, err);
+}
+
 inline Program linkProgram(
     const vector<Program>& inputPrograms,
     const char* options = nullptr,
@@ -6770,6 +7371,16 @@ inline Program linkProgram(
 
     return Program(prog);
 }
+
+inline Program linkProgram(
+    const vector<Program>& inputPrograms,
+    const string& options,
+    void (CL_CALLBACK * notifyFptr)(cl_program, void *) = nullptr,
+    void* data = nullptr,
+    cl_int* err = nullptr)
+{
+    return linkProgram(inputPrograms, options.c_str(), notifyFptr, data, err);
+}
 #endif // CL_HPP_TARGET_OPENCL_VERSION >= 120
 
 // Template specialization for CL_PROGRAM_BINARIES
@@ -6828,6 +7439,18 @@ inline cl_int cl::Program::setSpecializationConstant(cl_uint index, const bool &
 }
 #endif // CL_HPP_TARGET_OPENCL_VERSION >= 220
 
+inline Kernel::Kernel(const Program& program, const string& name, cl_int* err)
+{
+    cl_int error;
+
+    object_ = ::clCreateKernel(program(), name.c_str(), &error);
+    detail::errHandler(error, __CREATE_KERNEL_ERR);
+
+    if (err != nullptr) {
+        *err = error;
+    }
+}
+
 inline Kernel::Kernel(const Program& program, const char* name, cl_int* err)
 {
     cl_int error;
@@ -6838,27 +7461,24 @@ inline Kernel::Kernel(const Program& program, const char* name, cl_int* err)
     if (err != nullptr) {
         *err = error;
     }
-
 }
 
 #ifdef cl_khr_external_memory
 enum class ExternalMemoryType : cl_external_memory_handle_type_khr
 {
     None = 0,
-
+#ifdef cl_khr_external_memory_opaque_fd
     OpaqueFd = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR,
+#endif // cl_khr_external_memory_opaque_fd
+#ifdef cl_khr_external_memory_win32
     OpaqueWin32 = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR,
     OpaqueWin32Kmt = CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR,
-
-    D3D11Texture = CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KHR,
-    D3D11TextureKmt = CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KMT_KHR,
-
-    D3D12Heap = CL_EXTERNAL_MEMORY_HANDLE_D3D12_HEAP_KHR,
-    D3D12Resource = CL_EXTERNAL_MEMORY_HANDLE_D3D12_RESOURCE_KHR,
-
+#endif // cl_khr_external_memory_win32
+#ifdef cl_khr_external_memory_dma_buf
     DmaBuf = CL_EXTERNAL_MEMORY_HANDLE_DMA_BUF_KHR,
+#endif // cl_khr_external_memory_dma_buf
 };
-#endif
+#endif // cl_khr_external_memory
 
 enum class QueueProperties : cl_command_queue_properties
 {
@@ -6933,7 +7553,7 @@ private:
 
     static void initMemoryExtension(const cl::Device& device) 
     {
-        auto platform = device.getInfo<CL_DEVICE_PLATFORM>();
+        auto platform = device.getInfo<CL_DEVICE_PLATFORM>()();
 
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clEnqueueAcquireExternalMemObjectsKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clEnqueueReleaseExternalMemObjectsKHR);
@@ -8111,6 +8731,162 @@ public:
     }
 
 #if CL_HPP_TARGET_OPENCL_VERSION >= 200
+
+    /**
+    * Enqueues a command that copies a region of memory from the source pointer to the destination pointer.
+    * This function is specifically for transferring data between the host and a coarse-grained SVM buffer.
+    */
+    template<typename T>
+    cl_int enqueueMemcpySVM(
+            T *dst_ptr,
+            const T *src_ptr,
+            cl_bool blocking,
+            size_type size,
+            const vector<Event> *events = nullptr,
+            Event *event = nullptr) const {
+        cl_event tmp;
+        cl_int err = detail::errHandler(::clEnqueueSVMMemcpy(
+                object_, blocking, static_cast<void *>(dst_ptr), static_cast<const void *>(src_ptr), size,
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != nullptr) ? &tmp : nullptr), __ENQUEUE_COPY_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
+    /**
+    *Enqueues a command that will copy data from one coarse-grained SVM buffer to another.
+    *This function takes two cl::pointer instances representing the destination and source buffers.
+    */
+    template<typename T, class D>
+    cl_int enqueueMemcpySVM(
+            cl::pointer<T, D> &dst_ptr,
+            const cl::pointer<T, D> &src_ptr,
+            cl_bool blocking,
+            size_type size,
+            const vector<Event> *events = nullptr,
+            Event *event = nullptr) const {
+        cl_event tmp;
+        cl_int err = detail::errHandler(::clEnqueueSVMMemcpy(
+                object_, blocking, static_cast<void *>(dst_ptr.get()), static_cast<const void *>(src_ptr.get()),
+                size,
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != nullptr) ? &tmp : nullptr), __ENQUEUE_COPY_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
+    /**
+    * Enqueues a command that will allow the host to update a region of a coarse-grained SVM buffer.
+    * This variant takes a cl::vector instance.
+    */
+    template<typename T, class Alloc>
+    cl_int enqueueMemcpySVM(
+            cl::vector<T, Alloc> &dst_container,
+            const cl::vector<T, Alloc> &src_container,
+            cl_bool blocking,
+            const vector<Event> *events = nullptr,
+            Event *event = nullptr) const {
+        cl_event tmp;
+        if(src_container.size() != dst_container.size()){
+            return detail::errHandler(CL_INVALID_VALUE,__ENQUEUE_COPY_SVM_ERR);
+        }
+        cl_int err = detail::errHandler(::clEnqueueSVMMemcpy(
+                object_, blocking, static_cast<void *>(dst_container.data()),
+                static_cast<const void *>(src_container.data()),
+                dst_container.size() * sizeof(T),
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != NULL) ? &tmp : nullptr), __ENQUEUE_COPY_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
+    /**
+    * Enqueues a command to fill a SVM buffer with a pattern.
+    *
+    */
+    template<typename T, typename PatternType>
+    cl_int enqueueMemFillSVM(
+            T *ptr,
+            PatternType pattern,
+            size_type size,
+            const vector<Event> *events = nullptr,
+            Event *event = nullptr) const {
+        cl_event tmp;
+        cl_int err = detail::errHandler(::clEnqueueSVMMemFill(
+                object_, static_cast<void *>(ptr), static_cast<void *>(&pattern),
+                sizeof(PatternType), size,
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != nullptr) ? &tmp : nullptr), __ENQUEUE_FILL_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
+    /**
+    * Enqueues a command that fills a region of a coarse-grained SVM buffer with a specified pattern.
+    * This variant takes a cl::pointer instance.
+    */
+    template<typename T, class D, typename PatternType>
+    cl_int enqueueMemFillSVM(
+            cl::pointer<T, D> &ptr,
+            PatternType pattern,
+            size_type size,
+            const vector<Event> *events = nullptr,
+            Event *event = nullptr) const {
+        cl_event tmp;
+        cl_int err = detail::errHandler(::clEnqueueSVMMemFill(
+                object_, static_cast<void *>(ptr.get()), static_cast<void *>(&pattern),
+                sizeof(PatternType), size,
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != nullptr) ? &tmp : nullptr), __ENQUEUE_FILL_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
+    /**
+    * Enqueues a command that will allow the host to fill a region of a coarse-grained SVM buffer with a specified pattern.
+    * This variant takes a cl::vector instance.
+    */
+    template<typename T, class Alloc, typename PatternType>
+    cl_int enqueueMemFillSVM(
+            cl::vector<T, Alloc> &container,
+            PatternType pattern,
+            const vector<Event> *events = nullptr,
+            Event* event = nullptr) const
+    {
+        cl_event tmp;
+        cl_int err = detail::errHandler(::clEnqueueSVMMemFill(
+                object_, static_cast<void *>(container.data()), static_cast<void *>(&pattern),
+                sizeof(PatternType), container.size() * sizeof(T),
+                (events != nullptr) ? (cl_uint) events->size() : 0,
+                (events != nullptr && events->size() > 0) ? (cl_event *) &events->front() : nullptr,
+                (event != nullptr) ? &tmp : NULL), __ENQUEUE_FILL_SVM_ERR);
+
+        if (event != nullptr && err == CL_SUCCESS)
+            *event = tmp;
+
+        return err;
+    }
+
     /**
      * Enqueues a command that will allow the host to update a region of a coarse-grained SVM buffer.
      * This variant takes a raw SVM pointer.
@@ -8130,7 +8906,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_MAP_BUFFER_ERR);
+            __ENQUEUE_MAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -8158,7 +8934,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_MAP_BUFFER_ERR);
+            __ENQUEUE_MAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -8184,7 +8960,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_MAP_BUFFER_ERR);
+            __ENQUEUE_MAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -8233,7 +9009,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+            __ENQUEUE_UNMAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -8258,7 +9034,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+            __ENQUEUE_UNMAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -8283,7 +9059,7 @@ public:
             (events != nullptr) ? (cl_uint)events->size() : 0,
             (events != nullptr && events->size() > 0) ? (cl_event*)&events->front() : nullptr,
             (event != nullptr) ? &tmp : nullptr),
-            __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+            __ENQUEUE_UNMAP_SVM_ERR);
 
         if (event != nullptr && err == CL_SUCCESS)
             *event = tmp;
@@ -9301,7 +10077,7 @@ inline cl_int enqueueMapSVM(
     cl_int error;
     CommandQueue queue = CommandQueue::getDefault(&error);
     if (error != CL_SUCCESS) {
-        return detail::errHandler(error, __ENQUEUE_MAP_BUFFER_ERR);
+        return detail::errHandler(error, __ENQUEUE_MAP_SVM_ERR);
     }
 
     return queue.enqueueMapSVM(
@@ -9348,7 +10124,7 @@ inline cl_int enqueueMapSVM(
     cl_int error;
     CommandQueue queue = CommandQueue::getDefault(&error);
     if (error != CL_SUCCESS) {
-        return detail::errHandler(error, __ENQUEUE_MAP_BUFFER_ERR);
+        return detail::errHandler(error, __ENQUEUE_MAP_SVM_ERR);
     }
 
     return queue.enqueueMapSVM(
@@ -9400,11 +10176,11 @@ inline cl_int enqueueUnmapSVM(
     cl_int error;
     CommandQueue queue = CommandQueue::getDefault(&error);
     if (error != CL_SUCCESS) {
-        return detail::errHandler(error, __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        return detail::errHandler(error, __ENQUEUE_UNMAP_SVM_ERR);
     }
 
     return detail::errHandler(queue.enqueueUnmapSVM(ptr, events, event), 
-        __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        __ENQUEUE_UNMAP_SVM_ERR);
 
 }
 
@@ -9422,11 +10198,11 @@ inline cl_int enqueueUnmapSVM(
     cl_int error;
     CommandQueue queue = CommandQueue::getDefault(&error);
     if (error != CL_SUCCESS) {
-        return detail::errHandler(error, __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        return detail::errHandler(error, __ENQUEUE_UNMAP_SVM_ERR);
     }
 
     return detail::errHandler(queue.enqueueUnmapSVM(ptr, events, event),
-        __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        __ENQUEUE_UNMAP_SVM_ERR);
 }
 
 /**
@@ -9443,11 +10219,11 @@ inline cl_int enqueueUnmapSVM(
     cl_int error;
     CommandQueue queue = CommandQueue::getDefault(&error);
     if (error != CL_SUCCESS) {
-        return detail::errHandler(error, __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        return detail::errHandler(error, __ENQUEUE_UNMAP_SVM_ERR);
     }
 
     return detail::errHandler(queue.enqueueUnmapSVM(container, events, event),
-        __ENQUEUE_UNMAP_MEM_OBJECT_ERR);
+        __ENQUEUE_UNMAP_SVM_ERR);
 }
 
 #endif // #if CL_HPP_TARGET_OPENCL_VERSION >= 200
@@ -9523,15 +10299,15 @@ inline cl_int copy( const CommandQueue &queue, IteratorType startIterator, Itera
     if( error != CL_SUCCESS ) {
         return error;
     }
-#if defined(_MSC_VER)
+#if defined(_MSC_VER) && _MSC_VER < 1920
     std::copy(
-        startIterator, 
-        endIterator, 
+        startIterator,
+        endIterator,
         stdext::checked_array_iterator<DataType*>(
             pointer, length));
 #else
     std::copy(startIterator, endIterator, pointer);
-#endif
+#endif // defined(_MSC_VER) && _MSC_VER < 1920
     Event endEvent;
     error = queue.enqueueUnmapMemObject(buffer, pointer, 0, &endEvent);
     // if exceptions enabled, enqueueUnmapMemObject will throw
@@ -10400,6 +11176,24 @@ namespace compatibility {
 } // namespace compatibility
 
 #ifdef cl_khr_semaphore
+
+#ifdef cl_khr_external_semaphore
+enum ExternalSemaphoreType : cl_external_semaphore_handle_type_khr
+{
+    None = 0,
+#ifdef cl_khr_external_semaphore_opaque_fd
+    OpaqueFd = CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR,
+#endif // cl_khr_external_semaphore_opaque_fd
+#ifdef cl_khr_external_semaphore_sync_fd
+    SyncFd = CL_SEMAPHORE_HANDLE_SYNC_FD_KHR,
+#endif // cl_khr_external_semaphore_sync_fd
+#ifdef cl_khr_external_semaphore_win32
+    OpaqueWin32 = CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR,
+    OpaqueWin32Kmt = CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR,
+#endif // cl_khr_external_semaphore_win32
+};
+#endif // cl_khr_external_semaphore
+
 class Semaphore : public detail::Wrapper<cl_semaphore_khr>
 {
 public:
@@ -10447,7 +11241,7 @@ public:
         }
 
         return detail::errHandler(
-            detail::getInfo(&pfn_clGetSemaphoreInfoKHR, object_, name, param),
+            detail::getInfo(pfn_clGetSemaphoreInfoKHR, object_, name, param),
             __GET_SEMAPHORE_KHR_INFO_ERR);
     }
     template <cl_semaphore_info_khr name> typename
@@ -10463,6 +11257,36 @@ public:
         return param;      
     }
 
+#ifdef cl_khr_external_semaphore
+    template <typename T>
+    cl_int getHandleForTypeKHR(
+        const Device& device, cl_external_semaphore_handle_type_khr name, T* param) const
+    {
+        if (pfn_clGetSemaphoreHandleForTypeKHR == nullptr) {
+            return detail::errHandler(CL_INVALID_OPERATION,
+                                      __GET_SEMAPHORE_HANDLE_FOR_TYPE_KHR_ERR);
+        }
+
+        return detail::errHandler(
+            detail::getInfo(
+                pfn_clGetSemaphoreHandleForTypeKHR, object_, device(), name, param),
+                __GET_SEMAPHORE_HANDLE_FOR_TYPE_KHR_ERR);
+    }
+
+    template <cl_external_semaphore_handle_type_khr type> typename
+    detail::param_traits<detail::cl_external_semaphore_handle_type_khr, type>::param_type
+        getHandleForTypeKHR(const Device& device, cl_int* err = nullptr) const
+    {
+        typename detail::param_traits<
+        detail::cl_external_semaphore_handle_type_khr, type>::param_type param;
+        cl_int result = getHandleForTypeKHR(device, type, &param);
+        if (err != nullptr) {
+            *err = result;
+        }
+        return param;
+    }
+#endif // cl_khr_external_semaphore
+
     cl_int retain()
     { 
         if (pfn_clRetainSemaphoreKHR == nullptr) {
@@ -10488,13 +11312,17 @@ private:
     {
 #if CL_HPP_TARGET_OPENCL_VERSION >= 120
         Device device = context.getInfo<CL_CONTEXT_DEVICES>().at(0);
-        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>();
+        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>()();
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clCreateSemaphoreWithPropertiesKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clReleaseSemaphoreKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clRetainSemaphoreKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clEnqueueWaitSemaphoresKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clEnqueueSignalSemaphoresKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clGetSemaphoreInfoKHR);
+#ifdef cl_khr_external_semaphore
+        CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clGetSemaphoreHandleForTypeKHR);
+#endif // cl_khr_external_semaphore
+
 #else
         CL_HPP_INIT_CL_EXT_FCN_PTR_(clCreateSemaphoreWithPropertiesKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_(clReleaseSemaphoreKHR);
@@ -10502,12 +11330,19 @@ private:
         CL_HPP_INIT_CL_EXT_FCN_PTR_(clEnqueueWaitSemaphoresKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_(clEnqueueSignalSemaphoresKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_(clGetSemaphoreInfoKHR);
+#ifdef cl_khr_external_semaphore
+        CL_HPP_INIT_CL_EXT_FCN_PTR_(clGetSemaphoreHandleForTypeKHR);
+#endif // cl_khr_external_semaphore
+
 #endif
         if ((pfn_clCreateSemaphoreWithPropertiesKHR == nullptr) &&
             (pfn_clReleaseSemaphoreKHR              == nullptr) &&
             (pfn_clRetainSemaphoreKHR               == nullptr) &&
             (pfn_clEnqueueWaitSemaphoresKHR         == nullptr) &&
             (pfn_clEnqueueSignalSemaphoresKHR       == nullptr) &&
+#ifdef cl_khr_external_semaphore
+            (pfn_clGetSemaphoreHandleForTypeKHR     == nullptr) &&
+#endif // cl_khr_external_semaphore
             (pfn_clGetSemaphoreInfoKHR              == nullptr))
         {
             detail::errHandler(CL_INVALID_VALUE, __CREATE_SEMAPHORE_KHR_WITH_PROPERTIES_ERR);
@@ -10689,6 +11524,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandBarrierWithWaitListKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 (sync_points_vec != nullptr) ? (cl_uint) sync_points_vec->size() : 0,
                 (sync_points_vec != nullptr && sync_points_vec->size() > 0) ? &sync_points_vec->front() : nullptr,
                 (sync_point != nullptr) ? &tmp_sync_point : nullptr,
@@ -10720,6 +11558,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandCopyBufferKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 src(),
                 dst(),
                 src_offset,
@@ -10760,6 +11601,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandCopyBufferRectKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 src(),
                 dst(),
                 src_origin.data(),
@@ -10800,6 +11644,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandCopyBufferToImageKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 src(),
                 dst(),
                 src_offset,
@@ -10836,6 +11683,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandCopyImageKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 src(),
                 dst(),
                 src_origin.data(),
@@ -10872,6 +11722,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandCopyImageToBufferKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 src(),
                 dst(),
                 src_origin.data(),
@@ -10908,6 +11761,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandFillBufferKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 buffer(),
                 static_cast<void*>(&pattern),
                 sizeof(PatternType),
@@ -10943,6 +11799,9 @@ public:
         cl_int error = detail::errHandler(
             pfn_clCommandFillImageKHR(object_,
                 (command_queue != nullptr) ? (*command_queue)() : nullptr,
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+                nullptr, // Properties
+#endif
                 image(),
                 static_cast<void*>(&fillColor),
                 origin.data(),
@@ -10959,7 +11818,12 @@ public:
         return error;
     }
 
-    cl_int commandNDRangeKernel(const cl::vector<cl_ndrange_kernel_command_properties_khr> &properties,
+    cl_int commandNDRangeKernel(
+#if CL_KHR_COMMAND_BUFFER_EXTENSION_VERSION > CL_MAKE_VERSION(0, 9, 4)
+            const cl::vector<cl_command_properties_khr> &properties,
+#else
+            const cl::vector<cl_ndrange_kernel_command_properties_khr> &properties,
+#endif
         const Kernel& kernel,
         const NDRange& offset,
         const NDRange& global,
@@ -10997,6 +11861,8 @@ public:
     }
 
 #if defined(cl_khr_command_buffer_mutable_dispatch)
+#if CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_EXTENSION_VERSION <                 \
+    CL_MAKE_VERSION(0, 9, 2)
     cl_int updateMutableCommands(const cl_mutable_base_config_khr* mutable_config)
     {
         if (pfn_clUpdateMutableCommandsKHR == nullptr) {
@@ -11006,6 +11872,21 @@ public:
         return detail::errHandler(pfn_clUpdateMutableCommandsKHR(object_, mutable_config),
                         __UPDATE_MUTABLE_COMMANDS_KHR_ERR);
     }
+#else
+    template <int ArrayLength>
+    cl_int updateMutableCommands(std::array<cl_command_buffer_update_type_khr,
+                                            ArrayLength> &config_types,
+                                 std::array<void *, ArrayLength> &configs) {
+        if (pfn_clUpdateMutableCommandsKHR == nullptr) {
+            return detail::errHandler(CL_INVALID_OPERATION,
+                                      __UPDATE_MUTABLE_COMMANDS_KHR_ERR);
+        }
+        return detail::errHandler(
+            pfn_clUpdateMutableCommandsKHR(object_, configs.length(),
+                                           config_types.data().configs.data()),
+            __UPDATE_MUTABLE_COMMANDS_KHR_ERR);
+    }
+#endif /* CL_KHR_COMMAND_BUFFER_MUTABLE_DISPATCH_EXTENSION_VERSION */
 #endif /* cl_khr_command_buffer_mutable_dispatch */
 
 private:
@@ -11014,7 +11895,7 @@ private:
     static void initExtensions(const cl::Device& device)
     {
 #if CL_HPP_TARGET_OPENCL_VERSION >= 120
-        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>();
+        cl_platform_id platform = device.getInfo<CL_DEVICE_PLATFORM>()();
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clCreateCommandBufferKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clFinalizeCommandBufferKHR);
         CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_(platform, clRetainCommandBufferKHR);
@@ -11212,7 +12093,12 @@ public:
 #undef __ENQUEUE_FILL_IMAGE_ERR            
 #undef __ENQUEUE_COPY_IMAGE_TO_BUFFER_ERR  
 #undef __ENQUEUE_COPY_BUFFER_TO_IMAGE_ERR  
-#undef __ENQUEUE_MAP_BUFFER_ERR            
+#undef __ENQUEUE_MAP_BUFFER_ERR
+#undef __ENQUEUE_MAP_IMAGE_ERR
+#undef __ENQUEUE_MAP_SVM_ERR
+#undef __ENQUEUE_FILL_SVM_ERR
+#undef __ENQUEUE_COPY_SVM_ERR
+#undef __ENQUEUE_UNMAP_SVM_ERR              
 #undef __ENQUEUE_MAP_IMAGE_ERR             
 #undef __ENQUEUE_UNMAP_MEM_OBJECT_ERR      
 #undef __ENQUEUE_NDRANGE_KERNEL_ERR        
@@ -11249,8 +12135,12 @@ public:
 #undef __GET_DEVICE_AND_HOST_TIMER_ERR
 #undef __GET_SEMAPHORE_KHR_INFO_ERR
 #undef __CREATE_SEMAPHORE_KHR_WITH_PROPERTIES_ERR
+#undef __GET_IMAGE_REQUIREMENT_INFO_EXT_ERR
 #undef __ENQUEUE_WAIT_SEMAPHORE_KHR_ERR
 #undef __ENQUEUE_SIGNAL_SEMAPHORE_KHR_ERR
+#undef __RETAIN_SEMAPHORE_KHR_ERR
+#undef __RELEASE_SEMAPHORE_KHR_ERR
+#undef __GET_SEMAPHORE_HANDLE_FOR_TYPE_KHR_ERR
 
 #endif //CL_HPP_USER_OVERRIDE_ERROR_STRINGS
 
@@ -11259,7 +12149,6 @@ public:
 #undef CL_HPP_INIT_CL_EXT_FCN_PTR_
 #undef CL_HPP_INIT_CL_EXT_FCN_PTR_PLATFORM_
 
-#undef CL_HPP_NOEXCEPT_
 #undef CL_HPP_DEFINE_STATIC_MEMBER_
 
 } // namespace cl
diff --git a/lib/CL/clCommandBarrierWithWaitListKHR.c b/lib/CL/clCommandBarrierWithWaitListKHR.c
index 6036d78044699c5fb90c0397f542e47be4928f64..55be79d7491bc502a1fc44457cc176d726331745 100644
--- a/lib/CL/clCommandBarrierWithWaitListKHR.c
+++ b/lib/CL/clCommandBarrierWithWaitListKHR.c
@@ -30,6 +30,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandBarrierWithWaitListKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr *sync_point_wait_list,
     cl_sync_point_khr *sync_point,
diff --git a/lib/CL/clCommandCopyBufferKHR.c b/lib/CL/clCommandCopyBufferKHR.c
index 0a2e961a90e186c1a2af92af2568a7fac932fc39..95f92adb65b9527cf908c3d13160ded5b9214b65 100644
--- a/lib/CL/clCommandCopyBufferKHR.c
+++ b/lib/CL/clCommandCopyBufferKHR.c
@@ -31,6 +31,7 @@
 extern CL_API_ENTRY cl_int CL_API_CALL
 POname (clCommandCopyBufferKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset,
     size_t size, cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr *sync_point_wait_list,
diff --git a/lib/CL/clCommandCopyBufferRectKHR.c b/lib/CL/clCommandCopyBufferRectKHR.c
index 952388a030101152ce188535b6f062addce6811f..dd3d60943bfdc96231177b609efb77570c568bde 100644
--- a/lib/CL/clCommandCopyBufferRectKHR.c
+++ b/lib/CL/clCommandCopyBufferRectKHR.c
@@ -29,6 +29,7 @@
 CL_API_ENTRY cl_int CL_API_CALL
 POname (clCommandCopyBufferRectKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer, cl_mem dst_buffer, const size_t *src_origin,
     const size_t *dst_origin, const size_t *region, size_t src_row_pitch,
     size_t src_slice_pitch, size_t dst_row_pitch, size_t dst_slice_pitch,
diff --git a/lib/CL/clCommandCopyBufferToImageKHR.c b/lib/CL/clCommandCopyBufferToImageKHR.c
index 277a28f62e8dedae1d58136c113088368b133cf3..4324a9271a684d2dca6910bdb5ca8baeb29b1084 100644
--- a/lib/CL/clCommandCopyBufferToImageKHR.c
+++ b/lib/CL/clCommandCopyBufferToImageKHR.c
@@ -29,6 +29,7 @@
 CL_API_ENTRY cl_int CL_API_CALL
 POname (clCommandCopyBufferToImageKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_buffer, cl_mem dst_image, size_t src_offset,
     const size_t *dst_origin, const size_t *region,
     cl_uint num_sync_points_in_wait_list,
diff --git a/lib/CL/clCommandCopyImageKHR.c b/lib/CL/clCommandCopyImageKHR.c
index dcb1604d6c109d381d2d9c0b7c04d86f8aba917f..386915c914d25dec88dc21410f9c082a6e533830 100644
--- a/lib/CL/clCommandCopyImageKHR.c
+++ b/lib/CL/clCommandCopyImageKHR.c
@@ -31,6 +31,7 @@
 extern CL_API_ENTRY cl_int CL_API_CALL
 POname (clCommandCopyImageKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image, cl_mem dst_image, const size_t *src_origin,
     const size_t *dst_origin, const size_t *region,
     cl_uint num_sync_points_in_wait_list,
diff --git a/lib/CL/clCommandCopyImageToBufferKHR.c b/lib/CL/clCommandCopyImageToBufferKHR.c
index 62732d845f055171792ffd6a9f14e39f32a8d6b8..9b278d9ad970584b04321d6043e9ecb8f6871851 100644
--- a/lib/CL/clCommandCopyImageToBufferKHR.c
+++ b/lib/CL/clCommandCopyImageToBufferKHR.c
@@ -29,6 +29,7 @@
 extern CL_API_ENTRY cl_int CL_API_CALL
 POname (clCommandCopyImageToBufferKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem src_image, cl_mem dst_buffer, const size_t *src_origin,
     const size_t *region, size_t dst_offset,
     cl_uint num_sync_points_in_wait_list,
diff --git a/lib/CL/clCommandFillBufferKHR.c b/lib/CL/clCommandFillBufferKHR.c
index aefe51966025b6b808e9c87cec0033bef4515aac..bc569a2e67f288d294f161a5d70edb50e6f93174 100644
--- a/lib/CL/clCommandFillBufferKHR.c
+++ b/lib/CL/clCommandFillBufferKHR.c
@@ -29,6 +29,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandFillBufferKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem buffer, const void *pattern, size_t pattern_size, size_t offset,
     size_t size, cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr *sync_point_wait_list,
diff --git a/lib/CL/clCommandFillImageKHR.c b/lib/CL/clCommandFillImageKHR.c
index af8cd58efc4f70f524fc9ee1044edd9339e55bc2..87b39c12a9ed6694a855b41b3cc1477f7718a544 100644
--- a/lib/CL/clCommandFillImageKHR.c
+++ b/lib/CL/clCommandFillImageKHR.c
@@ -32,6 +32,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandFillImageKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     cl_mem image, const void *fill_color, const size_t *origin,
     const size_t *region, cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr *sync_point_wait_list,
diff --git a/lib/CL/clCommandNDRangeKernelKHR.c b/lib/CL/clCommandNDRangeKernelKHR.c
index 8c94ecc0cd3038909e1b76290ebb5ca0df6839ac..46d5578164d0679ece981bb0b19a670ee7be7d95 100644
--- a/lib/CL/clCommandNDRangeKernelKHR.c
+++ b/lib/CL/clCommandNDRangeKernelKHR.c
@@ -29,7 +29,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandNDRangeKernelKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
-    const cl_ndrange_kernel_command_properties_khr *properties,
+    const cl_command_properties_khr* properties,
     cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset,
     const size_t *global_work_size, const size_t *local_work_size,
     cl_uint num_sync_points_in_wait_list,
diff --git a/lib/CL/clCommandSVMMemFillKHR.c b/lib/CL/clCommandSVMMemFillKHR.c
index 06fb71f75ab6ccd1ca785342e36acb29fbd278ff..0fa9ee738b54b9f2034d402567210e0cc9f8f6e5 100644
--- a/lib/CL/clCommandSVMMemFillKHR.c
+++ b/lib/CL/clCommandSVMMemFillKHR.c
@@ -29,6 +29,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandSVMMemFillKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void *svm_ptr, const void *pattern,
     size_t pattern_size, size_t size,
     cl_uint num_sync_points_in_wait_list,
diff --git a/lib/CL/clCommandSVMMemcpyKHR.c b/lib/CL/clCommandSVMMemcpyKHR.c
index 84d29e21bca5b790d5afed8b01f4719a04f94f13..1836bc4df6d2870960f9358c8db49e6f70ab5bf6 100644
--- a/lib/CL/clCommandSVMMemcpyKHR.c
+++ b/lib/CL/clCommandSVMMemcpyKHR.c
@@ -29,6 +29,7 @@
 CL_API_ENTRY cl_int
 POname (clCommandSVMMemcpyKHR) (
     cl_command_buffer_khr command_buffer, cl_command_queue command_queue,
+    const cl_command_properties_khr* properties,
     void *dst_ptr, const void *src_ptr, size_t size,
     cl_uint num_sync_points_in_wait_list,
     const cl_sync_point_khr *sync_point_wait_list,
diff --git a/lib/CL/clCreateCommandQueue.c b/lib/CL/clCreateCommandQueue.c
index c7f1e1ca45e137a9888a44e2bfc3e9bca62dd96a..66d825a1bb14e41bfe35efa31f5fba5c0029a121 100644
--- a/lib/CL/clCreateCommandQueue.c
+++ b/lib/CL/clCreateCommandQueue.c
@@ -53,6 +53,17 @@ POname(clCreateCommandQueue)(cl_context context,
   POCL_GOTO_ERROR_ON ((properties & (~all_properties)), CL_INVALID_VALUE,
                       "Unknown properties requested\n");
 
+  cl_command_queue_properties supported_device_props;
+  if (properties & (CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT))
+    supported_device_props = device->on_dev_queue_props;
+  else
+    supported_device_props = device->on_host_queue_props | CL_QUEUE_HIDDEN;
+
+  POCL_GOTO_ERROR_ON (((properties & supported_device_props) != properties),
+                      CL_INVALID_QUEUE_PROPERTIES, "properties (%zu) are "
+                      "not supported by the device (%zu)\n",
+                      (size_t)properties, (size_t)supported_device_props);
+
   if (POCL_DEBUGGING_ON || pocl_cq_profiling_enabled)
     properties |= CL_QUEUE_PROFILING_ENABLE;
 
diff --git a/lib/CL/clCreateCommandQueueWithProperties.c b/lib/CL/clCreateCommandQueueWithProperties.c
index f111323ed47c2f6f8b6463332a186736925e1c70..19de0de6f3fd8fc0b6bf5212cd5e9fbdcb5f00da 100644
--- a/lib/CL/clCreateCommandQueueWithProperties.c
+++ b/lib/CL/clCreateCommandQueueWithProperties.c
@@ -74,7 +74,7 @@ POname(clCreateCommandQueueWithProperties)(cl_context context,
         case CL_QUEUE_SIZE:
           {
             POCL_GOTO_ERROR_ON ((queue_size_set > 0), CL_INVALID_VALUE,
-                                "CL_QUEUE_PROPERTIES was already set");
+                                "CL_QUEUE_SIZE was already set");
             queue_size = (cl_uint)properties[i + 1];
             queue_size_set++;
             i += 2;
@@ -83,6 +83,8 @@ POname(clCreateCommandQueueWithProperties)(cl_context context,
         case CL_QUEUE_PRIORITY_KHR:
           {
             cl_queue_properties value = properties[i + 1];
+            POCL_GOTO_ERROR_ON ((queue_priority_set > 0), CL_INVALID_VALUE,
+                                "CL_QUEUE_PRIORITY_KHR was already set");
             POCL_GOTO_ERROR_ON ((value != CL_QUEUE_PRIORITY_HIGH_KHR
                                  && value != CL_QUEUE_PRIORITY_MED_KHR
                                  && value != CL_QUEUE_PRIORITY_LOW_KHR),
@@ -98,6 +100,8 @@ POname(clCreateCommandQueueWithProperties)(cl_context context,
         case CL_QUEUE_THROTTLE_KHR:
           {
             cl_queue_properties value = properties[i + 1];
+            POCL_GOTO_ERROR_ON ((queue_throttle_set > 0), CL_INVALID_VALUE,
+                                "CL_QUEUE_THROTTLE_KHR was already set");
             POCL_GOTO_ERROR_ON ((value != CL_QUEUE_THROTTLE_HIGH_KHR
                                  && value != CL_QUEUE_THROTTLE_MED_KHR
                                  && value != CL_QUEUE_THROTTLE_LOW_KHR),
@@ -129,13 +133,29 @@ POname(clCreateCommandQueueWithProperties)(cl_context context,
           POCL_GOTO_ERROR_COND((queue_throttle_set), CL_INVALID_QUEUE_PROPERTIES);
 
          // create a device side queue
-          POCL_GOTO_ERROR_ON (1, CL_INVALID_QUEUE_PROPERTIES,
+          POCL_GOTO_ERROR_ON ((device->on_dev_queue_props == 0),
+                              CL_INVALID_QUEUE_PROPERTIES,
                               "Device-side enqueue is not supported "
                               "by any device\n");
         }
       else
-        POCL_GOTO_ERROR_ON((queue_size > 0), CL_INVALID_VALUE,
-                           "To specify queue size, you must use CL_QUEUE_ON_DEVICE in flags\n");
+        {
+          POCL_GOTO_ERROR_ON (
+            (queue_size > 0), CL_INVALID_QUEUE_PROPERTIES,
+            "Queue size can only be specified for on-device queues\n");
+          POCL_GOTO_ERROR_ON (
+            (queue_priority_set
+             && (strstr (device->extensions, "cl_khr_priority_hints")
+                 == NULL)),
+            CL_INVALID_QUEUE_PROPERTIES,
+            "device does not support cl_khr_priority_hints\n");
+          POCL_GOTO_ERROR_ON (
+            (queue_throttle_set
+             && (strstr (device->extensions, "cl_khr_throttle_hints")
+                 == NULL)),
+            CL_INVALID_QUEUE_PROPERTIES,
+            "device does not support cl_khr_throttle_hints\n");
+        }
 
       /* validate flags */
       POCL_GOTO_ERROR_ON ((queue_props & (~valid_prop_flags)),
@@ -169,4 +189,17 @@ ERROR:
     }
   return NULL;
 }
-POsym(clCreateCommandQueueWithProperties)
+POsym (clCreateCommandQueueWithProperties)
+
+  CL_API_ENTRY cl_int CL_API_CALL POname (clSetCommandQueueProperty) (
+    cl_command_queue command_queue,
+    cl_command_queue_properties properties,
+    cl_bool enable,
+    cl_command_queue_properties *old_properties)
+    CL_API_SUFFIX__VERSION_1_0_DEPRECATED
+{
+  /* CL_INVALID_OPERATION if no devices in the context associated with
+   * command_queue support modifying the properties of a command-queue */
+  return CL_INVALID_OPERATION;
+}
+POsym (clSetCommandQueueProperty)
diff --git a/lib/CL/clEnqueueCopyBufferToImage.c b/lib/CL/clEnqueueCopyBufferToImage.c
index 0f02d4aa52cea41b537941ff59bc1e33d3eff361..ac4dbe9c8b0eb2bc9546a5cf86385d778d6873fb 100644
--- a/lib/CL/clEnqueueCopyBufferToImage.c
+++ b/lib/CL/clEnqueueCopyBufferToImage.c
@@ -73,7 +73,7 @@ pocl_copy_buffer_to_image_common (
       else
         {
           return POname (clCommandCopyBufferRectKHR) (
-              command_buffer, command_queue, src_buffer, dst_image->buffer,
+              command_buffer, command_queue, NULL, src_buffer, dst_image->buffer,
               src_origin, i1d_origin, i1d_region, dst_image->image_row_pitch,
               0, dst_image->image_row_pitch, 0, num_items_in_wait_list,
               sync_point_wait_list, sync_point, mutable_handle);
diff --git a/lib/CL/clEnqueueCopyImageToBuffer.c b/lib/CL/clEnqueueCopyImageToBuffer.c
index 12e30343712a01be2d1e700b6ae400764cae3dce..11f915086f90f35591df585b28313c02f1b61bac 100644
--- a/lib/CL/clEnqueueCopyImageToBuffer.c
+++ b/lib/CL/clEnqueueCopyImageToBuffer.c
@@ -68,7 +68,7 @@ pocl_copy_image_to_buffer_common (
       else
         {
           return POname (clCommandCopyBufferRectKHR) (
-              command_buffer, command_queue, src_image->buffer, dst_buffer,
+              command_buffer, command_queue, NULL, src_image->buffer, dst_buffer,
               i1d_origin, dst_origin, i1d_region, src_image->image_row_pitch,
               0, src_image->image_row_pitch, 0, num_items_in_wait_list,
               sync_point_wait_list, sync_point, mutable_handle);
diff --git a/lib/CL/clEnqueueSVMMemFill.c b/lib/CL/clEnqueueSVMMemFill.c
index 8b7e67ac050645a2caa39d303d7a3d6243d95145..da5804071425ee4dbe46bae6c71871399e479270 100644
--- a/lib/CL/clEnqueueSVMMemFill.c
+++ b/lib/CL/clEnqueueSVMMemFill.c
@@ -95,7 +95,7 @@ pocl_svm_memfill_common (cl_command_buffer_khr command_buffer,
   size_t offset = svm_ptr - dst_svm_ptr->vm_ptr;
   if (command_buffer)
     errcode = POname (clCommandFillBufferKHR) (
-        command_buffer, NULL, dst_svm_ptr->shadow_cl_mem, pattern,
+        command_buffer, NULL, NULL, dst_svm_ptr->shadow_cl_mem, pattern,
         pattern_size, offset, size, num_items_in_wait_list,
         sync_point_wait_list, sync_point, NULL);
   else
diff --git a/lib/CL/clEnqueueSVMMemcpy.c b/lib/CL/clEnqueueSVMMemcpy.c
index 4d589bca0500b38c6fe97a3ae3ff1fd4d3de071e..b8debf16ef5a3a908a5d0fdef5aa35afbb48d76d 100644
--- a/lib/CL/clEnqueueSVMMemcpy.c
+++ b/lib/CL/clEnqueueSVMMemcpy.c
@@ -75,7 +75,7 @@ pocl_svm_memcpy_common (cl_command_buffer_khr command_buffer,
          buffers. */
       if (command_buffer)
         errcode = POname (clCommandCopyBufferKHR) (
-            command_buffer, NULL, src_svm_ptr->shadow_cl_mem,
+            command_buffer, NULL, NULL, src_svm_ptr->shadow_cl_mem,
             dst_svm_ptr->shadow_cl_mem, src_ptr - src_svm_ptr->vm_ptr,
             dst_ptr - dst_svm_ptr->vm_ptr, size, num_items_in_wait_list,
             sync_point_wait_list, sync_point, NULL);
diff --git a/lib/CL/clGetPlatformIDs.c b/lib/CL/clGetPlatformIDs.c
index c0101000e746b739ce42fa042ea8c912de127151..e5a431ccc02a2d2fae57b2bb4dbff228b10e2e60 100644
--- a/lib/CL/clGetPlatformIDs.c
+++ b/lib/CL/clGetPlatformIDs.c
@@ -50,7 +50,7 @@ struct _cl_icd_dispatch pocl_dispatch = {
   &POname(clRetainCommandQueue), /* 10 */
   &POname(clReleaseCommandQueue),
   &POname(clGetCommandQueueInfo),
-  NULL /*clSetCommandQueueProperty*/,
+  &POname(clSetCommandQueueProperty),
   &POname(clCreateBuffer),
   &POname(clCreateImage2D),
   &POname(clCreateImage3D),
diff --git a/lib/CL/clRemapCommandBufferKHR.c b/lib/CL/clRemapCommandBufferKHR.c
index 0fe8006d92516abc68d928c15bd5ac81a5813239..8a79434f7df8adb2182e86b1de9acb1a87add3d4 100644
--- a/lib/CL/clRemapCommandBufferKHR.c
+++ b/lib/CL/clRemapCommandBufferKHR.c
@@ -73,13 +73,13 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
       {
       case CL_COMMAND_BARRIER:
         errcode = POname (clCommandBarrierWithWaitListKHR) (
-          new_cmdbuf, new_queue,
+          new_cmdbuf, new_queue, NULL,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
           cmd->sync.syncpoint.sync_point_wait_list, NULL, NULL);
         break;
       case CL_COMMAND_COPY_BUFFER:
         errcode = POname (clCommandCopyBufferKHR) (
-          new_cmdbuf, new_queue, cmd->command.copy.src, cmd->command.copy.dst,
+          new_cmdbuf, new_queue, NULL, cmd->command.copy.src, cmd->command.copy.dst,
           cmd->command.copy.src_offset, cmd->command.copy.dst_offset,
           cmd->command.copy.size,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -87,7 +87,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
         break;
       case CL_COMMAND_COPY_BUFFER_RECT:
         errcode = POname (clCommandCopyBufferRectKHR) (
-          new_cmdbuf, new_queue, cmd->command.copy_rect.src,
+          new_cmdbuf, new_queue, NULL, cmd->command.copy_rect.src,
           cmd->command.copy_rect.dst, cmd->command.copy_rect.src_origin,
           cmd->command.copy_rect.dst_origin, cmd->command.copy_rect.region,
           cmd->command.copy_rect.src_row_pitch,
@@ -99,7 +99,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
         break;
       case CL_COMMAND_COPY_BUFFER_TO_IMAGE:
         errcode = POname (clCommandCopyBufferToImageKHR) (
-          new_cmdbuf, new_queue, cmd->command.write_image.src,
+          new_cmdbuf, new_queue, NULL, cmd->command.write_image.src,
           cmd->command.write_image.dst, cmd->command.write_image.src_offset,
           cmd->command.write_image.origin, cmd->command.write_image.region,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -107,7 +107,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
         break;
       case CL_COMMAND_COPY_IMAGE_TO_BUFFER:
         errcode = POname (clCommandCopyImageToBufferKHR) (
-          new_cmdbuf, new_queue, cmd->command.read_image.src,
+          new_cmdbuf, new_queue, NULL, cmd->command.read_image.src,
           cmd->command.read_image.dst, cmd->command.read_image.origin,
           cmd->command.read_image.region, cmd->command.read_image.dst_offset,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -115,7 +115,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
         break;
       case CL_COMMAND_COPY_IMAGE:
         errcode = POname (clCommandCopyImageKHR) (
-          new_cmdbuf, new_queue, cmd->command.copy_image.src,
+          new_cmdbuf, new_queue, NULL, cmd->command.copy_image.src,
           cmd->command.copy_image.dst, cmd->command.copy_image.src_origin,
           cmd->command.copy_image.dst_origin, cmd->command.copy_image.region,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -124,7 +124,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
 
       case CL_COMMAND_FILL_BUFFER:
         errcode = POname (clCommandFillBufferKHR) (
-          new_cmdbuf, new_queue, cmd->command.memfill.dst,
+          new_cmdbuf, new_queue, NULL, cmd->command.memfill.dst,
           cmd->command.memfill.pattern, cmd->command.memfill.pattern_size,
           cmd->command.memfill.offset, cmd->command.memfill.size,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -132,7 +132,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
         break;
       case CL_COMMAND_FILL_IMAGE:
         errcode = POname (clCommandFillImageKHR) (
-          new_cmdbuf, new_queue, cmd->command.fill_image.dst,
+          new_cmdbuf, new_queue, NULL, cmd->command.fill_image.dst,
           cmd->command.fill_image.fill_pixel, cmd->command.fill_image.origin,
           cmd->command.fill_image.region,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
@@ -195,7 +195,7 @@ POname (clRemapCommandBufferKHR) (cl_command_buffer_khr command_buffer,
 
       case CL_COMMAND_SVM_MEMCPY:
         errcode = POname (clCommandSVMMemcpyKHR) (
-          new_cmdbuf, new_queue, cmd->command.svm_memcpy.dst,
+          new_cmdbuf, new_queue, NULL, cmd->command.svm_memcpy.dst,
           cmd->command.svm_memcpy.src, cmd->command.svm_memcpy.size,
           cmd->sync.syncpoint.num_sync_points_in_wait_list,
           cmd->sync.syncpoint.sync_point_wait_list, NULL, NULL);
diff --git a/lib/CL/devices/common.c b/lib/CL/devices/common.c
index 59e7c0e87bd65045764c1443740a88603aa1f6ed..eec255d86571ad4b16c752ae4e7e8d548b932293 100644
--- a/lib/CL/devices/common.c
+++ b/lib/CL/devices/common.c
@@ -1902,7 +1902,7 @@ static const cl_name_version OPENCL_EXTENSIONS[]
       { CL_MAKE_VERSION (1, 0, 0), "cl_khr_pci_bus_info" },
       { CL_MAKE_VERSION (1, 0, 0), "cl_khr_device_uuid" },
 
-      { CL_MAKE_VERSION (0, 9, 4), "cl_khr_command_buffer" },
+      { CL_MAKE_VERSION (0, 9, 5), "cl_khr_command_buffer" },
       { CL_MAKE_VERSION (0, 9, 1), "cl_khr_command_buffer_multi_device" },
       { CL_MAKE_VERSION (1, 0, 0), "cl_ext_float_atomics" },
       { CL_MAKE_VERSION (0, 1, 0), "cl_ext_buffer_device_address" },
diff --git a/lib/CL/devices/level0/level0-driver.cc b/lib/CL/devices/level0/level0-driver.cc
index 3388d0d7d00aada26e91d09141efe076de6b964d..dd08fcb07c71d514be099f0188fbdd562ae47f07 100644
--- a/lib/CL/devices/level0/level0-driver.cc
+++ b/lib/CL/devices/level0/level0-driver.cc
@@ -2060,7 +2060,8 @@ bool Level0Device::setupDeviceProperties(bool HasIPVersionExt) {
 #else
     ClDev->supported_spir_v_versions = "SPIR-V_1.2 SPIR-V_1.1 SPIR-V_1.0";
 #endif
-    ClDev->on_host_queue_props = CL_QUEUE_PROFILING_ENABLE;
+    ClDev->on_host_queue_props
+        = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE;
     ClDev->version_of_latest_passed_cts = "v2000-12-31-01";
   }
 
@@ -2248,6 +2249,15 @@ bool Level0Device::setupModuleProperties(bool &SupportsInt64Atomics,
   SupportsInt64Atomics = (ModuleProperties.flags &
                           ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS) != 0u;
 #endif
+  // clear flags set in setupDeviceProperties
+  if (ClDev->double_fp_config == 0) {
+    ClDev->preferred_vector_width_double = 0;
+    ClDev->native_vector_width_double = 0;
+  }
+  if (ClDev->half_fp_config == 0) {
+    ClDev->preferred_vector_width_half = 0;
+    ClDev->native_vector_width_half = 0;
+  }
 
   KernelUUID = ModuleProperties.nativeKernelSupported;
   SupportsDP4A = (ModuleProperties.flags & ZE_DEVICE_MODULE_FLAG_DP4A) > 0;
diff --git a/lib/CL/devices/proxy/pocl_proxy.c b/lib/CL/devices/proxy/pocl_proxy.c
index 233164569a2d9988b703b5bf6c1bc70a626bfe9b..a9e3f461ebcc78321fbafade7416ea8926a11b95 100644
--- a/lib/CL/devices/proxy/pocl_proxy.c
+++ b/lib/CL/devices/proxy/pocl_proxy.c
@@ -515,7 +515,6 @@ pocl_proxy_get_device_info (cl_device_id device, proxy_device_data_t *d)
   DIstring (profile, CL_DEVICE_PROFILE);
 
   // TODO queue properties
-  device->queue_properties = CL_QUEUE_PROFILING_ENABLE;
   DIbool (available[0], CL_DEVICE_AVAILABLE);
   DIbool (compiler_available, CL_DEVICE_COMPILER_AVAILABLE);
   DIbool (linker_available, CL_DEVICE_LINKER_AVAILABLE);
diff --git a/lib/CL/devices/remote/communication.c b/lib/CL/devices/remote/communication.c
index 3f3faf6b3887e0d1565b524a08489e97ba0daddd..3a5b1b30dbfd81cb5b810206f5c43c0ccdc133c0 100644
--- a/lib/CL/devices/remote/communication.c
+++ b/lib/CL/devices/remote/communication.c
@@ -1987,7 +1987,7 @@ pocl_network_fetch_devinfo (cl_device_id device,
 
   device->profile
       = (devinfo->full_profile ? "FULL_PROFILE" : "EMBEDDED_PROFILE");
-  device->queue_properties = CL_QUEUE_PROFILING_ENABLE;
+  device->on_host_queue_props = CL_QUEUE_PROFILING_ENABLE;
   device->compiler_available = 1;
   device->linker_available = 1;
 
diff --git a/lib/CL/devices/tce/ttasim/ttasim.cc b/lib/CL/devices/tce/ttasim/ttasim.cc
index cf74c974904c714e1e2b55565c0651816aec033c..e9305e6cbcb6b154d16ee400784527a7fb18afd0 100644
--- a/lib/CL/devices/tce/ttasim/ttasim.cc
+++ b/lib/CL/devices/tce/ttasim/ttasim.cc
@@ -700,7 +700,7 @@ pocl_ttasim_init (unsigned j, cl_device_id dev, const char* parameters)
   dev->spmd = CL_FALSE;
   dev->run_workgroup_pass = CL_TRUE;
   dev->execution_capabilities = CL_EXEC_KERNEL;
-  dev->queue_properties = CL_QUEUE_PROFILING_ENABLE;
+  dev->on_host_queue_props = CL_QUEUE_PROFILING_ENABLE;
   dev->vendor = "TTA-Based Co-design Environment";
   dev->profile = "EMBEDDED_PROFILE";
   dev->extensions = TCE_DEVICE_EXTENSIONS;
diff --git a/lib/CL/devices/vulkan/pocl-vulkan.c b/lib/CL/devices/vulkan/pocl-vulkan.c
index d3b3095e934338efbb2b5d8e35e82e7a7979f94c..1dcacf9d8e3da93c4ca6dd8d1d9aa4f5e10dc88d 100644
--- a/lib/CL/devices/vulkan/pocl-vulkan.c
+++ b/lib/CL/devices/vulkan/pocl-vulkan.c
@@ -1460,6 +1460,7 @@ pocl_vulkan_init (unsigned j, cl_device_id dev, const char *parameters)
     strcat (extensions, " cl_khr_fp64");
 
   dev->extensions = strdup (extensions);
+  dev->on_host_queue_props = CL_QUEUE_PROFILING_ENABLE;
 
   if (dev->vendor_id == 0x10de)
     {
diff --git a/lib/CL/pocl_cl.h b/lib/CL/pocl_cl.h
index 25856528ac031da79966deb7e82d2775979f5a26..7550aa7f22c7140d9eaec9b26460a190aab5bcac 100644
--- a/lib/CL/pocl_cl.h
+++ b/lib/CL/pocl_cl.h
@@ -1102,7 +1102,6 @@ struct _cl_device_id {
    * _clang_opencl.h. For most drivers, this should default to CL_FALSE. */
   cl_bool use_only_clang_opencl_headers;
   cl_device_exec_capabilities execution_capabilities;
-  cl_command_queue_properties queue_properties;
   cl_platform_id platform;
   cl_uint max_sub_devices;
   size_t num_partition_properties;
diff --git a/lib/CL/pocl_fill_memobj.c b/lib/CL/pocl_fill_memobj.c
index 59817a9161116112bd2c9e58e13d6f7875d3fec0..ff4c734c496b998b2cf4e0d820b2f8c2e0b0ac45 100644
--- a/lib/CL/pocl_fill_memobj.c
+++ b/lib/CL/pocl_fill_memobj.c
@@ -207,7 +207,7 @@ pocl_fill_image_common (cl_command_buffer_khr command_buffer,
       else
         {
           return POname (clCommandFillBufferKHR) (
-              command_buffer, command_queue, image->buffer, fill_pattern, px,
+              command_buffer, command_queue, NULL, image->buffer, fill_pattern, px,
               origin[0] * px, region[0] * px, num_items_in_wait_list,
               sync_point_wait_list, sync_point, mutable_handle);
         }
diff --git a/lib/CL/pocl_intfn.h b/lib/CL/pocl_intfn.h
index af771c8e2d51f57187077c8768d144ec32fbe778..6c0c71d728d3f93d9025bc67ca806874707e2158 100644
--- a/lib/CL/pocl_intfn.h
+++ b/lib/CL/pocl_intfn.h
@@ -134,6 +134,7 @@ POdeclsym(clSVMAlloc)
 POdeclsym(clSetKernelArgSVMPointer)
 POdeclsym(clSetKernelExecInfo)
 POdeclsym(clCreateCommandQueueWithProperties)
+POdeclsym(clSetCommandQueueProperty)
 POdeclsym(clCreateFromGLBuffer)
 POdeclsym(clCreateFromGLTexture)
 POdeclsym(clCreateFromGLTexture2D)
diff --git a/lib/CL/pocl_ndrange_kernel.c b/lib/CL/pocl_ndrange_kernel.c
index 7931a832fe5838cb9183c99bb19697c605d69c9b..b6fa3b8ca78a00b2869ede7f186a3a41c3c3a243 100644
--- a/lib/CL/pocl_ndrange_kernel.c
+++ b/lib/CL/pocl_ndrange_kernel.c
@@ -388,13 +388,13 @@ pocl_kernel_copy_args (cl_kernel kernel,
 
 static int
 process_command_ndrange_properties (
-  const cl_ndrange_kernel_command_properties_khr *properties)
+  const cl_command_properties_khr *properties)
 {
   if (properties == NULL)
     return CL_SUCCESS;
 
   cl_uint num_properties = 0;
-  const cl_ndrange_kernel_command_properties_khr *key = NULL;
+  const cl_command_properties_khr *key = NULL;
   for (key = properties; *key != 0; key += 2)
     num_properties += 1;
   POCL_RETURN_ERROR_ON ((num_properties == 0), CL_INVALID_VALUE,
@@ -413,7 +413,7 @@ process_command_ndrange_properties (
         default:
           POCL_RETURN_ERROR_ON (1, CL_INVALID_VALUE,
                                 "Unknown property value in "
-                                "cl_ndrange_kernel_command_properties_khr\n");
+                                "cl_command_properties_khr\n");
         }
     }
 
@@ -421,10 +421,9 @@ process_command_ndrange_properties (
 }
 
 cl_int
-pocl_record_ndrange_kernel (
-  cl_command_buffer_khr command_buffer,
+pocl_record_ndrange_kernel (cl_command_buffer_khr command_buffer,
   cl_command_queue command_queue,
-  const cl_ndrange_kernel_command_properties_khr *properties,
+  const cl_command_properties_khr *properties,
   cl_kernel kernel,
   struct pocl_argument *src_arguments,
   cl_uint work_dim,
@@ -473,7 +472,7 @@ cl_int
 pocl_ndrange_kernel_common (
   cl_command_buffer_khr command_buffer,
   cl_command_queue command_queue,
-  const cl_ndrange_kernel_command_properties_khr *properties,
+  const cl_command_properties_khr *properties,
   cl_kernel kernel,
   struct pocl_argument *src_arguments,
   cl_uint work_dim,
diff --git a/lib/CL/pocl_shared.h b/lib/CL/pocl_shared.h
index 6f3434629f966533a6b63919a35a2b93d98a8bb7..65320d88ec854cc965dbbc62982d2b6b0db3e35b 100644
--- a/lib/CL/pocl_shared.h
+++ b/lib/CL/pocl_shared.h
@@ -87,7 +87,7 @@ cl_int pocl_kernel_copy_args (cl_kernel kernel,
 cl_int pocl_ndrange_kernel_common (
   cl_command_buffer_khr command_buffer,
   cl_command_queue command_queue,
-  const cl_ndrange_kernel_command_properties_khr *properties,
+  const cl_command_properties_khr *properties,
   cl_kernel kernel,
   struct pocl_argument *src_arguments,
   cl_uint work_dim,
@@ -104,7 +104,7 @@ cl_int pocl_ndrange_kernel_common (
 cl_int pocl_record_ndrange_kernel (
   cl_command_buffer_khr command_buffer,
   cl_command_queue command_queue,
-  const cl_ndrange_kernel_command_properties_khr *properties,
+    const cl_command_properties_khr *properties,
   cl_kernel kernel,
   struct pocl_argument *src_arguments,
   cl_uint work_dim,
diff --git a/lib/CL/pocl_util.c b/lib/CL/pocl_util.c
index 5e5bc86f849cd195d3d2b1a8c5b115c4d08039f8..7e3572e8617e5eb3d70a2503d9105533037eb81e 100644
--- a/lib/CL/pocl_util.c
+++ b/lib/CL/pocl_util.c
@@ -1602,11 +1602,12 @@ pocl_setup_context (cl_context context)
       if (dev->ops->init_context)
         dev->ops->init_context (dev, context);
 
-      context->default_queues[i] = POname (clCreateCommandQueue) (
-          context, dev,
-          (CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_HIDDEN
-           | CL_QUEUE_PROFILING_ENABLE),
-          &err);
+      cl_command_queue_properties props
+        = CL_QUEUE_HIDDEN | CL_QUEUE_PROFILING_ENABLE;
+      if (dev->on_host_queue_props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+        props |= CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
+      context->default_queues[i]
+        = POname (clCreateCommandQueue) (context, dev, props, &err);
       assert (err == CL_SUCCESS);
       assert (context->default_queues[i]);
     }
diff --git a/poclu/misc.c b/poclu/misc.c
index fc9383a85215a2aa0169d13b98a234d21772c3dd..7206fe06e56dc50987c7575972bec2b7f35447bf 100644
--- a/poclu/misc.c
+++ b/poclu/misc.c
@@ -110,10 +110,13 @@ poclu_get_any_device2 (cl_context *context, cl_device_id *device,
 }
 
 cl_int
-poclu_get_multiple_devices (cl_platform_id *platform, cl_context *context,
-                            cl_char include_custom_dev, cl_uint *num_devices,
-                            cl_device_id **devices, cl_command_queue **queues,
-                            int ooo_queues)
+poclu_get_multiple_devices (cl_platform_id *platform,
+                            cl_context *context,
+                            cl_char include_custom_dev,
+                            cl_uint *num_devices,
+                            cl_device_id **devices,
+                            cl_command_queue **queues,
+                            cl_command_queue_properties optional_props)
 {
   cl_int err;
   cl_uint num_dev_all = 0;
@@ -184,11 +187,13 @@ poclu_get_multiple_devices (cl_platform_id *platform, cl_context *context,
   if (err != CL_SUCCESS)
     goto ERROR;
 
-  cl_command_queue_properties props = CL_QUEUE_PROFILING_ENABLE;
-  if (ooo_queues)
-    props |= CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
   for (i = 0; i < *num_devices; ++i)
     {
+      cl_command_queue_properties dev_props = 0;
+      CHECK_CL_ERROR (clGetDeviceInfo (devs[i],
+                                       CL_DEVICE_QUEUE_ON_HOST_PROPERTIES,
+                                       sizeof (dev_props), &dev_props, NULL));
+      cl_command_queue_properties props = dev_props & optional_props;
       ques[i] = clCreateCommandQueue (*context, devs[i], props, &err);
       if (err != CL_SUCCESS)
         goto ERROR;
@@ -200,6 +205,11 @@ poclu_get_multiple_devices (cl_platform_id *platform, cl_context *context,
 
 ERROR:
   free (devs);
+  for (i = 0; i < *num_devices; ++i)
+    {
+      if (ques[i])
+        clReleaseCommandQueue (ques[i]);
+    }
   free (ques);
   return err;
 }
diff --git a/poclu/poclu.h b/poclu/poclu.h
index 7cbc90e165ced435938f8403d2c07995f7dc398c..0b98a37f1854921302da85bccc4bedc0375a0d0e 100644
--- a/poclu/poclu.h
+++ b/poclu/poclu.h
@@ -220,10 +220,14 @@ POCLU_API cl_int POCLU_CALL poclu_get_any_device (cl_context *context,
  * @return CL_SUCCESS on success, or a descriptive OpenCL error code upon
  * failure.
  */
-POCLU_API cl_int POCLU_CALL poclu_get_multiple_devices (
-    cl_platform_id *platform, cl_context *context, cl_char include_custom_dev,
-    cl_uint *num_devices, cl_device_id **devices, cl_command_queue **queues,
-    int ooo_queues);
+POCLU_API cl_int POCLU_CALL
+poclu_get_multiple_devices (cl_platform_id *platform,
+                            cl_context *context,
+                            cl_char include_custom_dev,
+                            cl_uint *num_devices,
+                            cl_device_id **devices,
+                            cl_command_queue **queues,
+                            cl_command_queue_properties optional_props);
 
 /**
  * \brief read the contents of a file.
diff --git a/tests/runtime/test_buffer_migration.c b/tests/runtime/test_buffer_migration.c
index b7f41c9b15e6379cd05d38801f0ba93601413035..21299ef8a353fa5f7fafc14fe8ebb4fc49f87059 100644
--- a/tests/runtime/test_buffer_migration.c
+++ b/tests/runtime/test_buffer_migration.c
@@ -55,7 +55,8 @@ main (int argc, char **argv)
   cl_event ev1, ev2;
 
   err = poclu_get_multiple_devices (&platform, &context, 0, &num_devices,
-                                    &devices, &queues, 1);
+                                    &devices, &queues,
+                                    CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
   CHECK_OPENCL_ERROR_IN ("poclu_get_multiple_devices");
 
   printf ("NUM DEVICES: %u \n", num_devices);
diff --git a/tests/runtime/test_cl_pocl_content_size_migration.c b/tests/runtime/test_cl_pocl_content_size_migration.c
index 940b46fa289a156dedc6d33b89e81c8cee8fb8f5..c8907ae057770f17098dff4405f5203c299fb0a1 100644
--- a/tests/runtime/test_cl_pocl_content_size_migration.c
+++ b/tests/runtime/test_cl_pocl_content_size_migration.c
@@ -63,7 +63,7 @@ main (void)
   uint64_t content_size;
 
   poclu_get_multiple_devices (&platform, &context, CL_FALSE, &num_devices,
-                              &devices, &queues, CL_FALSE);
+                              &devices, &queues, 0);
   if (num_devices < 2)
     {
       printf ("Not enough devices (2 required), skipping");
diff --git a/tests/runtime/test_command_buffer.c b/tests/runtime/test_command_buffer.c
index fca17585d648993c0f31f85e17e496b266a6c23a..61009724d3648439aa2635d23ff41d016be3076c 100644
--- a/tests/runtime/test_command_buffer.c
+++ b/tests/runtime/test_command_buffer.c
@@ -118,8 +118,13 @@ main (int _argc, char **_argv)
   CHECK_CL_ERROR (
       clSetKernelArg (kernel, 2, sizeof (buffer_res), &buffer_res));
 
-  cl_command_queue command_queue = clCreateCommandQueue (
-      context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
+  cl_command_queue_properties props = 0;
+  CHECK_CL_ERROR (clGetDeviceInfo (device, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES,
+                                   sizeof (props), &props, NULL));
+  if (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+    props = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
+  cl_command_queue command_queue
+    = clCreateCommandQueue (context, device, props, &error);
   CHECK_CL_ERROR (error);
 
   cl_command_buffer_khr command_buffer
@@ -141,7 +146,7 @@ main (int _argc, char **_argv)
     {
       cl_sync_point_khr copy_sync_points[2];
       CHECK_CL_ERROR (ext.clCommandCopyBufferKHR (
-          command_buffer, NULL, buffer_src1, buffer_tile1,
+          command_buffer, NULL, NULL, buffer_src1, buffer_tile1,
           tile_index * tile_size, 0, tile_size, tile_sync_point ? 1 : 0,
           tile_sync_point ? &tile_sync_point : NULL, &copy_sync_points[0],
           NULL));
@@ -151,7 +156,7 @@ main (int _argc, char **_argv)
       size_t dst_origin[3] = { 0, 0, 0 };
       size_t tile_region[3] = { 8 * sizeof (cl_int), 8, 1 };
       CHECK_CL_ERROR (ext.clCommandCopyBufferRectKHR (
-          command_buffer, NULL, buffer_src2, buffer_tile2, src_origin,
+          command_buffer, NULL, NULL, buffer_src2, buffer_tile2, src_origin,
           dst_origin, tile_region, tile_region[0], 0, tile_region[0], 0,
           tile_sync_point ? 1 : 0, tile_sync_point ? &tile_sync_point : NULL,
           &copy_sync_points[1], NULL));
@@ -163,24 +168,24 @@ main (int _argc, char **_argv)
 
       cl_sync_point_khr res_copy_sync_point;
       CHECK_CL_ERROR (ext.clCommandCopyBufferKHR (
-          command_buffer, NULL, buffer_res, buffer_dst, 0,
+          command_buffer, NULL, NULL, buffer_res, buffer_dst, 0,
           tile_index * tile_size, tile_size, 1, &nd_sync_point,
           &res_copy_sync_point, NULL));
 
       char zero = 0;
       cl_sync_point_khr fill_sync_points[2];
       CHECK_CL_ERROR (ext.clCommandFillBufferKHR (
-          command_buffer, NULL, buffer_tile1, &zero, sizeof (zero), 0,
+          command_buffer, NULL, NULL, buffer_tile1, &zero, sizeof (zero), 0,
           tile_size, 1, &nd_sync_point, &fill_sync_points[0], NULL));
       CHECK_CL_ERROR (ext.clCommandFillBufferKHR (
-          command_buffer, NULL, buffer_tile2, &zero, sizeof (zero), 0,
+          command_buffer, NULL, NULL, buffer_tile2, &zero, sizeof (zero), 0,
           tile_size, 1, &nd_sync_point, &fill_sync_points[1], NULL));
 
       cl_sync_point_khr barrier_deps[4]
           = { nd_sync_point, res_copy_sync_point, fill_sync_points[0],
               fill_sync_points[1] };
       CHECK_CL_ERROR (ext.clCommandBarrierWithWaitListKHR (
-          command_buffer, NULL, 4, barrier_deps, &tile_sync_point, NULL));
+          command_buffer, NULL, NULL, 4, barrier_deps, &tile_sync_point, NULL));
     }
 
   CHECK_CL_ERROR (ext.clFinalizeCommandBufferKHR (command_buffer));
diff --git a/tests/runtime/test_command_buffer_images.c b/tests/runtime/test_command_buffer_images.c
index 7310600dc6e7fab58791960622b9e0d35a40e685..a647292ed502d916b7052bc7488ef48a6e6cc4a3 100644
--- a/tests/runtime/test_command_buffer_images.c
+++ b/tests/runtime/test_command_buffer_images.c
@@ -116,8 +116,13 @@ main (int _argc, char **_argv)
                                &img_desc, NULL, &error);
   CHECK_CL_ERROR (error);
 
-  cl_command_queue command_queue = clCreateCommandQueue (
-      context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
+  cl_command_queue_properties props = 0;
+  CHECK_CL_ERROR (clGetDeviceInfo (device, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES,
+                                   sizeof (props), &props, NULL));
+  if (props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+    props = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
+  cl_command_queue command_queue
+    = clCreateCommandQueue (context, device, props, &error);
   CHECK_CL_ERROR (error);
 
   /**** Command buffer creation ****/
@@ -129,21 +134,21 @@ main (int _argc, char **_argv)
     cl_sync_point_khr fill_syncpt;
     size_t origin[3] = { 0, 0, 0 };
     size_t region[3] = { img_width, img_height, img_depth };
-    CHECK_CL_ERROR (ext.clCommandFillImageKHR (command_buffer, NULL, img1,
+    CHECK_CL_ERROR (ext.clCommandFillImageKHR (command_buffer, NULL, NULL, img1,
                                                fill_pixel, origin, region, 0,
                                                NULL, &fill_syncpt, NULL));
     cl_sync_point_khr buf2img_syncpt;
     CHECK_CL_ERROR (ext.clCommandCopyBufferToImageKHR (
-        command_buffer, NULL, buffer, img2, 0, origin, region, 0, NULL,
+        command_buffer, NULL, NULL, buffer, img2, 0, origin, region, 0, NULL,
         &buf2img_syncpt, NULL));
     cl_sync_point_khr img2img_syncpt;
     cl_sync_point_khr img2img_deps[2] = { fill_syncpt, buf2img_syncpt };
     CHECK_CL_ERROR (ext.clCommandCopyImageKHR (
-        command_buffer, NULL, img2, img1, img2img_origin, img2img_origin,
+        command_buffer, NULL, NULL, img2, img1, img2img_origin, img2img_origin,
         img2img_region, 2, img2img_deps, &img2img_syncpt, NULL));
 
     CHECK_CL_ERROR (ext.clCommandCopyImageToBufferKHR (
-        command_buffer, NULL, img1, buffer, origin, region, 0, 1,
+        command_buffer, NULL, NULL, img1, buffer, origin, region, 0, 1,
         &img2img_syncpt, NULL, NULL));
   }
 
diff --git a/tests/runtime/test_command_buffer_multi_device.c b/tests/runtime/test_command_buffer_multi_device.c
index 36bc3ecf14ea423eb00d41263d6c9d88ac8bfc3e..67f06e8e70073f2b5145fce5b042ea718e93ac47 100644
--- a/tests/runtime/test_command_buffer_multi_device.c
+++ b/tests/runtime/test_command_buffer_multi_device.c
@@ -232,7 +232,7 @@ main (int _argc, char **_argv)
     {
       cl_sync_point_khr copy_sync_points[2];
       CHECK_CL_ERROR (ext.clCommandCopyBufferKHR (
-        command_buffer, queues[tile_index % num_devices], buffer_src1,
+        command_buffer, queues[tile_index % num_devices], NULL, buffer_src1,
         buffer_tile1, tile_index * tile_size, 0, tile_size,
         tile_sync_point ? 1 : 0, tile_sync_point ? &tile_sync_point : NULL,
         &copy_sync_points[0], NULL));
@@ -242,7 +242,7 @@ main (int _argc, char **_argv)
       size_t dst_origin[3] = { 0, 0, 0 };
       size_t tile_region[3] = { 8 * sizeof (cl_int), 8, 1 };
       CHECK_CL_ERROR (ext.clCommandCopyBufferRectKHR (
-        command_buffer, queues[tile_index % num_devices], buffer_src2,
+        command_buffer, queues[tile_index % num_devices], NULL, buffer_src2,
         buffer_tile2, src_origin, dst_origin, tile_region, tile_region[0], 0,
         tile_region[0], 0, tile_sync_point ? 1 : 0,
         tile_sync_point ? &tile_sync_point : NULL, &copy_sync_points[1],
@@ -256,18 +256,18 @@ main (int _argc, char **_argv)
 
       cl_sync_point_khr res_copy_sync_point;
       CHECK_CL_ERROR (ext.clCommandCopyBufferKHR (
-        command_buffer, queues[tile_index % num_devices], buffer_res,
+        command_buffer, queues[tile_index % num_devices], NULL, buffer_res,
         buffer_dst, 0, tile_index * tile_size, tile_size, 1, &nd_sync_point,
         &res_copy_sync_point, NULL));
 
       char zero = 0;
       cl_sync_point_khr fill_sync_points[2];
       CHECK_CL_ERROR (ext.clCommandFillBufferKHR (
-        command_buffer, queues[tile_index % num_devices], buffer_tile1, &zero,
+        command_buffer, queues[tile_index % num_devices], NULL, buffer_tile1, &zero,
         sizeof (zero), 0, tile_size, 1, &nd_sync_point, &fill_sync_points[0],
         NULL));
       CHECK_CL_ERROR (ext.clCommandFillBufferKHR (
-        command_buffer, queues[tile_index % num_devices], buffer_tile2, &zero,
+        command_buffer, queues[tile_index % num_devices], NULL, buffer_tile2, &zero,
         sizeof (zero), 0, tile_size, 1, &nd_sync_point, &fill_sync_points[1],
         NULL));
 
@@ -275,7 +275,7 @@ main (int _argc, char **_argv)
         = { nd_sync_point, res_copy_sync_point, fill_sync_points[0],
             fill_sync_points[1] };
       CHECK_CL_ERROR (ext.clCommandBarrierWithWaitListKHR (
-        command_buffer, queues[tile_index % num_devices], 4, barrier_deps,
+        command_buffer, queues[tile_index % num_devices], NULL, 4, barrier_deps,
         &tile_sync_point, NULL));
     }