Coalesced Memory Access
要实现合并内存访问,需要满足以下条件:
连续内存地址: warp 中的相邻线程必须访问连续的全局内存地址。这意味着线程 0 访问地址 A,线程 1 访问地址 A+B,线程 2 访问地址 A+2B,以此类推,其中 B 是每个线程访问的内存大小。
对齐: 访问的内存地址起始位置必须按照一定大小对齐。对齐大小取决于每个线程访问的数据大小和 GPU 的计算能力。例如,对于 32 位数据,通常需要 128 字节对齐。上面的 A = 32.
数据大小: 访问的数据大小也影响合并的效率。较大的数据访问(例如 128 位或 256 位)通常比较小的数据访问(例如 8 位或 16 位)更容易合并。
合并访问的理想情况是一个 warp 中的 32 个线程应该能够在一个或少数几个内存事务中完成所有数据的读取或写入。如果不是合并访存,GPU 可能需要多个内存事务才能完成所有数据的读取或写入。
上述中的:“对于 32 位数据,通常需要 128 字节对齐”,如何理解:每个线程访问 32位 数据,就是 4 字节数据,一个 warp 32 线程,那么一个 warp 会访问 32 * 4 = 128 字节数据。如果这 128 字节的数据在内存中是 128 字节对齐的(起始地址是 128 的倍数),那么 GPU 就可以在一个内存事务中完成所有数据的读取或写入,从而实现合并访问。
连续的 threads 访问连续的地址,效率是最高的。保证在一个内存事务可以传输多个数据字。
一个比喻:想象从一个装满巧克力的大盒子里取巧克力。如果每个线程都从盒子的同一层取(连续的内存),那么取的过程会非常高效。但如果每个线程都从盒子的不同层取(非连续的内存),那么取糖果的过程就会变得低效,因为需要多次移动到不同的层。
关于上述的 对齐
在 CPU 系统上,对于 4 字节数据,通常需要 4 字节对齐。
所有线程访问数据的起始地址必须是这个对齐大小的倍数。
- 4 字节数据: 如果每个线程访问 4 字节的数据(例如一个 float),那么所有线程访问数据的起始地址必须是 4 的倍数 (0, 4, 8, 12, …)。
 - 8 字节数据: 如果每个线程访问 8 字节的数据(例如一个 double),那么所有线程访问数据的起始地址必须是 8 的倍数 (0, 8, 16, 24, …)。
 - 16 字节数据: 以此类推。
 
如何保证 线程访问数据的起始位置是上述的情况?
struct MyData {
  float value;
  char padding[3]; // 填充 3 个字节,使结构体大小为 4 的倍数
};
结构体填充 (Padding): 如果你的数据存储在结构体中,你可以通过在结构体中添加填充成员来确保对齐。编译器通常会自动进行填充以满足对齐要求,但为了确保,你可以手动添加填充成员。 在上述例子中,padding 数组确保 MyData 结构体的大小是 4 的倍数,即使 value 只有 4 字节。 这样,即使你有多个 MyData 实例,它们在内存中的起始地址也自然会是 4 的倍数。
数组: 如果你的数据存储在数组中,那么数组元素的地址本身就自然满足对齐要求。假设你有一个 float 类型的数组:
float data[100];data[0]的地址将是 4 的倍数,data[1]的地址将是 8 的倍数,以此类推。cudaMallocPitch(对于二维数据): 如果你处理的是二维数据(例如矩阵),cudaMallocPitch函数可以帮助你分配内存并控制行对齐。 它允许你指定每行的字节数,从而确保每行的起始地址满足对齐要求。 这对于处理矩阵等二维数据非常有用,可以确保行访问的合并性。使用
alignas()等强制内存对齐。详见 SIMD
编译器通常会进行优化以确保数据对齐,但为了确保你的代码在不同编译器和平台上都能正确运行,最好手动进行对齐处理。***
给个实例,打印连续10个对象的地址和与前一个对象的地址差:
#include <iostream>
using namespace std;
struct MyData {
  float value;
  char padding[3]; // 填充 3 个字节,使结构体大小为 4 的倍数
};
int main() {
  MyData myDataArray[10];
  // 打印每个对象的起始地址和与前一个对象的地址差
  for (int i = 0; i < 10; ++i) {
    std::uintptr_t address = reinterpret_cast<std::uintptr_t>(&myDataArray[i]);
    std::cout << "MyData[" << i << "] address: " << address << std::endl;
    if (i > 0) {
      std::uintptr_t prevAddress = reinterpret_cast<std::uintptr_t>(&myDataArray[i - 1]);
      std::cout << "Difference from previous: " << address - prevAddress << " bytes" << std::endl;
    }
  }
  return 0;
}
编译器会对齐,返回:
MyData[0] address: 140724588819808  # 所有地址都是 8 的倍数
MyData[1] address: 140724588819816  
Difference from previous: 8 bytes
MyData[2] address: 140724588819824
Difference from previous: 8 bytes
MyData[3] address: 140724588819832
Difference from previous: 8 bytes
MyData[4] address: 140724588819840
Difference from previous: 8 bytes
MyData[5] address: 140724588819848
Difference from previous: 8 bytes
MyData[6] address: 140724588819856
Difference from previous: 8 bytes
MyData[7] address: 140724588819864
Difference from previous: 8 bytes
MyData[8] address: 140724588819872
Difference from previous: 8 bytes
MyData[9] address: 140724588819880
Difference from previous: 8 bytes
编译器会自动填充。