CUDA访问内存的效率将决定程序的效率。在此将介绍CUDA的几种存储和合理使用方式。

全局内存

对全局内存的访问将触发数据传输。传输的规则如下》

• 每次传输的大小为32字节。
• 首地址为128的整数倍。

可以用合并度这个量度来描述“被浪费的数据传输次数”。假设线程需要A Byte数据，结果因为数据对齐等原因传输了B Byte， 合并度 = A / B。

如果合并度不为1， 则称为非合并访问。

在非合并访问不可避免的情况下，CUDA对非合并的读操作有缓存机制，所以应尽量保证写操作的合并度。

共享内存

共享内存与线程块对应，每个线程块都有自己的副本；不同线程块的值可以不同；相当于在线程块级别上的“全局内存”缓存。

共享内存大小有一定限制，可以声明固定大小或动态大小的共享内存：

void __global__ something() {

__shared__ float s_y[128]; //在核函数内声明固定大小的共享内存

}

<<>> // 在核函数的执行配置指定第三个参数，给出shared大小

void __global__ something() {

extern __shared__ float s_y[]; //在核函数内用extern修饰共享内存大小

}

不同线程对同一共享内存的访问，需要__syncthreads()来保证先后访问顺序。

例如：

for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) {

if (tid < offset) {

s_y[tid] += s_y[tid + offset];

}

__syncthreads();

}

这是一个数组reduce的例子。数组里的每个元素需要累加指定的offset值。syncthread保证每个线程都会先完成 offset=blockDim.x>>1累加。等所有线程都完成这一步后，才会累加blockDim.x >> 2；

利用共享内存可以改善全局内存的访问。可以以合并访问的形式，将数据拷贝到共享内存；然后在共享内存上进行操作，最后再以合并访问的形式写入全局内存；通过添加一层缓存，可以绕过输入和输出总会有一个非合并的问题。

另外，共享内存本身划分为多个Bank。多个线程同时访问一个Bank也将造成冲突。这也应该尽量避免。

原子函数

考虑如下代码：

if (tid == 0) {

d_y[0] += s_y[0];

}

在每个block的Thread0里都会执行这一行，它被多个线程访问，但是访问顺序不确定。同CPU代码一样，这里需要有取出元素-相加-写入三步，会出现读写竞争。

和CPU端一样，CUDA提供了atomic操作，可以直接调用。

一个函数声明的例子如下：

T atomicAdd(T *address, T val)

address代表待累加的地址，val代表需要累加的函数值。