fix : cuda order of synchronization when setting a buffer (ggml/679)

Green-Sky · slaren · ggerganov · commit 11b1b63b14bb · 2024-01-05T17:01:59.000+02:00
* fix : cuda order of synchronization when setting a buffer

* also sync before memcpy

---------

Co-authored-by: slaren &lt;slarengh@gmail.com&gt;
diff --git a/extra/sync-ggml.last b/extra/sync-ggml.last
@@ -1 +1 @@
-f96711108d55bdbbd277e6be07204dce6a94fb93
+3eace58911ea8d2cf35defdc59848d99b91a57f5
diff --git a/ggml-cuda.cu b/ggml-cuda.cu
@@ -9689,8 +9689,8 @@ static void ggml_backend_cuda_buffer_set_tensor(ggml_backend_buffer_t buffer, gg
 
     ggml_cuda_set_device(ctx->device);
     CUDA_CHECK(cudaDeviceSynchronize());
-
     CUDA_CHECK(cudaMemcpy((char *)tensor->data + offset, data, size, cudaMemcpyHostToDevice));
+    CUDA_CHECK(cudaDeviceSynchronize());
 }
 
 static void ggml_backend_cuda_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) {

Original file line number	Diff line number	Diff line change
`@@ -9689,8 +9689,8 @@ static void ggml_backend_cuda_buffer_set_tensor(ggml_backend_buffer_t buffer, gg`
`9689`	`9689`
`9690`	`9690`	`ggml_cuda_set_device(ctx->device);`
`9691`	`9691`	`CUDA_CHECK(cudaDeviceSynchronize());`
`9692`		`-`
`9693`	`9692`	`CUDA_CHECK(cudaMemcpy((char *)tensor->data + offset, data, size, cudaMemcpyHostToDevice));`
	`9693`	`+ CUDA_CHECK(cudaDeviceSynchronize());`
`9694`	`9694`	`}`
`9695`	`9695`
`9696`	`9696`	`static void ggml_backend_cuda_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) {`