Merge remote-tracking branch 'origin/sycl' into sycl_bundle_multi_device_core_version

againull · againull · commit 0198f0e8269c · 2024-10-01T08:14:23.000-07:00
diff --git a/.github/workflows/sycl-linux-precommit.yml b/.github/workflows/sycl-linux-precommit.yml
@@ -83,22 +83,22 @@ jobs:
           - name: Intel
             runner: '["Linux", "gen12"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu;opencl:cpu
             reset_intel_gpu: true
             extra_lit_opts: --param gpu-intel-gen12=True
           - name: E2E tests on Intel Arc A-Series Graphics
             runner: '["Linux", "arc"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu
             reset_intel_gpu: true
             extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
             env: '{"LIT_FILTER":${{ needs.determine_arc_tests.outputs.arc_tests }} }'
           - name: E2E tests with dev igc on Intel Arc A-Series Graphics
             runner: '["Linux", "arc"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:devigc
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu;opencl:gpu
             reset_intel_gpu: true
             extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
diff --git a/.github/workflows/sycl-nightly.yml b/.github/workflows/sycl-nightly.yml
@@ -53,7 +53,7 @@ jobs:
           - name: Intel L0 GPU
             runner: '["Linux", "gen12"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: level_zero:gpu
             reset_intel_gpu: true
             tests_selector: e2e
@@ -62,7 +62,7 @@ jobs:
           - name: Intel OCL GPU
             runner: '["Linux", "gen12"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest
-            image_options: -u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN
+            image_options: -u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN
             target_devices: opencl:gpu
             reset_intel_gpu: true
             tests_selector: e2e
diff --git a/.github/workflows/sycl-post-commit.yml b/.github/workflows/sycl-post-commit.yml
@@ -87,7 +87,7 @@ jobs:
       name: ${{ matrix.name }}
       runner: ${{ matrix. runner }}
       image: ${{ matrix.image || 'ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest' }}
-      image_options: ${{ matrix.image_options || '-u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN' }}
+      image_options: ${{ matrix.image_options || '-u 1001 --device=/dev/dri -v /dev/dri/by-path:/dev/dri/by-path --privileged --cap-add SYS_ADMIN' }}
       target_devices: ${{ matrix.target_devices || 'level_zero:gpu' }}
       reset_intel_gpu: ${{ matrix.reset_intel_gpu }}
 
diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake
@@ -117,7 +117,7 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/againull/unified-runtime")
-  set(UNIFIED_RUNTIME_TAG 88f2f3b89d0e0e92c552432b2f8ef4a271c7ab32)
+  set(UNIFIED_RUNTIME_TAG 550df8a622a3bae8113d8e8d9c6deafba0464a87)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need
diff --git a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp
@@ -19,16 +19,16 @@ template <typename Set>
 inline constexpr indirectly_callable_key::value_t<Set> indirectly_callable_in;
 
 struct calls_indirectly_key {
-  template <typename First = void, typename... SetIds>
+  template <typename... SetIds>
   using value_t =
       sycl::ext::oneapi::experimental::property_value<calls_indirectly_key,
-                                                      First, SetIds...>;
+                                                      SetIds...>;
 };
 
 inline constexpr calls_indirectly_key::value_t<void> assume_indirect_calls;
 
-template <typename First, typename... Rest>
-inline constexpr calls_indirectly_key::value_t<First, Rest...>
+template <typename... SetIds>
+inline constexpr calls_indirectly_key::value_t<SetIds...>
     assume_indirect_calls_to;
 
 template <> struct is_property_key<indirectly_callable_key> : std::true_type {};
@@ -60,16 +60,57 @@ struct PropertyMetaInfo<indirectly_callable_key::value_t<Set>> {
 #endif
 };
 
