我刚开始在 CUDA 中编写代码，我正试图了解线程如何执行和内存访问的概念，以便充分利用 GPU。我通读了 CUDA 最佳实践指南、CUDA by Example 一书以及这里的几篇文章。我还发现 Mark Harris 的简化示例非常有趣和有用，但尽管有所有信息，但我对细节感到相当困惑。
假设我们有一个大型二维数组 (N*M)，我们在其上进行列操作。我将数组拆分为块，以便每个块都有多个线程数，这些线程数是 32 的倍数（所有线程都适合多个线程）。每个块中的第一个线程分配额外的内存（初始数组的副本，但仅针对其自身维度的大小）并使用_shared_共享指针变量，以便同一块的所有线程可以访问相同的内存。由于线程数是 32 的倍数，因此内存也应该是 32 的倍数，以便在单次读取中访问。但是，我需要在内存块周围有一个额外的填充，一个边框，以便我的数组的宽度变为 (32*x)+2 列。边界来自分解大数组，因此我有一个重叠区域，其中其邻居的副本暂时可用。
共享内存访问：
想象一个块的线程正在访问本地内存块
1 int x = threadIdx.x;
2
3 for (int y = 0; y < height; y++)
4 {
5 double value_centre = array[y*width + x+1]; // remeber we have the border so we need an offset of + 1
6 double value_left = array[y*width + x ]; // hence the left element is at x
7 double value_right = array[y*width + x+2]; // and the right element at x+2
8
9 // .. do something
10 }
现在，我的理解是，由于我确实有一个不可避免的偏移量 (+1,+2)，我将在每个 warp 和每个分配中至少读取两次（左侧元素除外），或者这无关紧要只要第一个线程后的内存完全对齐，我从哪里开始阅读？另请注意，如果不是这种情况，那么我将对第一行之后的每一行的数组进行非对齐访问，因为我的数组的宽度是 (32*x)+2，因此不是 32 字节对齐的。然而，进一步的填充将解决每个新行的问题。
问题：我的理解是否正确，在上面的示例中，只有第一行允许共同访问，并且只允许数组中的左元素，因为这是唯一没有任何偏移量访问的元素？
在扭曲中执行的线程：
只有当且仅当所有指令都相同时（根据link），warp 中的线程才会并行执行。如果我确实有条件语句/发散执行，那么该特定线程将自行执行，而不是与其他线程一起执行。
例如，如果我初始化数组，我可以做类似的事情
1 int x = threadIdx.x;
2
3 array[x+1] = globalArray[blockIdx.x * blockDim.x + x]; // remember the border and therefore use +1
4
5 if (x == 0 || x == blockDim.x-1) // border
6 {
7 array[x] = DBL_MAX;
8 }
扭曲的大小是否为 32 并并行执行直到第 3 行，然后停止所有其他线程，并且只有第一个和最后一个线程进一步执行以初始化边界，或者这些线程是否会在开始时与所有其他线程分开，因为是否有一个所有其他线程都不满足的 if 语句？
问题：如何将线程收集到单个经线中？warp 中的每个线程都需要共享相同的指令。需要这对整个功能有效吗？线程 1 (x=0) 不是这种情况，因为它也初始化边界，因此与其他边界不同。据我了解，线程 1 在单个 warp 中执行，线程（2-33 等）在另一个 warp 中执行，然后由于未对齐而不会在单次读取中访问内存，然后再次访问最终由于另一个边界，在一条经线中穿线。那是对的吗？
我想知道最佳实践是什么，让每一行的内存完全对齐（在这种情况下，我将使用 (32*x-2) 线程运行每个块，以便带边框的数组为 (32*x-2)+ 2 每个新行是 32 的倍数）或者按照我上面演示的方式进行操作，每个块的线程都是 32 的倍数，并且只使用未对齐的内存。我知道这类问题并不总是直截了当，而且通常取决于特定情况，但有时某些事情是一种不好的做法，不应该成为习惯。
当我进行一些实验时，我并没有真正注意到执行时间的差异，但也许我的示例太简单了。我试图从视觉分析器获取信息，但我并没有真正理解它给我的所有信息。然而，我收到警告说我的入住率是 17%，我认为这一定非常低，因此我做错了。我没有找到有关线程如何并行执行以及我的内存访问效率如何的信息。
-编辑-
添加并突出显示了 2 个问题，一个是关于内存访问的，另一个是关于如何将线程收集到单个 warp 的问题。