我刚开始学习CUDA编程，试图理解如何执行线程和访问内存以充分利用GPU。我阅读了CUDA最佳实践指南、《CUDA示例》一书和这里的几篇帖子。我还发现Mark Harris的归约示例非常有趣和有用，但是尽管我获得了所有信息，仍然对细节感到困惑。
假设我们有一个大的二维数组（N*M），我们对其进行逐列操作。我将数组分成块，使得每个块都有多个线程，这些线程的数量是32的倍数（所有线程都适合于几个warp中）。每个块中的第一个线程分配额外的内存（初始数组的副本，但仅为自身维度的大小），并使用_shared _变量共享指针，以便同一块中的所有线程都可以访问相同的内存。由于线程数是32的倍数，因此内存也应该是32的倍数，才能在单次读取中被访问。但是，我需要在内存块周围增加额外的填充边框，使我的数组宽度变为（32*x）+2列。边框来自于分解大数组，因此我有重叠区域，在其中邻居的副本暂时可用。
现在，我的理解是，由于我有一个偏移量（+1，+2），这是不可避免的，每个warp和每个赋值至少会有两个读取（除了左侧元素），或者无论从哪里开始读取都没有关系，只要第一个线程后面的内存完全对齐即可。另请注意，如果不是这种情况，那么对于第一行之后的每一行，我将对数组进行未对齐访问，因为我的数组宽度为（32 * x）+2，因此不是32字节对齐的。进一步的填充可以解决每一行的问题。
问题：我的理解是否正确，以上示例中只有第一行才允许联合访问，而且仅限于数组中的左侧元素，因为这是唯一一个没有任何偏移量访问的元素？
在一个warp中执行的线程：
当且仅当所有指令相同时（根据link），warp中的线程才会并行执行。如果我有一个条件语句/分支执行，那么该特定线程将单独执行，而不是与其他线程一起在warp中执行。
例如，如果我初始化数组，我可以这样做：
当我进行了一些实验时，我没有真正注意到执行时间上的差异，但也许我的例子太简单了。我试图从可视化分析器中获取信息，但我并没有真正理解它给我提供的所有信息。然而，我收到了一个警告，说我的占用率为17％，我认为这一定非常低，因此我做错了什么。我没有成功地找到关于如何并行执行线程和我的内存访问效率的信息。
-编辑-
添加并突出显示了两个问题，一个关于内存访问，另一个关于如何将线程收集到单个warp中。
假设我们有一个大的二维数组（N*M），我们对其进行逐列操作。我将数组分成块，使得每个块都有多个线程，这些线程的数量是32的倍数（所有线程都适合于几个warp中）。每个块中的第一个线程分配额外的内存（初始数组的副本，但仅为自身维度的大小），并使用_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字节对齐的。进一步的填充可以解决每一行的问题。
问题：我的理解是否正确，以上示例中只有第一行才允许联合访问，而且仅限于数组中的左侧元素，因为这是唯一一个没有任何偏移量访问的元素？
在一个warp中执行的线程：
当且仅当所有指令相同时（根据link），warp中的线程才会并行执行。如果我有一个条件语句/分支执行，那么该特定线程将单独执行，而不是与其他线程一起在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 }
这个warp的大小是32吗？并且在第三行之前以并行方式执行，然后停止所有其他线程，只有第一个和最后一个线程继续执行以初始化边框，还是这些线程已经在一开始就与所有其他线程分开了，因为有一个if语句，所有其他线程都不符合条件？
问题：如何将线程收集到单个warp中？warp中的每个线程需要共享相同的指令。需要整个函数都有效吗？这对于线程1（x=0）不适用，因为它还初始化了边框，因此与其他线程不同。据我理解，线程1在单个warp中执行，线程（2-33等）在另一个warp中执行，由于未对齐，因此不会在单次读取中访问内存，然后再次在单个warp中执行最终线程，因为存在另一个边框。这正确吗？
我想知道最佳实践是什么，是为每一行都完美对齐内存（在这种情况下，我将使用（32*x-2）个线程运行每个块，以便带有边框的数组为（32*x-2）+2的倍数，以适应每一行），还是像我上面演示的那样，使用每个块的线程为32的倍数，并且只接受不对齐的内存。我知道这些问题通常并不直接，而且往往要视具体情况而定，但有时候某些事情是不好的做法，不应该成为习惯。当我进行了一些实验时，我没有真正注意到执行时间上的差异，但也许我的例子太简单了。我试图从可视化分析器中获取信息，但我并没有真正理解它给我提供的所有信息。然而，我收到了一个警告，说我的占用率为17％，我认为这一定非常低，因此我做错了什么。我没有成功地找到关于如何并行执行线程和我的内存访问效率的信息。
-编辑-
添加并突出显示了两个问题，一个关于内存访问，另一个关于如何将线程收集到单个warp中。