本地内存#

全局内存不用说了，前面动态分配的就是。下面是本地内存。

本地内存逻辑上属于某个线程私有，实际上和全局内存一样慢。什么是本地内存？简单讲，就是单个线程的一些变量，在寄存器里放不下了，就会用到本地内存。

（写作本地内存，逻辑上我们当成单个线程私有的，其实就是给物理上扔到显存里，读取很慢）

共享内存#

共享内存是用__shared__修饰的，类似CPU一级缓存，但是不同的是，是可以自己开辟空间的。也就是说，我可以直接指定在共享内存中开一个数组用来存东西，同一个Block中读取速度比较快。

常量内存#

我们知道，全局内存虽然大，但是很慢啊。但是如果我们要给所有线程一块数据，只读不写，有没有办法加速？

常量内存。

常量内存用__constant__修饰，声明在核函数外面，也就是全局作用域。

1
__constant__ float filter[16];  // 声明常量内存

我们知道，正常的C语言，写常量是直接在声明里写。

1
const int a = 16;

不过，这里声明的是在显存里的，就不能直接像C语言一样直接写。

所以要用专门的函数写入：

1
cudaMemcpyToSymbol(filter, h_filter, 16 * sizeof(float));

这写完后，在kernel执行期间，GPU只能读不能写。

这有什么好处呢？

快。不过这里的快是有要求的。GPU有专门的常量缓存。如果同一个Warp里的32个线程，同时读的是同一个地址，那只需要读一次，广播给所有线程，速度接近寄存器。

NOTE

我们可能会想，在用CUDA C的时候，诸如神经网络的层数layer_size之类的，从前我们会这样写：

1
train_kernel_skipgram<<<grid, block>>>(
2
    d_syn0, d_syn1neg, d_batch,
3
    BATCH_SIZE,
4
    layer1_size,      // ← 这些标量参数
5
    alpha
6
);

有了常量内存后，直接放常量内存，会不会更快？

不会。

因为CUDA的核函数参数，本来就自动存在常量内存里。

那什么时候才值得手动用 __constant__？

假设我们有一个什么东西要查表，就可以这样扔进常量内存里。

而且，常量内存要求的比较苛刻，如果是随机采样的访问查表，也不优，反而可能比全局内存更慢（常量内存不擅长处理分散读取）。

所以常量内存的使用条件其实挺苛刻的，同时满足三点才划算：

1. 只读
2. 数据小于64KB（常量内存只有这么大）
3. 访问模式是”所有线程读同一地址”（广播）

全局内存#

前面我们学过可以动态分配显存，那可以静态分配吗？

可以。

__device__ 可以用来声明变量，这样声明的变量住在全局内存里，但是是静态分配的（编译时就确定大小）。

1
__device__ float d_alpha;           // 单个标量
2
__device__ float d_table[1000];     // 固定大小的数组

从CPU读写 __device__ 变量，不能直接赋值，要用专门的函数：

1
__device__ float d_alpha = 0.025f;
2

3
// CPU往里写
4
float val = 0.01f;
5
cudaMemcpyToSymbol(d_alpha, &val, sizeof(float));
6

7
// CPU从里读
8
float result;
9
cudaMemcpyFromSymbol(&result, d_alpha, sizeof(float));

固定内存#

我们学操作系统的时候就知道，操作系统为了管理内存，普通内存都是可以换页的。

当内存紧张时，OS可以把这块数据临时挪到硬盘上（swap），腾出物理内存给别的程序用。

如果需要大量的GPU向CPU传输，可以改成固定内存。也就是用前面的cudaMallocHost来分配内存。

比如：

1
int *corpus = (int *)malloc(corpus_size * sizeof(int));
2

3
// 改成
4
int *corpus;
5
cudaMallocHost(&corpus, corpus_size * sizeof(int));

零拷贝内存#

普通的数据传输流程是这样的：

1
CPU内存 → cudaMemcpy → GPU显存 → 核函数使用

