我在学习this guide的同时正在学习CUDA。
我还没有完成，但是我决定玩一些到目前为止所看到的。
我试图重写第一个使用256个线程的示例。我想这样做，以便每个线程在数组的连续切片上进行操作。
目标是将2个数组与1,048,576个项相加。
为了进行比较，这是原始代码，其中根据跨步访问每个数组项:
这是我的功能:
__global__
void add2(int n, float* x, float* y) {
int sliceSize = n / blockDim.x;
int lower = threadIdx.x * sliceSize;
int upper = lower + sliceSize;
for (int i = lower; i < upper; i++) {
y[i] = x[i] + y[i];
}
}
事实证明，最后一个摘要的执行速度比上一个摘要慢了近7倍(22ms对3ms)。我认为，通过在连续的片上访问它们，它将执行相同或更快的操作。
我正在使用
add<<<1, threads>>>(n, x, y)和add<<<1, threads>>>(n, x, y)(256个线程)调用该函数。sliceSize的值始终为4096。在这种情况下，应该发生的是:threadIdx.x = 0从0到4095 threadIdx.x = 1从4096变为8191 threadIdx.x = 255从1044480变为1048576 我打开了NVidia Visual Profiler，然后了解到我的内存访问模式效率不高(全局内存加载/存储效率低)。第一个代码段中没有此警告。为什么会这样呢？
我以为第一个被删除的对象会在整个数组中跳转，从而造成错误的访问模式。实际上，这似乎很好。
我已经阅读了可视分析器随附的一些有关内存优化的文档，但是我不太明白为什么它这么慢。
最佳答案
您正在探索合并和未合并的内存访问之间的区别。或者我们可以简单地说“最有效”和“效率较低”的内存访问。
在GPU上，所有指令都在整个warp中执行。因此，当warp中的一个线程正在读取内存中的位置时，warp中的所有线程都正在从内存中读取。粗略地说，最佳模式是当经线中的所有线程都从相邻位置读取时。这导致以下情况:GPU内存 Controller 在检查特定时间段内warp中每个线程请求的内存地址后，可以将合并地址在一起，从而导致需要从内存请求的行数最少。高速缓存(或要从DRAM请求的最小段数)。
在幻灯片36(或37)here上以图形方式描绘了这种情况。
第一个代码段中表示了100％合并的大小写。从全局内存读取的示例如下:
y[i] = x[i] + y[i];
^
reading from the vector x in global memory
让我们考虑循环的第一遍，并考虑第一次扭曲的情况(即线程块中的前32个线程)。在这种情况下，
i由threadIdx.x给出。因此，线程0的索引为0，线程1的索引为1，依此类推。因此，每个线程都在读取全局内存中的相邻位置。假设我们错过了所有缓存，这将转化为DRAM读取请求，并且存储器 Controller 可以针对DRAM中的段(或等效地，针对缓存中的行)生成最少数量的请求(更准确地说:事务)。从“总线带宽利用率”为100％的角度来看，这是最佳的。在该读取周期上，请求的每个字节实际上都被扭曲中的线程使用。“不公开”访问通常可以指任何不符合以上描述的情况。转化为上述更细粒度的“总线带宽利用率”数字，根据具体情况和不同情况，非强制访问的程度可能有所不同，从最好的情况(略低于100％)到最坏的情况为12.5％(或3.125％)。 GPU。
在幻灯片44(或45)here中给出了根据此描述的最坏情况的不分时段访问模式示例。这不能完全描述您的最坏情况的代码段，但是对于足够大的
sliceSize来说，它是等效的。代码行是相同的。考虑到相同的读取请求(对于x，在循环的第一次迭代中，对于warp 0，为warp 0)，唯一的区别在于i在整个warp中取值:int sliceSize = n / blockDim.x;
int lower = threadIdx.x * sliceSize;
...
for (int i = lower; i < upper; i++) {
y[i] = x[i] + y[i];
因此，
i从lower开始，即threadIdx.x * sliceSize。假设sliceSize大于1。然后，第一个线程将读取位置0。第二个线程将读取位置sliceSize。第三个线程将读取位置2*sliceSize等。这些位置由sliceSize距离分隔。即使sliceSize只有2，该模式的效率仍然较低，因为内存 Controller 现在必须请求两倍的行数或段数才能满足遍历0的特定读取周期。如果sliceSize足够大，则内存 Controller 必须请求a每个线程的唯一行或段，这是最坏的情况。作为最后的注释/总结，可以对“快速分析”做出有益的观察:
threadIdx.x乘以除1以外的任何数量。threadIdx.x乘以不等于1的某个数字，则无论其他考虑如何，这几乎都是普遍的指示，即所生成的访问模式将不是最佳的。 为了清楚起见，请重复此操作:
index = any_constant_across_the_warp + threadIdx.x;
通常将是最佳的访问模式。
index = any_constant_across_the_warp + C*threadIdx.x;
通常将不是最佳的访问模式。注意
any_constant_across_the_warp可以由数量上的任意算术组成，例如:循环索引，blockIdx.?，blockDim.?，gridDim.?和任何其他常量。必须考虑到2D或3D线程块模式，其中将考虑threadIdx.y，但通常不难将这种理解扩展到2D情况。对于典型的线程块形状，为了进行快速分析，通常不希望在threadIdx.x或threadIdx.y上使用常数乘数。整个讨论适用于全局内存读/写。共享内存还具有用于最佳访问的规则，这些规则在某种程度上类似于上述说明，但在某些方面却大不相同。但是，通常的情况是，针对全局内存的完全最佳的100％合并模式也将是共享内存读/写的最佳模式。换句话说，扭曲中的相邻访问通常对于共享内存也是最佳的(但它不是共享内存的唯一可能的最佳模式)。
与here链接的演示文稿将对此主题进行更全面的处理，网络上的许多其他演示文稿和处理也将对此进行更全面的处理。