cudaThreadSynchronize & 性能答案

【问题标题】：cudaThreadSynchronize & performancecudaThreadSynchronize & 性能
【发布时间】：2013-03-26 14:42:19
【问题描述】：

几天前，我在比较我的一些代码的性能，其中我执行了相同算法的非常简单的替换和推力实现。我发现一个数量级的不匹配（！）有利于 Thrust，所以我开始让我的调试器“冲浪”到他们的代码中，以发现魔法发生的地方。

令人惊讶的是，我发现我非常直接的实现实际上与他们的非常相似，一旦我摆脱了所有仿函数的东西并进入了本质。我看到 Thrust 有一个聪明的方法来决定 block _size 和 grid_size （顺便说一句：确切地说，它是如何工作的？！），所以我只是接受他们的设置并再次执行我的代码，因为它们非常相似。我获得了一些微秒，但情况几乎相同。然后，最后，我不知道为什么，只是“尝试”我在内核和 BINGO 之后删除了 cudaThreadSynchronize()！我将差距归零（并且更好）并获得了整个数量级的执行时间。访问我的数组的值，我看到它们完全符合我的预期，因此执行正确。

现在的问题是：我什么时候可以摆脱 cudaThreadSynchronize (et similia)？为什么会造成如此巨大的开销？我看到 Thrust 本身最后没有同步（如果未定义宏 __THRUST_SYNCHRONOUS 且未定义，则为 NOP，synchronize_if_enabled(const char* message)）。详细信息和代码如下。

// my replace code
template <typename T>
__global__ void replaceSimple(T* dev, const int n, const T oldval, const T newval)
{
    const int gridSize = blockDim.x * gridDim.x;
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    while(index < n)
    {
        if(dev[index] == oldval)
            dev[index] = newval;
        index += gridSize;
    }
}

// replace invocation - not in main because of cpp - cu separation
template <typename T>
void callReplaceSimple(T* dev, const int n, const T oldval, const T newval)
{       
    replaceSimple<<<30,768,0>>>(dev,n,oldval,newval);
    cudaThreadSynchronize();
}

// thrust replace invocation
template <typename T>
void callReplace(thrust::device_vector<T>& dev, const T oldval, const T newval)
{
    thrust::replace(dev.begin(), dev.end(), oldval, newval);
}

参数详细信息：数组：n=10,000,000 个元素设置为 2，oldval=2，newval=3

执行推力调用替换（推力）的时间：0.057 毫秒
同步执行 callReplaceSimple 的时间：0.662 毫秒
在不同步的情况下执行 callReplaceSimple 的时间：0.011 毫秒

我使用包含 Thrust 的 CUDA 5.0，我的显卡是 GeForce GTX 570，我有一个四核 Q9550 2.83 GHz 和 2 GB RAM。

【问题讨论】：

标签： cuda synchronization thrust

【解决方案1】：

内核启动是异步的。如果删除cudaThreadSynchronize() 调用，则仅测量内核启动时间，而不是直到内核完成的时间。

【讨论】：

此外，推力操作也不一定是阻塞的，所以我建议在内核或推力调用之后使用cudaDeviceSynchronize() 调用来定时这两个操作，以获得良好的比较。（cudaThreadSynchronize() 是 deprecated）。同步调用不会显着扭曲任一操作的时序，但对于基于主机的时序方法，将确保准确标记完成时间。
非常感谢您的回答。因此，我应该始终将 cudaDeviceSynchronize() 用于计时目的，但在“正常行为”中是否需要它？我的意思是：如果我试图按顺序执行 n 个内核并且从不强制同步，它们是否可以毫无问题地执行并且最终将控制权交给 CPU 而无需等待完成？此外：当我在同一设备上调用后续内核（&流！）时，某种同步是否隐含？提前谢谢你。
如果没有cudaDeviceSynchronize()，等待将发生在cudaMemcpy()，您将在该处将结果复制回主机。因此，除非您使用映射内存 cudaDeviceSynchronize() 等高级功能，否则对于正确的结果将不是必需的。
内核之间没有隐式同步，因为正确的结果不需要同步。保证内核在同一流中的前一个内核完成之前不会启动。然而，主机上的执行不受影响，并且不等待任何内核。