pycuda编程之线程同步

持续创作，加速成长！这是我参与「掘金日新计划 · 10 月更文挑战」的第24天，点击查看活动详情

GPU编程也有线程同步和线程通信的概念，在某些时候我们需要确保线程块中的线程都执行完成之后再进行下一轮，这样的描述我们称之为线程同步。一般来说线程同步和线程通信密不可分，在不同的线程之间会相互传递信息，这样一想，那要实现线程同步岂不是非常麻烦？核函数都是C写的，再加上线程的操作，那直接从入门到放弃！其实，问题不大，线程同步没有想象的那么麻烦，不就是等待一个线程块的线程都执行完成之后再一起去执行别的操作嘛，那我们让先执行完的线程等着就是了，CUDA的设备函数__syncthreads，这个函数用于同步内核函数中的单个线程块。

就拿倒数第三篇博客来说，如果们想让grid中的每个block指定迭代次数，那我们用个for直接就解决了，还这么麻烦干啥，但是如果真的不使用CUDA的设备函数，直接上for循环会导致什么问题呢？据说，会导致竞态条件，即多个线程同时访问同一个地址，并执行读写操作，那岂不是乱套了，那不成脏读、幻读了嘛。而__syncthreads(两个_，别写一个，不然报错哈)就是一个线程块级别的同步屏障，也就是说某线程块内的每个线程执行到__syncthreads调用时都会暂停，并等待同一线程块内的其它线程都执行到函数的同一调用，这些线程才会继续执行后面的代码。

下面我们就把之前的代码改写，让其进行1000000一次迭代。相比于之前的代码，主要提供下面改动的部分。

__global__ void conway_ker(int *lattice, int  iters)
{
    int x = _X, y = _Y;
    for(int i=0;i<iters;i++)
    {
       int n = nbrs(x, y, lattice); 
       int cell_value;
       if( lattice[_INDEX(x,y)] == 1 )
        switch(n)
        {
            case 2:
            case 3:cell_value=1;break;
            default: cell_value=0;
        }
    else if( lattice[ _INDEX(x,y) ] == 0 )
        switch(n)
        {
            case 3:cell_value=1;break;
            default:cell_value=0;
        }
        __syncthreads();
    }
}

然后就是主体部分的调用

N = 32
lattice = np.int32(np.random.choice([1, 0], N * N, p=[0.25, 0.75]).reshape(N, N))
lattice_gpu = gpuarray.to_gpu(lattice)
conway_ker(lattice_gpu, np.int32(100000000), grid=(1, 1, 1), block=(32, 32, 1))
fig = plt.figure(1)
plt.imshow(lattice_gpu.get())
plt.show()

最后就是结果了

分别是100x100和32x32的迭代结果。