[SYCL] Fixed crash when a non-nullptr_t 0x0 value is passed to the set_final_data method (#1004)

Signed-off-by: Sergey Kanaev <[email protected]>

Author: s-kanaev

sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

@@ -165,11 +165,14 @@ template <typename AllocatorT> class SYCLMemObjT : public SYCLMemObjI {
 
   template <typename Destination>
   EnableIfOutputPointerT<Destination> set_final_data(Destination FinalData) {
-    MUploadDataFunctor = [this, FinalData]() {
-      EventImplPtr Event = updateHostMemory(FinalData);
-      if (Event)
-        Event->wait(Event);
-    };
+    if (!FinalData)
+      MUploadDataFunctor = nullptr;
+    else
+      MUploadDataFunctor = [this, FinalData]() {
+        EventImplPtr Event = updateHostMemory(FinalData);
+        if (Event)
+          Event->wait(Event);
+      };
   }
 
   template <typename Destination>

sycl/test/basic_tests/image.cpp

@@ -73,6 +73,31 @@ int main() {
     });
   }
 
+  // image with write accessor to it in kernel
+  {
+    int NX = 32;
+    int NY = 32;
+
+    sycl::image<2> Img(sycl::image_channel_order::rgba,
+                       sycl::image_channel_type::fp32,
+                       sycl::range<2>(NX, NY));
+
+    sycl::queue Q;
+    Q.submit([&](sycl::handler &CGH) {
+        auto ImgAcc = Img.get_access<sycl::float4, sycl::access::mode::write>(
+              CGH);
+
+        sycl::nd_range<2> Rng(sycl::range<2>(NX, NY), sycl::range<2>(16, 16));
+
+        CGH.parallel_for<class sample>(Rng, [=](sycl::nd_item<2> Item) {
+            sycl::id<2> Idx = Item.get_global_id();
+            sycl::float4 C(0.5f, 0.5f, 0.2f, 1.0f);
+            ImgAcc.write(sycl::int2(Idx[0], Idx[1]), C);
+        });
+
+    }).wait_and_throw();
+  }
+
   std::cout << "Success" << std::endl;
   return 0;
 }