Move global/constant memory init() outside of cgh lambda

Not allowed to nest q.submit for the memcpy when using buffer support.

Author: joeatodd

sycl/test-e2e/syclcompat/memory/global_memory_usmnone.cpp

@@ -117,12 +117,12 @@ void test1(volatile int *acc_d1, int *acc_d2, int const *c1, int c2) {
 }
 
 int main() {
+  d1_a.init();
+  d2_a.init();
+  c1_a.init();
+  c2_a.init();
   syclcompat::get_default_queue().submit(
     [&](sycl::handler &cgh) {
-      d1_a.init();
-      d2_a.init();
-      c1_a.init();
-      c2_a.init();
       auto d1_acc = d1_a.get_access(cgh);
       auto d2_acc = d2_a.get_access(cgh);
       auto c1_acc = c1_a.get_access(cgh);
@@ -136,10 +136,10 @@ int main() {
                   c2_acc);
           });
     });
+  c3_a.init();
+  c4_a.init();
   syclcompat::get_default_queue().submit(
     [&](sycl::handler &cgh) {
-      c3_a.init();
-      c4_a.init();
       auto c3_acc = c3_a.get_access(cgh);
       auto c4_acc = c4_a.get_access(cgh);
       cgh.parallel_for<syclcompat_kernel_name<class kernel_test2>>(
@@ -150,10 +150,10 @@ int main() {
     });
 
   sycl::queue *q = syclcompat::get_current_device().create_queue();
+  d3_a.init(*q);
+  d4_a.init(*q);
   q->submit(
     [&](sycl::handler &cgh) {
-      d3_a.init(*q);
-      d4_a.init(*q);
       auto d3_acc = d3_a.get_access(cgh);
       auto d4_acc = d4_a.get_access(cgh);
       cgh.parallel_for<syclcompat_kernel_name<class kernel_test3>>(