0

几天前，我正在比较我的一些代码的性能，其中我执行了相同算法的非常简单的替换和推力实现。我发现一个数量级的不匹配（！）有利于 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 ms
  • 同步执行 callReplaceSimple 的时间：0.662 毫秒
  • 不同步执行 callReplaceSimple 的时间：0.011 毫秒

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

4

1 回答 1

6

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

于 2013-03-26T01:38:43.747 回答