fix gpups error when upgrade cuda12.3 (#63391)

danleifeng · web-flow · commit de54be3edeb3 · 2024-04-11T11:03:49.000+08:00
diff --git a/paddle/fluid/framework/fleet/heter_ps/feature_value.cu b/paddle/fluid/framework/fleet/heter_ps/feature_value.cu
@@ -355,22 +355,26 @@ void AccessorWrapper<GPUAccessor>::CopyForPushImpl(
   int64_t* gpu_len = reinterpret_cast<int64_t*>(buf_length->ptr());
   int* d_slot_vector = reinterpret_cast<int*>(buf_slot_vector->ptr());
   int* d_mf_dim_vector = reinterpret_cast<int*>(buf_mf_dim_vector->ptr());
-  cudaMemcpy(gpu_values,
-             grad_values.data(),
-             grad_values.size() * sizeof(float*),
-             cudaMemcpyHostToDevice);
-  cudaMemcpy(gpu_len,
-             slot_lengths_lod.data(),
-             slot_lengths.size() * sizeof(int64_t),
-             cudaMemcpyHostToDevice);
-  cudaMemcpy(d_slot_vector,
-             slot_vector.data(),
-             slot_lengths_lod.size() * sizeof(int),
-             cudaMemcpyHostToDevice);
-  cudaMemcpy(d_mf_dim_vector,
-             slot_mf_dim_vector.data(),
-             slot_lengths_lod.size() * sizeof(int),
-             cudaMemcpyHostToDevice);
+  cudaMemcpyAsync(gpu_values,
+                  grad_values.data(),
+                  grad_values.size() * sizeof(float*),
+                  cudaMemcpyHostToDevice,
+                  stream);
+  cudaMemcpyAsync(gpu_len,
+                  slot_lengths_lod.data(),
+                  slot_lengths.size() * sizeof(int64_t),
+                  cudaMemcpyHostToDevice,
+                  stream);
+  cudaMemcpyAsync(d_slot_vector,
+                  slot_vector.data(),
+                  slot_lengths_lod.size() * sizeof(int),
+                  cudaMemcpyHostToDevice,
+                  stream);
+  cudaMemcpyAsync(d_mf_dim_vector,
+                  slot_mf_dim_vector.data(),
+                  slot_lengths_lod.size() * sizeof(int),
+                  cudaMemcpyHostToDevice,
+                  stream);
   PushCopyWithPool<<<(total_length + 1024 - 1) / 1024, 1024, 0, stream>>>(
       total_grad_values_gpu,
       gpu_values,
diff --git a/paddle/fluid/framework/fleet/heter_ps/heter_comm_inl.h b/paddle/fluid/framework/fleet/heter_ps/heter_comm_inl.h
@@ -1631,6 +1631,7 @@ void HeterComm<KeyType, ValType, GradType, GPUAccessor>::pull_merge_sparse(
                 val_type_size);
   }
 
+  AnyDeviceGuard guard2(dev_id);
   auto d_merged_vals = MemoryAlloc(place, uniq_len * val_type_size);
   auto d_merged_vals_ptr = reinterpret_cast<float *>(d_merged_vals->ptr());
   heter_comm_kernel_->dy_mf_fill_dvals(d_shard_vals_ptr,

Original file line number	Diff line number	Diff line change
`@@ -1631,6 +1631,7 @@ void HeterComm<KeyType, ValType, GradType, GPUAccessor>::pull_merge_sparse(`
`1631`	`1631`	`val_type_size);`
`1632`	`1632`	`}`
`1633`	`1633`
	`1634`	`+ AnyDeviceGuard guard2(dev_id);`
`1634`	`1635`	`auto d_merged_vals = MemoryAlloc(place, uniq_len * val_type_size);`
`1635`	`1636`	`auto d_merged_vals_ptr = reinterpret_cast<float *>(d_merged_vals->ptr());`
`1636`	`1637`	`heter_comm_kernel_->dy_mf_fill_dvals(d_shard_vals_ptr,`