9

我正在使用一个天真的素数生成函数。这段代码大约需要 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.707 秒内生成 10k。我只是想知道为什么通过这种简单的修改来加快速度，破坏那么糟糕吗？

4

1 回答 1

2

正如 Tony 所建议的那样，不同的代码执行可能会导致 gpu 代码的速度大幅下降，从而迫使某些代码以串行而不是并行的方式运行。在上述代码的慢版本中，遇到中断的线程与继续执行的代码不同。

cuda c 编程指南是一个很好的gpu 编程技术资源。 这是关于控制流的内容：

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

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

于 2012-11-25T23:07:15.083 回答