CUDA并行化嵌套的for循环

Question

我是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++;
}

但是这段代码填充数组的数据并没有意义。此外，许多值都是零。提前感谢任何帮助，非常感激。

- Nikita K

Answer 1

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

int num = i + 2;

该表达式将交互2分配给线程0，将迭代3分配给线程1，依此类推。问题在于线程将计算的下一个迭代是基于表达式num++;。因此，线程0将计算下一个已经被线程1计算过的迭代3，从而导致冗余计算。此外，我认为对于这个问题，只使用一个维度而不是两个(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中质数的第一个位置，0是给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 是全局的，我们需要保证互斥。请注意，您可能需要使用 -arch sm_20 标志编译。 优化： 从代码角度来看，使用较少或不需要同步的方法将更好。此外，通过考虑质数的某些属性，可以减少计算次数，如Sieve of Eratosthenes中所示。

- dreamcrash