6

我是 CUDA 的新手。我正在尝试并行化以下代码。现在它位于内核上，但根本不使用线程，因此速度很慢。我试图使用这个答案，但到目前为止无济于事。

内核应该生成前 n 个素数，将它们放入 device_primes 数组，稍后从主机访问该数组。代码是正确的，并且在串行版本中工作正常，但我需要加快速度，也许使用共享内存。

//CUDA kernel code
__global__ void generatePrimes(int* device_primes, int n) 
{
//int i = blockIdx.x * blockDim.x + threadIdx.x;
//int j = blockIdx.y * blockDim.y + threadIdx.y;

int counter = 0;
int c = 0;

for (int num = 2; counter < n; num++)
{       
    for (c = 2; c <= num - 1; c++)
    { 
        if (num % c == 0) //not prime
        {
            break;
        }
    }
    if (c == num) //prime
    {
        device_primes[counter] = num;
        counter++;
    }
}
}

我当前的、初步的、绝对错误的并行化尝试如下所示：

//CUDA kernel code
__global__ void generatePrimes(int* device_primes, int n) 
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int num = i + 2; 
    int c = j + 2;
    int counter = 0;

    if ((counter >= n) || (c > num - 1))
    {
        return;
    }
    if (num % c == 0) //not prime
    {
    
    }
    if (c == num) //prime
    {
       device_primes[counter] = num;
       counter++;
    }
    num++;
    c++;
}

但是这段代码用没有意义的数据填充了数组。此外，许多值为零。提前感谢您的帮助，不胜感激。

4

1 回答 1

4

您的代码中存在一些问题，例如：

int num = i + 2;

此表达式分配给thread 0交互 2、thread 1迭代 3，依此类推。问题是线程将计算的下一次迭代是基于表达式的num++;。因此，thread 0将计算下一次迭代 3，该迭代已由 计算thread 1。因此，导致冗余计算。此外，我认为对于这个问题，只使用一维而不是二维会更容易(x,y)。因此，考虑到这一点，您必须更改num++为：

num += blockDim.x * gridDim.x;

另一个问题是您没有考虑到变量counter必须在线程之间共享。否则，每个线程将尝试找到“n”个素数，并且所有这些素数都将填充整个数组。因此，您必须更改int counter = 0;为共享或全局变量。让我们使用一个全局变量，以便它可以在所有块的所有线程中可见。我们可以使用数组的零位置device_primes来保存变量counter。

您还必须初始化此值。让我们将这个作业只分配给一个线程，即 `id = 0 的线程，所以：

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

但是，这个变量是全局的，它会被所有线程写入。因此，我们必须保证所有线程，在写入该全局变量之前，将看到该变量counter为 1（device_primes带素数的第一个位置，零用于counter），因此您还必须在末尾添加一个屏障，所以：

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

所以一个可能的解决方案（尽管效率低下）：

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

    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       
    }
}

以下行将atomicAdd(&device_primes[0], 1);基本执行device_primes[0]++;。我们使用原子操作，因为变量counter是全局的，我们需要保证互斥。请注意，您可能必须使用flag -arch sm_20.

优化：在代码方面，最好使用较少/不同步的方法。此外，还可以通过考虑素数的一些属性来减少计算次数，如http://en.wikipedia.org/wiki/Sieve_of_Eratosthenes中的示例所示。

于 2012-11-04T19:35:49.723 回答