cuda - 如何正确地将全局内存中的读取合并到具有 short 或 char 类型元素的共享内存中(假设每个元素一个线程)？

Question

我对 CUDA 中合并的全局内存负载有疑问。目前我需要能够在具有计算能力 CUDA 1.1 或 1.3 的 CUDA 设备上执行。

我正在编写一个 CUDA 内核函数，它将 T 类型的数组从全局内存读取到共享内存，进行一些计算，然后将 T 类型的数组写回全局内存。我正在使用共享内存，因为每个输出元素的计算实际上不仅取决于相应的输入元素，还取决于附近的输入元素。我只想加载每个输入元素一次，因此我想将输入元素缓存在共享内存中。

我的计划是让每个线程将一个元素读入共享内存，然后在开始计算之前调用 __syncthreads()。在这种情况下，每个线程加载、计算和存储一个元素(尽管计算取决于其他线程加载到共享内存中的元素)。

对于这个问题，我想关注从全局内存到共享内存的读取。

假设数组中有N个元素，我配置CUDA一共执行N个线程。对于 sizeof(T) == 4 的情况，根据我对 CUDA 的理解，这应该很好地合并，因为线程 K 将读取单词 K(其中 K 是线程索引)。

但是，在 sizeof(T) < 4 的情况下，例如，如果 T=unsigned char 或 T=short，那么我认为可能存在问题。在这种情况下，我的(天真的)计划是:

• 计算 numElementsPerWord = 4/sizeof(T)
• if(K % numElementsPerWord == 0), then read 让线程 K 读取下一个完整的 32 位字
• 将 32 位字存储在共享内存中
• 在填充共享内存(并调用 __syncthreads())之后，每个线程 K 都可以处理计算输出元素 K 的工作

我担心的是它不会合并，因为(例如，在 T=short 的情况下)

• 线程 0 从全局内存中读取字 0
• 线程 1 不读取
• 线程 2 从全局内存中读取字 1
• 线程 3 不读
• 等...

换句话说，线程K读取了word(K/sizeof(T))。这似乎无法正确合并。

我考虑的另一种方法是:

• 以线程数 = (N + 3)/4 启动，这样每个线程将负责加载和处理 (4/sizeof(T)) 个元素(每个线程处理一个 32 位字 - 可能是 1 个， 2 个或 4 个元素，具体取决于 sizeof(T))。然而，我担心这种方法不会尽可能快，因为每个线程必须执行两次(如果 T=short)甚至四倍(如果 T=unsigned char)处理量。

有人可以告诉我我对我的计划的假设是否正确:即:它不会正确合并？

您能评论一下我的替代方法吗？

您能否推荐一种更优化的正确合并方法？

Answer 1

你是对的，你必须进行至少 32 位大小的加载才能合并，而你描述的方案(让所有其他线程执行加载)不会合并。只需将偏移量右移 2 位，让每个线程执行连续的 32 位加载，并使用条件代码来禁止对超出范围的地址进行操作的线程执行。

由于您的目标是 SM 1.x，还请注意 1) 为了进行合并，给定 warp 的线程 0(32 个线程的集合)必须为 4 的 64、128 或 256 字节对齐-，分别是 8 字节和 16 字节操作数，以及 2) 一旦您的数据位于共享内存中，您可能希望将循环展开 2x(简称)或 4x(对于 char)，以便相邻线程引用相邻的 32 位词，以避免共享内存库冲突。