Remove useless wait from SYCL backend

The default queue is in order so many synchronization with the host are
useless.

Author: s-Nick

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -355,7 +355,7 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer,
         if (padded_size > original_size && tensor->view_src == nullptr) {
             SYCL_CHECK(CHECK_TRY_ERROR(ctx->stream->memset(
                 (char *)tensor->data + original_size, 0,
-                padded_size - original_size).wait()));
+                padded_size - original_size)));
         }
     }
     return GGML_STATUS_SUCCESS;
@@ -489,7 +489,7 @@ static void ggml_backend_sycl_buffer_clear(ggml_backend_buffer_t buffer,
 
     SYCL_CHECK(CHECK_TRY_ERROR((*stream)
                                     .memset(ctx->dev_ptr, value, buffer->size)
-                                    .wait()));
+                                    ));
 }
 catch (sycl::exception const &exc) {
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
@@ -511,7 +511,6 @@ static void ggml_backend_sycl_buffer_memset_tensor(ggml_backend_buffer_t buffer,
     }
     void * target_ptr = static_cast<char *>(tensor->data) + offset;
     SYCL_CHECK(CHECK_TRY_ERROR((*stream).memset(target_ptr, value, size)));
-    SYCL_CHECK(CHECK_TRY_ERROR((*stream).wait()));
 }
 
 static void ggml_backend_sycl_buffer_reset(ggml_backend_buffer_t buffer) {
@@ -833,7 +832,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer,
             SYCL_CHECK(CHECK_TRY_ERROR(
                 (*stream)
                     .memset(buf + original_size, 0, size - original_size)
-                    .wait()));
+                    ));
         }
 
         extra->data_device[i] = buf;
@@ -901,7 +900,7 @@ ggml_backend_sycl_split_buffer_set_tensor(ggml_backend_buffer_t buffer,
         SYCL_CHECK(CHECK_TRY_ERROR(
             (*stream)
                 .memcpy(extra->data_device[i], buf_host, original_size)
-                .wait()));
+                ));
     }
 }
 catch (sycl::exception const &exc) {
@@ -954,7 +953,7 @@ ggml_backend_sycl_split_buffer_get_tensor(ggml_backend_buffer_t buffer,
         SYCL_CHECK(CHECK_TRY_ERROR(
             (*stream)
                 .memcpy(buf_host, extra->data_device[i], original_size)
-                .wait()));
+                ));
     }
 }
 catch (sycl::exception const &exc) {
@@ -2487,7 +2486,7 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten
                           SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(
                                 src1_ddq_i, src1_ddq_i_source,
                                 src1_ncols * src1_padded_col_size * q8_1_ts /
-                                    q8_1_bs).wait()));
+                                    q8_1_bs)));
                         } else {
 
                             float * src1_ddf_i_source = (float *) src1_extra->data_device[ctx.device];
@@ -2554,7 +2553,7 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten
                         dhf_dst_i += src1_col_0*ne0;
                         SYCL_CHECK(CHECK_TRY_ERROR(
                             stream->memcpy(dhf_dst_i, dst_dd_i,
-                                           src1_ncols * ne0 * sizeof(float)).wait()));
+                                           src1_ncols * ne0 * sizeof(float))));
                     }
                 }
 
@@ -3682,7 +3681,7 @@ static void ggml_backend_sycl_get_tensor_async(ggml_backend_t backend,
     GGML_ASSERT(buf->buft == ggml_backend_sycl_buffer_type(sycl_ctx->device) && "unsupported buffer type");
     const queue_ptr stream = sycl_ctx->stream(sycl_ctx->device, 0);
     SYCL_CHECK(CHECK_TRY_ERROR((stream)->memcpy(
-        data, (const char *)tensor->data + offset, size).wait()));
+        data, (const char *)tensor->data + offset, size)));
 }
 catch (sycl::exception const &exc) {
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
@@ -3702,7 +3701,7 @@ static bool ggml_backend_sycl_cpy_tensor_async(ggml_backend_t backend,
         */
         const queue_ptr stream = sycl_ctx->stream(sycl_ctx->device, 0);
         SYCL_CHECK(CHECK_TRY_ERROR((stream)->memcpy(
-            dst->data, src->data, ggml_nbytes(dst)).wait()));
+            dst->data, src->data, ggml_nbytes(dst))));
         return true;
     }