2

我有一个代码来计算我使用 OpenMP 并行化的素数：

    #pragma omp parallel for private(i,j) reduction(+:pcount) schedule(dynamic)
    for (i = sqrt_limit+1; i < limit; i++)
    {
            check = 1;
            for (j = 2; j <= sqrt_limit; j++)
            {

                    if ( !(j&1) && (i&(j-1)) == 0 )
                    {
                            check = 0;
                            break;
                    }

                    if ( j&1 && i%j == 0 )
                    {
                            check = 0;
                            break;
                    }

            }
            if (check)
                pcount++;

    }

我正在尝试将它移植到 GPU，并且我想像上面的 OpenMP 示例一样减少计数。以下是我的代码，除了给出不正确的结果外，它的速度也较慢：

__global__ void sieve ( int *flags, int *o_flags, long int sqrootN, long int N) 
{
    long int gid = blockIdx.x*blockDim.x+threadIdx.x, tid = threadIdx.x, j;
    __shared__ int s_flags[NTHREADS];

    if (gid > sqrootN && gid < N)
            s_flags[tid] = flags[gid];
    else
            return;
    __syncthreads();

    s_flags[tid] = 1;

    for (j = 2; j <= sqrootN; j++)
    {
            if ( gid%j == 0 )
            {
                    s_flags[tid] = 0;
                    break;
            }
    }
    //reduce        
    for(unsigned int s=1; s < blockDim.x; s*=2)
    {
            if( tid % (2*s) == 0 )
            {
                    s_flags[tid] += s_flags[tid + s];
            }
            __syncthreads();
    }
    //write results of this block to the global memory
    if (tid == 0)
            o_flags[blockIdx.x] = s_flags[0];

}

首先，我如何让这个内核快，我认为瓶颈是for循环，我不知道如何替换它。接下来，我的计数不正确。我确实更改了 '%' 运算符并注意到了一些好处。

在flags数组中，我标记了从 2 到 sqroot(N) 的素数，在这个内核中我正在计算从 sqroot(N) 到 N 的素数，但我需要检查 {sqroot(N),N} 中的每个数字是否可被 {2,sqroot(N)} 中的素数整除。该o_flags数组存储每个块的部分和。

编辑：按照建议，我修改了我的代码（我现在更好地理解了关于同步线程的评论）；我意识到我不需要标志数组，只需要全局索引就可以了。在这一点上，我担心的是代码的缓慢（不仅仅是正确性），这可能归因于 for 循环。此外，在特定数据大小（100000）之后，内核对后续数据大小产生了不正确的结果。即使对于小于 100000 的数据大小，GPU 缩减结果也不正确（NVidia 论坛的一位成员指出，这可能是因为我的数据大小不是 2 的幂）。所以还有三个（可能是相关的）问题——

  1. 我怎样才能让这个内核更快？在我必须遍历每个 tid 的情况下使用共享内存是个好主意吗？

  2. 为什么它只对某些数据大小产生正确的结果？

  3. 我怎样才能修改减少？

    __global__ void sieve ( int *o_flags, long int sqrootN, long int N )
{
unsigned int gid = blockIdx.x*blockDim.x+threadIdx.x, tid = threadIdx.x;
volatile __shared__ int s_flags[NTHREADS];

s_flags[tid] = 1;
for (unsigned int j=2; j<=sqrootN; j++)
{
       if ( gid % j == 0 )
            s_flags[tid] = 0;
}

__syncthreads();
//reduce
reduce(s_flags, tid, o_flags);
}
4

1 回答 1

3

虽然我自称对筛选素数一无所知，但您的 GPU 版本中存在许多正确性问题，无论您实施的算法是否正确，都会阻止它正常工作：

  1. __syncthreads()调用必须是无条件的。编写代码分支发散可能导致同一扭曲中的某些线程无法执行__syncthreads()调用的代码是不正确的。底层 PTX 是bar.sync，PTX 指南是这样说的：

    屏障是在每个 warp 的基础上执行的，就好像一个 warp 中的所有线程都处于活动状态一样。因此，如果 warp 中的任何线程执行 bar 指令，就好像 warp 中的所有线程都执行了 bar 指令。warp 中的所有线程都将停止，直到屏障完成，并且屏障的到达计数按 warp 大小（而不是 warp 中的活动线程数）递增。在有条件执行的代码中，只有当已知所有线程都以相同的方式评估条件时才应该使用 bar 指令（warp 不会发散）。由于屏障是在每个 warp 的基础上执行的，因此可选的线程数必须是 warp 大小的倍数。

  2. s_flags在有条件地从全局内存中加载一些值后，您的代码无条件地设置为 1。这肯定不是代码的意图吗？
  3. 代码在筛选代码和归约之间缺乏同步屏障，这可能导致共享内存竞争和归约的错误结果。
  4. 如果您计划在 Fermi 类卡上运行此代码，则应声明共享内存数组volatile以防止编译器优化可能破坏共享内存减少。

If you fix those things, the code might work. Performance is a completely different issue. Certainly on older hardware, the integer modulo operation was very, very slow and not recommended. I can recall reading some material suggesting that Sieve of Atkin was a useful approach to fast prime generation on GPUs.

于 2011-09-26T10:11:05.507 回答