将 __sync CUDA 与全局内存一起使用答案

【问题标题】：Using __sync CUDA with global memory将 __sync CUDA 与全局内存一起使用
【发布时间】：2013-09-22 07:44:07
【问题描述】：

我正在尝试实现一个异步 PSO。我这样做的方法如下：

__global__ void particle(double *pos, double *pbest, double *vpbest, double *vel, double *gbest){

    int thread = threadIdx.x + blockDim.x * blockIdx.x;
    int particle, i = 0;
    double tpbest;

    double l, r;
    int index, best, j;

    if(thread < DIMPAR){
      particle = thread / NDIM;
        do{
            best    = ring(vpbest, &particle);

            index = (best * NDIM) + (thread % NDIM);

            l = (double) 2.05 * (double) uniform(thread) * ( pbest[thread] -   pos[thread] );
            r = (double) 2.05 * (double) uniform(thread) * ( pbest[index]  -   pos[thread] );

            vel[thread] = vel[thread] + l + r;

                pos[thread] =  pos[thread] + vel[thread];

                __syncthreads(); // I am trying wait all threads write in global memory

            if( (thread % NDIM) == 0 ){ //only one thread replace the vector
                tpbest = rastrigin(pos, particle * NDIM, NDIM);
                if(tpbest < vpbest[particle]){
                    vpbest[particle] = tpbest;
                    for(j = 0 ; j < NDIM; j++){
                        pbest[(particle * NDIM) + j] = pos[(particle * NDIM) + j];
                    }

                }
            }

            i++;
        }while(i < 10000);
    }
}

电话：

particle<<<1,512>>>(d_pos, d_pbest, d_vpbest, d_velo, d_gbest);

有时同步会出现问题... pos[thread] 中的某些值会发散。在 B.6 节 CUDA_C_PROGRAMMING 指南：

等待直到线程块中的所有线程都到达这个点以及这些线程之前进行的所有全局和共享内存访问到 __syncthreads() 对块中的所有线程都是可见的。

pos向量是这样的：

p0 = [0,1,2] //粒子1

p1 = [3,4,5] //粒子 2

p2 = [6,7,8] //粒子 3

pos = [1,2,3,4,5,6,7,8] //pos向量，DIMPAR = 9; NPAR = 3； NDIM = 3

当我使用 NDIM >= 30 时，就会发生分歧

如何使用全局内存确保同步？

【问题讨论】：

__syncthreads() 保证以前对全局或共享内存的更新对块中的所有线程都是可见的。我不清楚你为什么将数值差异归因于__syncthreads() 的问题。你怎么知道数值差异不是由你的算术引起的，就像在你的 previous question 中那样？
@RobertCrovella，谢谢。当循环结束时，我在 C 代码中名为 rastrigin 的函数（已验证函数）中使用 pos 向量并与 vpbest[particle] 值进行比较。在这一点上，价值观是不同的。因此，如果我始终使用相同的值，我将收到相同的值。就好像 pos 矢量和 vpbest 之间存在异步...非常感谢您的帮助
在我看来，这个错误可能在任何地方。对于这样的问题，SO 期望：“有关您编写的代码问题的问题必须在问题本身中描述具体问题 - 并包括有效的代码来重现它。有关指导，请参阅 SSCCE.org。”您没有提供一个 SSCCE.org 代码，以及你怀疑 __syncthreads() 函数的理由对我来说充其量是模糊的。投票结束。请注意，我并不是建议您在此处转储一大堆代码。如果问题如您所描述，应该可以创建一个简单的复制器。

标签： cuda

【解决方案1】：

您的 _syncthread() 位于 if 语句中。请注意，如果块大小大于 DIMPAR，您的程序会停止。要使 _syncthread() 正常工作，块内的所有线程都必须能够访问它。

修改代码的一种方法如下（我不知道您的代码的用途，所以可能有更好的方法）：

particle = thread / NDIM;
    do{
        if(thread < DIMPAR){
            best    = ring(vpbest, &particle);

            index = (best * NDIM) + (thread % NDIM);

            l = (double) 2.05 * (double) uniform(thread) * ( pbest[thread] -   pos[thread] );
            r = (double) 2.05 * (double) uniform(thread) * ( pbest[index]  -   pos[thread] );

            vel[thread] = vel[thread] + l + r;

            pos[thread] =  pos[thread] + vel[thread];
        }

        __syncthreads();

现在所有线程都可以到达同步点。我在其余代码中注意到的另一个问题是 NDIM 的一个线程在 for 循环中工作。相反，您可以将 tpbest 作为块的线程之间的共享变量（或数组）。然后，再次同步线程后，您可以在已经空闲的线程之间分配 NDIM 工作，并让它们一起将效果写入全局内存，而不是使用 for 循环。通过这种方式，访问会更快且合并。

【讨论】：