cudaThreadSynchronize&performance

Question

几天前,我正在比较我的一些代码的性能,我执行一个非常简单的替换和Thrust实现相同的算法.我发现一个数量级(!)的不匹配有利于Thrust,所以我开始让我的调试器"冲浪"到他们的代码中以发现魔法发生的地方.

令人惊讶的是,我发现我的非常直接的实现实际上非常类似于他们的,一旦我摆脱了所有的functor东西并得到了细节.我看到Thrust有一个聪明的方法来决定块_size和grid_size(顺便说一句:确切地说,它是如何工作的？!),所以我只是采取了他们的设置并再次执行我的代码,因为它们非常相似.我增加了几微秒,但情况几乎相同.然后,最后,我不知道为什么,只是"尝试"我删除了我的内核和BINGO之后的cudaThreadSynchronize()!我将差距归零(并且更好),并获得了整个执行时间的数量级.访问我的数组值我发现它们完全符合我的预期,所以正确执行.

现在的问题是:我什么时候可以摆脱cudaThreadSynchronize(et similia)？为什么会造成如此巨大的开销？我看到Thrust本身在最后没有同步(synchronize_if_enabled(const char*message)是一个NOP,如果没有定义宏__THRUST_SYNCHRONOUS而且它不是).细节和代码如下.

// 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 ms
• 执行callReplaceSimple而不同步的时间:0.011毫秒

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

Answer 1

内核启动是异步的.如果删除cudaThreadSynchronize()调用,则只测量内核启动时间,而不是内核完成时间.