打破开销与控制标志

Question

我正在使用一个天真的素数生成函数.这段代码需要大约5.25秒来生成10k素数(device_primes [0]保存已经找到的数字素数,其余位置是素数找到的).

_global__ void getPrimes(int *device_primes,int n)
{ 
    int c = 0;
    int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
    int num = thread_id+2;

    if (thread_id == 0) device_primes[0] = 1;
    __syncthreads();

    while(device_primes[0] < n)
    {
        for (c = 2; c <= num - 1; c++)
        { 
            if (num % c == 0) //not prime
            {
                break;
            }
        }
        if (c == num) //prime
        {
            int pos = atomicAdd(&device_primes[0],1);
            device_primes[pos] = num;
        }
        num += blockDim.x * gridDim.x; // Next number for this thread       
    }
}

我刚开始优化代码,我做了以下修改,而不是:

for (c = 2; c <= num - 1; c++)
{ 
    if (num % c == 0) //not prime
         break;
}
 if (c == num) {...}

我现在有了 :

   int prime = 1;

   ...
   for (c = 2; c <= num - 1 && prime; c++)
    { 
        if (num % c == 0) prime = 0; // not prime
    }
     if (prime) {...} // if prime

现在我可以在0.707s中产生10k.我只是想知道为什么这种简单的修改加速,是打破那么糟糕？

Answer 1

正如 Tony 所建议的，不同的代码执行可能会导致 GPU 代码大幅减慢，迫使某些代码串行运行而不是并行运行。在上面代码的慢速版本中，遇到中断的线程与继续运行的代码不同。

cuda c 编程指南是GPU 编程技术的良好资源。 这是关于控制流的内容：

任何流控制指令（if、switch、do、for、while）都会导致同一 warp 的线程发散（即遵循不同的执行路径），从而显着影响有效指令吞吐量。如果发生这种情况，则必须对不同的执行路径进行序列化，从而增加为此扭曲执行的指令总数。当所有不同的执行路径完成后，线程会收敛回相同的执行路径。

较新的 nvidia 硬件和 cuda 版本可以比旧版本更好地处理某些分支，但最好还是尽可能避免分支。