这里的cudaMemcpy需要显式传输，占用时间。

零拷贝内存的思路是：跳过这个传输步骤，让GPU直接读CPU内存。

1
float *h_data;
2
cudaHostAlloc(&h_data, size, cudaHostAllocMapped);  // 分配零拷贝内存
3

4
float *d_data;
5
cudaHostGetDevicePointer(&d_data, h_data, 0);  // 获取GPU视角的指针
6

7
// 核函数直接用d_data，不需要cudaMemcpy
8
kernel<<<grid, block>>>(d_data);

不过，这里有个问题，零拷贝的意思是不需要显式调用cudaMemcpy，但数据还是要通过PCI-E总线传输的，这个物理过程跑不掉。

众所周知，PCI-E总线很慢，比GPU显存带宽慢，所以其实性能上是不优的。

什么时候用呢？

适合：

1. 数据量太大，GPU显存根本放不下
2. 数据只读一遍，没必要先传过去再读
3. CPU和GPU需要共享同一块数据，频繁同步

合并访问、非合并访问#

最理想的情况，是这样的：32个线程，访问连续的32个float（每个4字节，共128字节），硬件一次事务就全取完了。

1
// 假设 blockDim.x = 32，一个warp
2
int i = threadIdx.x;
3
float val = data[i];  // 线程0取data[0]，线程1取data[1]...线程31取data[31]

线程: 0 1 2 3 4 5 … 31 地址: [0][1][2][3][4][5]...[31] ，连续！一次搞定

这叫合并访问，效率最高。

非合并访问就很坏了。

比如，跨步访问：

1
int i = threadIdx.x;
2
float val = data[i * 2];  // 线程0取[0]，线程1取[2]，线程2取[4]...

这取的地址是[0] [2] [4] [6] ...，有间隔啊！

每隔一个元素，硬件没法合并（不能说完全无法合并，但是势必要浪费一半），需要多次事务，带宽浪费一半。

更坏的是随机访问。

1
float val = data[random_index[i]];  // 每个线程访问随机地址

32个线程触发32次独立事务，比合并访问慢32倍。

对齐访问#

先从CPU讲起。假设有一个结构体：

1
struct Example {
2
    char a;   // 1字节
3
    int b;    // 4字节
4
};

这个结构体占多少个字节？

看上去是5个字节是吗？但并不是，sizeof(Example)其实是8。

因为CPU访问内存时，硬件要求（或强烈偏好）一个N字节的数据类型，其起始地址必须是N的倍数。

int是4字节，所以它的地址必须是4的倍数。编译器会在a后面自动插入3字节的padding，让b对齐到偏移量4的位置。

（至于结构体的对齐要求，是其成员中最大的那个对齐要求。）

一般来说，在我们写普通C代码的时候，编译器会自动插入padding，自动把变量对齐。

如果要看见不对齐的效果，可以尝试用指针：

1
#include <stdio.h>
2
#include <string.h>
3

4
int main() {
5
    char buf[8];
6
    int *p = (int *)(buf + 1);  // 故意让int指针指向一个非4对齐的地址
7
    *p = 42;
8
    printf("%d\n", *p);
9
    return 0;
10
}

在x86上这段代码能跑，输出42，但访问速度会比对齐的慢——因为一次内存读取跨越了两个cache line边界，硬件需要读两次再拼起来。

在某些ARM芯片上（比如早期的嵌入式ARM），这段代码会直接触发SIGBUS崩溃。现代ARM（比如手机上的）一般也改成不崩但变慢了。

然后再看GPU上的对齐。

我们前面知道了，GPU一次内存事务是32字节、64字节或128字节，而且这个事务的起始地址必须是事务大小的整数倍。

比如一个warp的32个线程，每人读一个float（4字节），总共128字节。如果起始地址是128的倍数，硬件一次事务搞定。如果起始地址不对齐，比如偏移了4字节，那这128字节横跨了两个128字节的事务边界，硬件就要发两次事务，带宽浪费接近一半。

不同洗牌指令#