持续创作,加速成长!这是我参与「掘金日新计划 · 10 月更文挑战」的第25天,点击查看活动详情 之前的博客中已经接触到了cuda编程的几种特点,前面代码中使用的都是数组通信的方式,因为在全局内存中,受到线程块与网格结构的影响,不能很方便的像在主机中那样进行内存共享。其实,使用数组这样的全局内存已经可以解决大部分运算,但是共享内存可以提高计算的速度。在GPU编程中,共享内存指的是block级别的共享,在一个block中线程间可以进行通信的内存,需要注意的一点是,共享内存里面的数据如果想要在全局内存中使用,必须先通过内核函数将其复制到全局内存里面。
GPU的网格布局,使得每个线程都有固定的位置,就好比一个笛卡尔坐标系一样,其中threadIdx.x和threadIdx.y类比于笛卡尔坐标系的x和y。下面通过一个例子来说明一下共享内存和全局变量。在内核函数中声明了一个数组int a[10],然后函数中的参数int *b作为参数。比较这两个变量,对于a数组来说,在内核函数中所有的线程都有一个单独的a数组,其它线程无法读取a,除此之外每个线程也有一个指针b,但是与a不同的是指针b不仅具有相同的值,而且可以让所有线程访问。
下面我就举一个例子,这里与之前的线程同步一样,也有关键字函数修饰
__global__ void shared_demo(int *p_lattice, int iters)
{
int x = _X, y=_Y;
__shared__int lattice[32*32];
lattice[_INDEX(x,y)]=p_lattice[_INDEX(x,y)];
__syncthreads();
}
这个部分代码就是前面博客中代码的部分修改,只不过将原来的全局内存改为了共享内存,经过测试,当数据量小的时候时间变化不明显,当数据量较大的时候时间变化比较明显。在pycuda编程中还有并行前缀算法的概念,就是所谓的扫描设计模式,关于这部分的知识点,将在下一篇博客中详细的叙述。