-template <typename First, typename... Rest>
-struct PropertyMetaInfo<calls_indirectly_key::value_t<First, Rest...>> {
-  static_assert(
-      sizeof...(Rest) == 0,
-      "assume_indirect_calls_to property only supports a single set for now");
+#ifdef __SYCL_DEVICE_ONLY__
+// Helper to concatenate several lists of characters into a single string.
+// Lists are separated from each other with comma within the resulting string.
+template <typename List, typename... Rest> struct ConcatenateCharsToStr;
+
+// Specialization for a single list
+template <char... Chars> struct ConcatenateCharsToStr<CharList<Chars...>> {
+  static constexpr char value[] = {Chars..., '\0'};
+};
+
+// Specialization for two lists
+template <char... Chars, char... CharsToAppend>
+struct ConcatenateCharsToStr<CharList<Chars...>, CharList<CharsToAppend...>>
+    : ConcatenateCharsToStr<CharList<Chars..., ',', CharsToAppend...>> {};
+
+// Specialization for the case when there are more than two lists
+template <char... Chars, char... CharsToAppend, typename... Rest>
+struct ConcatenateCharsToStr<CharList<Chars...>, CharList<CharsToAppend...>,
+                             Rest...>
+    : ConcatenateCharsToStr<CharList<Chars..., ',', CharsToAppend...>,
+                            Rest...> {};
+
+// Helper to convert type T to a list of characters representing the type (its
+// mangled name).
+template <typename T, size_t... Indices> struct StableNameToCharsHelper {
+  using chars = CharList<__builtin_sycl_unique_stable_name(T)[Indices]...>;
+};
+
+// Wrapper helper for the struct above
+template <typename T, typename Sequence> struct StableNameToChars;
+
+// Specialization of that wrapper helper which accepts sequence of integers
+template <typename T, size_t... Indices>
+struct StableNameToChars<T, std::integer_sequence<size_t, Indices...>>
+    : StableNameToCharsHelper<T, Indices...> {};
+
+// Creates a comma-separated string with unique stable names for each type in
+// Ts.
+template <typename... Ts>
+struct UniqueStableNameListStr
+    : ConcatenateCharsToStr<typename StableNameToChars<
+          Ts, std::make_index_sequence<__builtin_strlen(
+                  __builtin_sycl_unique_stable_name(Ts))>>::chars...> {};
+#endif // __SYCL_DEVICE_ONLY__
+
+template <typename... SetIds>
+struct PropertyMetaInfo<calls_indirectly_key::value_t<SetIds...>> {
   static constexpr const char *name = "calls-indirectly";
   static constexpr const char *value =
 #ifdef __SYCL_DEVICE_ONLY__
-      // FIXME: we should handle Rest... here as well
-      __builtin_sycl_unique_stable_name(First);
+      UniqueStableNameListStr<SetIds...>::value;
 #else
       "";
 #endif
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
@@ -704,17 +704,26 @@ bool device_impl::has(aspect Aspect) const {
     return CallSuccessful && Result != nullptr;
   }
   case aspect::ext_oneapi_graph: {
-    bool SupportsCommandBufferUpdate = false;
+    ur_device_command_buffer_update_capability_flags_t UpdateCapabilities;
     bool CallSuccessful =
         getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
-            MDevice, UR_DEVICE_INFO_COMMAND_BUFFER_UPDATE_SUPPORT_EXP,
-            sizeof(SupportsCommandBufferUpdate), &SupportsCommandBufferUpdate,
+            MDevice, UR_DEVICE_INFO_COMMAND_BUFFER_UPDATE_CAPABILITIES_EXP,
+            sizeof(UpdateCapabilities), &UpdateCapabilities,
             nullptr) == UR_RESULT_SUCCESS;
     if (!CallSuccessful) {
       return false;
     }
 
-    return has(aspect::ext_oneapi_limited_graph) && SupportsCommandBufferUpdate;
+    /* The kernel handle update capability is not yet required for the
+     * ext_oneapi_graph aspect */
+    ur_device_command_buffer_update_capability_flags_t RequiredCapabilities =
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_ARGUMENTS |
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_LOCAL_WORK_SIZE |
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_SIZE |
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET;
+
+    return has(aspect::ext_oneapi_limited_graph) &&
+           (UpdateCapabilities & RequiredCapabilities) == RequiredCapabilities;
   }
   case aspect::ext_oneapi_limited_graph: {
     bool SupportsCommandBuffers = false;
diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp
@@ -1474,6 +1474,7 @@ void exec_graph_impl::updateImpl(std::shared_ptr<node_impl> Node) {
     }
   }
 
+  UpdateDesc.hNewKernel = UrKernel;
   UpdateDesc.numNewMemObjArgs = MemobjDescs.size();
   UpdateDesc.pNewMemObjArgList = MemobjDescs.data();
   UpdateDesc.numNewPointerArgs = PtrDescs.size();
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
@@ -2538,7 +2538,7 @@ ur_result_t enqueueImpCommandBufferKernel(
   ur_result_t Res =
       Adapter->call_nocheck<UrApiKind::urCommandBufferAppendKernelLaunchExp>(
           CommandBuffer, UrKernel, NDRDesc.Dims, &NDRDesc.GlobalOffset[0],
-          &NDRDesc.GlobalSize[0], LocalSize, SyncPoints.size(),
+          &NDRDesc.GlobalSize[0], LocalSize, 0, nullptr, SyncPoints.size(),
           SyncPoints.size() ? SyncPoints.data() : nullptr, OutSyncPoint,
           OutCommand);
 
diff --git a/sycl/test-e2e/DeviceCodeSplit/grf.cpp b/sycl/test-e2e/DeviceCodeSplit/grf.cpp
@@ -14,18 +14,16 @@
 //   compiler option
 
 // REQUIRES: arch-intel_gpu_pvc
-// https://github.com/intel/llvm/issues/14826
-// XFAIL: *
 
-// RUN: %{build} -o %t.out
+// RUN: %{build} -Wno-error=deprecated-declarations -o %t.out
 // Don't use SYCL_UR_TRACE as the output from the L0 adapter logging interferes
 // with the regular UR traces we are checking.
 // RUN: env UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-NO-VAR
 // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-WITH-VAR
 // RUN: %{build} -DUSE_NEW_API=1 -o %t.out
 // RUN: env UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-NO-VAR
 // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-WITH-VAR
-// RUN: %{build} -DUSE_AUTO_GRF=1 -o %t.out
+// RUN: %{build} -DUSE_AUTO_GRF=1 -Wno-error=deprecated-declarations -o %t.out
 // RUN: env UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-AUTO-NO-VAR
 // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" UR_LOG_TRACING="level:info;output:stdout;flush:info" UR_ENABLE_LAYERS=UR_LAYER_TRACING %{run} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK,CHECK-AUTO-WITH-VAR
 // RUN: %{build} -DUSE_NEW_API=1 -DUSE_AUTO_GRF=1 -o %t.out
diff --git a/sycl/test-e2e/KernelAndProgram/target_register_alloc_mode.cpp b/sycl/test-e2e/KernelAndProgram/target_register_alloc_mode.cpp
@@ -1,8 +1,5 @@
 // REQUIRES: arch-intel_gpu_pvc
 
-// https://github.com/intel/llvm/issues/14826
-// XFAIL: arch-intel_gpu_pvc
-
 // RUN: %{build} -ftarget-register-alloc-mode=pvc:auto -o %t_with.out
 // RUN: %{build} -o %t_without.out
 // RUN: %{build} -ftarget-register-alloc-mode=pvc:default -o %t_default.out
@@ -11,7 +8,7 @@
 // RUN: env SYCL_UR_TRACE=2 %{run} %t_without.out 2>&1 | FileCheck %if system-windows %{ --implicit-check-not=-ze-intel-enable-auto-large-GRF-mode %} %else %{ --check-prefix=CHECK-OPT %} %s
 // RUN: env SYCL_UR_TRACE=2 %{run} %t_default.out 2>&1 | FileCheck --implicit-check-not=-ze-intel-enable-auto-large-GRF-mode %s
 
-// CHECK-OPT: ---> urProgramBuild(
+// CHECK-OPT: ---> urProgramBuildExp(
 // CHECK-SAME-OPT: -ze-intel-enable-auto-large-GRF-mode
 
 #include <sycl/detail/core.hpp>
diff --git a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp
@@ -51,9 +51,18 @@ class IncrementBy8 : public BaseIncrement {
   void increment(int *Data) override { *Data += 8 + Mod; }
 };
 
+struct SetIncBy16;
+class IncrementBy16 : public BaseIncrement {
+public:
+  IncrementBy16(int Mod, int /* unused */) : BaseIncrement(Mod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable_in<SetIncBy16>)
+  void increment(int *Data) override { *Data += 16 + Mod; }
+};
+
 int main() try {
-  using storage_t =
-      obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>;
+  using storage_t = obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4,
+                                  IncrementBy8, IncrementBy16>;
 
   storage_t HostStorage;
   sycl::buffer<storage_t> DeviceStorage(sycl::range{1});
@@ -66,8 +75,9 @@ int main() try {
   sycl::queue q(asyncHandler);
 
   // TODO: cover uses case when objects are passed through USM
-  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
-  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+  constexpr oneapi::properties props{
+      oneapi::assume_indirect_calls_to<void, SetIncBy16>};
+  for (unsigned TestCase = 0; TestCase < 5; ++TestCase) {
     int HostData = 42;
     int Data = HostData;
     sycl::buffer<int> DataStorage(&Data, sycl::range{1});
diff --git a/sycl/test/virtual-functions/calls-indirectly-ir.cpp b/sycl/test/virtual-functions/calls-indirectly-ir.cpp
@@ -11,16 +11,12 @@
 // CHECK: define {{.*}}KInt{{.*}} #[[#ATTR_SET_INT:]]
 // CHECK: define {{.*}}KVoid{{.*}} #[[#ATTR_SET_DEFAULT]]
 // CHECK: define {{.*}}KUserDefined{{.*}} #[[#ATTR_SET_USER_DEFINED:]]
-// TODO: update the check below
-// As of now calls_indirectly_property takes into account only the first
-// template argument ignoring the rest. This will be fixed in a follow-up
-// patches and the test should be updated to reflect that, because current
-// behavior is not correct.
-// CHECK-disabled: define {{.*}}KMultiple{{.*}} #[[#ATTR_SET_INT]]
+// CHECK: define {{.*}}KMultiple{{.*}} #[[#ATTR_SET_MULTIPLE:]]
 //
 // CHECK-DAG: attributes #[[#ATTR_SET_DEFAULT]] {{.*}} "calls-indirectly"="_ZTSv"
 // CHECK-DAG: attributes #[[#ATTR_SET_INT]] {{.*}} "calls-indirectly"="_ZTSi"
 // CHECK-DAG: attributes #[[#ATTR_SET_USER_DEFINED]] {{.*}} "calls-indirectly"="_ZTS12user_defined"
+// CHECK-DAG: attributes #[[#ATTR_SET_MULTIPLE]] {{.*}} "calls-indirectly"="_ZTSi,_ZTS12user_defined"
 
 #include <sycl/sycl.hpp>
 
@@ -45,17 +41,14 @@ int main() {
   oneapi::properties props_void{oneapi::assume_indirect_calls_to<void>};
   oneapi::properties props_user_defined{
       oneapi::assume_indirect_calls_to<user_defined>};
-  // assume_indirect_calls_to is currently limited to a single set, so this test
-  // is disabled.
-  // FIXME: re-enable once the restriction is lifted.
-  // oneapi::properties props_multiple{
-  //    oneapi::assume_indirect_calls_to<int, user_defined>};
+  oneapi::properties props_multiple{
+     oneapi::assume_indirect_calls_to<int, user_defined>};
 
   q.single_task<KEmpty>(props_empty, [=]() {});
   q.single_task<KInt>(props_int, [=]() {});
   q.single_task<KVoid>(props_void, [=]() {});
   q.single_task<KUserDefined>(props_user_defined, [=]() {});
-  // q.single_task<KMultiple>(props_multiple, [=]() {});
+  q.single_task<KMultiple>(props_multiple, [=]() {});
 
   return 0;
 }
diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp
@@ -52,17 +52,14 @@ int main() {
   oneapi::properties props_void{oneapi::assume_indirect_calls_to<void>};
   oneapi::properties props_int{oneapi::assume_indirect_calls_to<int>};
   oneapi::properties props_base{oneapi::assume_indirect_calls_to<Base>};
-  // assume_indirect_calls_to is currently limited to a single set, so this test
-  // is disabled.
-  // FIXME: re-enable once the restriction is lifted.
-  // oneapi::properties props_multiple{
-  //    oneapi::assume_indirect_calls_to<int, Base>};
+  oneapi::properties props_multiple{
+     oneapi::assume_indirect_calls_to<int, Base>};
 
   q.single_task(props_empty, [=]() {});
   q.single_task(props_void, [=]() {});
   q.single_task(props_int, [=]() {});
   q.single_task(props_base, [=]() {});
-  // q.single_task(props_multiple, [=]() {});
+  q.single_task(props_multiple, [=]() {});
 
   return 0;
 }
diff --git a/sycl/unittests/helpers/UrMock.hpp b/sycl/unittests/helpers/UrMock.hpp
@@ -200,14 +200,27 @@ inline ur_result_t mock_urDeviceGetInfo(void *pParams) {
   case UR_DEVICE_INFO_AVAILABLE:
   case UR_DEVICE_INFO_LINKER_AVAILABLE:
   case UR_DEVICE_INFO_COMPILER_AVAILABLE:
-  case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP:
-  case UR_DEVICE_INFO_COMMAND_BUFFER_UPDATE_SUPPORT_EXP: {
+  case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: {
     if (*params->ppPropValue)
       *static_cast<ur_bool_t *>(*params->ppPropValue) = true;
     if (*params->ppPropSizeRet)
       **params->ppPropSizeRet = sizeof(true);
     return UR_RESULT_SUCCESS;
   }
+  case UR_DEVICE_INFO_COMMAND_BUFFER_UPDATE_CAPABILITIES_EXP: {
+    if (*params->ppPropValue)
+      *static_cast<ur_device_command_buffer_update_capability_flags_t *>(
+          *params->ppPropValue) =
+          UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_ARGUMENTS |
+          UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_SIZE |
+          UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_LOCAL_WORK_SIZE |
+          UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET |
+          UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_HANDLE;
+    if (*params->ppPropSizeRet)
+      **params->ppPropSizeRet =
+          sizeof(ur_device_command_buffer_update_capability_flags_t);
+    return UR_RESULT_SUCCESS;
+  }
   // This mock GPU device has no sub-devices
   case UR_DEVICE_INFO_SUPPORTED_PARTITIONS: {
     if (*params->ppPropSizeRet) {

Original file line number	Diff line number	Diff line change
`@@ -1474,6 +1474,7 @@ void exec_graph_impl::updateImpl(std::shared_ptr<node_impl> Node) {`
`1474`	`1474`	`}`
`1475`	`1475`	`}`
`1476`	`1476`
	`1477`	`+ UpdateDesc.hNewKernel = UrKernel;`
`1477`	`1478`	`UpdateDesc.numNewMemObjArgs = MemobjDescs.size();`
`1478`	`1479`	`UpdateDesc.pNewMemObjArgList = MemobjDescs.data();`
`1479`	`1480`	`UpdateDesc.numNewPointerArgs = PtrDescs.size();`