use cudaMemcpy3DPeerAsync

Author: slaren

ggml-cuda.cu

@@ -68,9 +68,13 @@
 #define cudaMallocHost(ptr, size) hipHostMalloc(ptr, size, hipHostMallocDefault)
 #endif
 #define cudaMemcpy hipMemcpy
-#define cudaMemcpy2DAsync hipMemcpy2DAsync
 #define cudaMemcpyAsync hipMemcpyAsync
 #define cudaMemcpyPeerAsync hipMemcpyPeerAsync
+#define cudaMemcpy2DAsync hipMemcpy2DAsync
+#define cudaMemcpy3DPeerAsync hipMemcpy3DPeerAsync
+#define cudaMemcpy3DPeerParms hipMemcpy3DPeerParms
+#define make_cudaPitchedPtr make_hipPitchedPtr
+#define make_cudaExtent make_hipExtent
 #define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice
 #define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost
 #define cudaMemcpyHostToDevice hipMemcpyHostToDevice
@@ -8258,14 +8262,15 @@ static void ggml_cuda_op_mul_mat(
                         float * dhf_dst_i = (float *) ((char *) dst_off_device + i02*nb2 + i03*nb3);
                         GGML_ASSERT(dst->nb[1] == ne0*sizeof(float));
                         dhf_dst_i += src1_col_0*ne0 + row_low[id];
-
                         if (kind == cudaMemcpyDeviceToDevice && id != g_main_device) {
-                            // there is no cudaMemcpy2DPeerAsync so we need to copy each row separately
-                            for (int64_t i = 0; i < src1_ncols; ++i) {
-                                CUDA_CHECK(cudaMemcpyPeerAsync(dhf_dst_i + i*ne0, g_main_device,
-                                                                dst_dd_i + i*row_diff, id,
-                                                                row_diff*sizeof(float), stream));
-                            }
+                            // there is no cudaMemcpy2DPeerAsync, use a 3D copy instead
+                            cudaMemcpy3DPeerParms p = {};
+                            p.dstDevice = g_main_device;
+                            p.dstPtr = make_cudaPitchedPtr(dhf_dst_i, ne0*sizeof(float), ne0, src1_ncols);
+                            p.srcDevice = id;
+                            p.srcPtr = make_cudaPitchedPtr(dst_dd_i, row_diff*sizeof(float), row_diff, src1_ncols);
+                            p.extent = make_cudaExtent(row_diff*sizeof(float), src1_ncols, 1);
+                            CUDA_CHECK(cudaMemcpy3DPeerAsync(&p, stream));
                         } else {
                             CUDA_CHECK(cudaMemcpy2DAsync(dhf_dst_i, ne0*sizeof(float), dst_dd_i, row_diff*sizeof(float),
                                                         row_diff*sizeof(float), src1_ncols, kind, stream));