我有一个关于 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)，然后让线程K读取下一个完整的 32 位字 - 将 32 位字存储在共享内存中
- 填充(并
__syncthreads()调用(共享内存后，每个线程K可以处理计算输出元素的工作K
我担心的是它不会合并，因为(例如，在 T = short
- 线程 0 从全局内存中读取单词 0
- 线程 1 不读取
- 线程 2 从全局内存中读取单词 1
- 线程 3 不读取
- 等。。。
换句话说，线程K读取单词K/sizeof(T)。这似乎没有正确合并。
我考虑的另一种方法是：
- 启动时线程数 =
(N + 3) / 4，这样每个线程将负责加载和处理4/sizeof(T)元素(每个线程处理一个 32 位字 - 根据sizeof(T)，可能是 1、2 或 4 个元素(。但是，我担心这种方法不会尽可能快，因为每个线程必须执行两次(如果T=short(甚至四倍(如果T=unsigned char(处理量。
有人可以告诉我我对我的计划的假设是否正确，即它不会正确合并吗？
你能评论一下我的替代方法吗？
您能否推荐一种正确合并的更优化的方法？
你是对的，你必须做至少 32 位大小的加载才能合并，并且你描述的方案(让每个其他线程做一个负载(不会合并。 只需将偏移量右移 2 位，让每个线程执行连续的 32 位加载，并使用条件代码来禁止执行将在超出范围地址上运行的线程。
由于您的目标是 SM 1.x，因此还请注意 1( 为了进行合并，给定 warp(32 个线程的集合(的线程 0 必须分别与 4、8 和 16 字节操作数对齐 64、128 或 256 字节，以及 2( 一旦您的数据在共享内存中，您可能希望将循环展开 2 倍(简称(或 4 倍(字符(，以便相邻线程引用相邻的 32 位字， 以避免共享内存库冲突。