c++ - CUDA:为什么特定的备忘录复制操作总是比其他类似操作花费 10 倍以上

Question

我相信下面的代码执行了一个典型的

• 复制到设备
• 调用内核
• 复制回主机

工作流程。

1. 我发现非常奇怪的是，当我使用 NSight Profiler 的“跟踪应用程序”选项时，在报告中，“堆栈跟踪”打开时，我发现最昂贵的操作是粗体行，就在那一行，而其他 memoCopy 操作的成本几乎仅为此 memoCopy 操作的 10% 或更少。

   这是因为它是调用内核之后的第一行，因此探查器以某种方式将一些同步的成本包括在这个特定的 memoCopy 操作的成本中吗？

2. 对于像我正在处理的问题，它需要非常频繁的同步并将结果“返回”给主机，任何人都可以提供一些关于最佳实践的一般性建议吗？我特别考虑了两个选项，我不太确定这两个选项最终是否会有所帮助

   • 使用“零拷贝”内存(示例 11.2 中的 CUDA)
   • 创建我如何使用原子操作进行同步

{

int numP = p_psPtr->P.size();
int numL = p_psPtr->L.size();

// Out partition is in Unit of the Number of Particles
int block_dim = BLOCK_DIM_X;
int grid_dim = numP/block_dim + (numP%block_dim == 0 ? 0:1);

vector<Particle> pVec(p_psPtr->P.begin(), p_psPtr->P.end());
Particle *d_part_arr = 0;
Particle *part_arr = pVec.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_part_arr, numP * sizeof(Particle)));
HANDLE_ERROR(cudaMemcpy(d_part_arr, part_arr, numP * sizeof(Particle), cudaMemcpyHostToDevice));

vector<SpringLink> lVec(p_psPtr->L.begin(), p_psPtr->L.end());
SpringLink *d_link_arr = 0;
SpringLink *link_arr = lVec.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_link_arr, numL * sizeof(SpringLink)));
HANDLE_ERROR(cudaMemcpy(d_link_arr, link_arr, numL * sizeof(SpringLink), cudaMemcpyHostToDevice));

Point3D *d_oriPos_arr = 0;
Point3D *oriPos_arr = p_originalPos.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_oriPos_arr, numP * sizeof(Point3D)));
HANDLE_ERROR(cudaMemcpy(d_oriPos_arr, oriPos_arr, numP * sizeof(Point3D), cudaMemcpyHostToDevice));

Vector3D *d_oriVel_arr = 0;
Vector3D *oriVel_arr = p_originalVel.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_oriVel_arr, numP * sizeof(Vector3D)));
HANDLE_ERROR(cudaMemcpy(d_oriVel_arr, oriVel_arr, numP * sizeof(Vector3D), cudaMemcpyHostToDevice));

Point3D *d_updPos_arr = 0;
Point3D *updPos_arr = p_updatedPos.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_updPos_arr, numP * sizeof(Point3D)));
HANDLE_ERROR(cudaMemcpy(d_updPos_arr, updPos_arr, numP * sizeof(Point3D), cudaMemcpyHostToDevice));

Vector3D *d_updVel_arr = 0;
Vector3D *updVel_arr = p_updatedVel.data(); 
HANDLE_ERROR(cudaMalloc((void**)&d_updVel_arr, numP * sizeof(Vector3D)));
HANDLE_ERROR(cudaMemcpy(d_updVel_arr, updVel_arr, numP * sizeof(Vector3D), cudaMemcpyHostToDevice));

int *d_converged_arr = 0;
int *converged_arr = &p_converged[0]; 
HANDLE_ERROR(cudaMalloc((void**)&d_converged_arr, numP * sizeof(int)));
HANDLE_ERROR(cudaMemcpy(d_converged_arr, converged_arr, numP * sizeof(int), cudaMemcpyHostToDevice));

// Run the function on the device
handleParticleKernel<<<grid_dim, block_dim>>>(d_part_arr, d_link_arr, numP,
    d_oriPos_arr, d_oriVel_arr, d_updPos_arr, d_updVel_arr, 
    d_converged_arr, p_innerLoopIdx, p_dt);

**HANDLE_ERROR(cudaMemcpy(oriPos_arr, d_oriPos_arr, numP * sizeof(Point3D), cudaMemcpyDeviceToHost));**
HANDLE_ERROR(cudaMemcpy(oriVel_arr, d_oriVel_arr, numP * sizeof(Vector3D), cudaMemcpyDeviceToHost));
HANDLE_ERROR(cudaMemcpy(updPos_arr, d_updPos_arr, numP * sizeof(Point3D), cudaMemcpyDeviceToHost));
HANDLE_ERROR(cudaMemcpy(updVel_arr, d_updVel_arr, numP * sizeof(Vector3D), cudaMemcpyDeviceToHost));
HANDLE_ERROR(cudaMemcpy(converged_arr, d_converged_arr, numP * sizeof(int), cudaMemcpyDeviceToHost));

Answer 1

那个特定的 cudaMemcpy 调用需要更长的时间，因为它要等到您的内核完成。如果您在内核之后添加一个 cudaDeviceSynchronize，您对该 cudaMemcpy 调用的感知执行时间应该与所有其他调用一致。 (当然，您看到的额外时间将花费在您的 cudaDeviceSynchronize 调用中)。

但是，您花费在 cudaDeviceSynchronize 上的时间是您无法真正避免的基本成本；如果您需要使用内核的输出，则必须等到内核执行完毕。由于内核启动是异步的，您可以在内核运行时执行不相关的语句；但是，在您的情况下，下一个调用是将内核的输出之一复制到主机内存，因此您必须等待内核完成才能获取数据。

如果您的程序允许，您可以尝试将内核启动和内存传输分解为多个 block 并使用不同的流启动它们，尽管这样做的可行性取决于几个因素(即您的内核可能无法很好地分解为独立的部分)。如果你真的走这条路，最好的情况就是这样(取自 the CUDA Best Practices Docs )

这将允许您将数据传输与内核执行重叠，从而隐藏一些数据传输成本。您可以使用零拷贝实现类似的异步，只是预先警告此类传输不会被缓存，因此根据您的内核访问模式，您最终可能会获得较低的吞吐量。