如何有效地从内核中访问设备的全局内存
在前面的 两 文章 中，我们研究了如何在主机和设备之间高效地移动数据。在我们的 cuda c / c ++系列的第六篇文章中，我们将讨论如何有效地从内核中访问设备存储器，特别是全局内存 。
在 cuda 设备上有几种内存，每种内存的作用域、生存期和缓存行为都不同。到目前为止，在本系列中，我们已经使用了驻留在设备 dram 中的 全局内存 ，用于主机和设备之间的传输，以及内核的数据输入和输出。这里的名称 global 是指作用域，因为它可以从主机和设备访问和修改。全局内存可以像下面代码片段的第一行那样使用 __device__ de clara 说明符在全局（变量）范围内声明，或者使用 cudamalloc() 动态分配并分配给一个常规的 c 指针变量，如第 7 行所示。全局内存分配可以在应用程序的生命周期内保持。根据设备的 计算能力 ，全局内存可能被缓存在芯片上，也可能不在芯片上缓存。
__device__ int globalarray[256];void foo(){ ... int *mydevicememory = 0; cudaerror_t result = cudamalloc(&mydevicememory, 256 * sizeof(int)); ... }在讨论全局内存访问性能之前，我们需要改进对 cuda 执行模型的理解。我们已经讨论了如何将 线程被分组为线程块 分配给设备上的多处理器。在执行过程中，有一个更精细的线程分组到 warps 。 gpu 上的多处理器以 simd （ 单指令多数据 ）方式为每个扭曲执行指令。所有当前支持 cuda – 的 gpus 的翘曲尺寸（实际上是 simd 宽度）是 32 个线程。
全局内存合并 将线程分组为扭曲不仅与计算有关，而且与全局内存访问有关。设备 coalesces 全局内存加载并存储由一个 warp 线程发出的尽可能少的事务，以最小化 dram 带宽（在计算能力小于 2 . 0 的旧硬件上，事务合并在 16 个线程的一半扭曲内，而不是整个扭曲中）。为了弄清楚 cuda 设备架构中发生聚结的条件，我们在三个 tesla 卡上进行了一些简单的实验： a tesla c870 （计算能力 1 . 0 ）、 tesla c1060 （计算能力 1 . 3 ）和 tesla c2050 （计算能力 2 . 0 ）。
我们运行两个实验，使用如下代码（ github 上也有 ）中所示的增量内核的变体，一个具有数组偏移量，这可能导致对输入数组的未对齐访问，另一个是对输入数组的跨步访问。
#include#include// convenience function for checking cuda runtime api results// can be wrapped around any runtime api call. no-op in release builds.inlinecudaerror_t checkcuda(cudaerror_t result){#if defined(debug) || defined(_debug) if (result != cudasuccess) { fprintf(stderr, cuda runtime error: %sn, cudageterrorstring(result)); assert(result == cudasuccess); }#endif return result;}template__global__ void offset(t* a, int s){ int i = blockdim.x * blockidx.x + threadidx.x + s; a[i] = a[i] + 1;}template__global__ void stride(t* a, int s){ int i = (blockdim.x * blockidx.x + threadidx.x) * s; a[i] = a[i] + 1;}templatevoid runtest(int deviceid, int nmb){ int blocksize = 256; float ms; t *d_a; cudaevent_t startevent, stopevent; int n = nmb*1024*1024/sizeof(t); // nb: d_a(33*nmb) for stride case checkcuda( cudamalloc(&d_a, n * 33 * sizeof(t)) ); checkcuda( cudaeventcreate(&startevent) ); checkcuda( cudaeventcreate(&stopevent) ); printf(offset, bandwidth (gb/s):n); offset(d_a, 0); // warm up for (int i = 0; i <= 32; i++) { checkcuda( cudamemset(d_a, 0.0, n * sizeof(t)) ); checkcuda( cudaeventrecord(startevent,0) ); offset(d_a, i); checkcuda( cudaeventrecord(stopevent,0) ); checkcuda( cudaeventsynchronize(stopevent) ); checkcuda( cudaeventelapsedtime(&ms, startevent, stopevent) ); printf(%d, %fn, i, 2*nmb/ms); } printf(n); printf(stride, bandwidth (gb/s):n); stride(d_a, 1); // warm up for (int i = 1; i <= 32; i++) { checkcuda( cudamemset(d_a, 0.0, n * sizeof(t)) ); checkcuda( cudaeventrecord(startevent,0) ); stride(d_a, i); checkcuda( cudaeventrecord(stopevent,0) ); checkcuda( cudaeventsynchronize(stopevent) ); checkcuda( cudaeventelapsedtime(&ms, startevent, stopevent) ); printf(%d, %fn, i, 2*nmb/ms); } checkcuda( cudaeventdestroy(startevent) ); checkcuda( cudaeventdestroy(stopevent) ); cudafree(d_a);}int main(int argc, char **argv){ int nmb = 4; int deviceid = 0; bool bfp64 = false; for (int i = 1; i < argc; i++) { if (!strncmp(argv[i], dev=, 4)) deviceid = atoi((char*)(&argv[i][4])); else if (!strcmp(argv[i], fp64)) bfp64 = true; } cudadeviceprop prop; checkcuda( cudasetdevice(deviceid) ) ; checkcuda( cudagetdeviceproperties(&prop, deviceid) ); printf(device: %sn, prop.name); printf(transfer size (mb): %dn, nmb); printf(%s precisionn, bfp64 ? double : single); if (bfp64) runtest(deviceid, nmb); else runtest(deviceid, nmb);
}此代码可以通过传递“ fp64 ”命令行选项以单精度（默认值）或双精度运行偏移量内核和跨步内核。每个内核接受两个参数，一个输入数组和一个表示访问数组元素的偏移量或步长的整数。内核在一系列偏移和跨距的循环中被称为。
未对齐的数据访问 下图显示了 tesla c870 、 c1060 和 c2050 上的偏移内核的结果。
设备内存中分配的数组由 cuda 驱动程序与 256 字节内存段对齐。该设备可以通过 32 字节、 64 字节或 128 字节的事务来访问全局内存。对于 c870 或计算能力为 1 . 0 的任何其他设备，半线程的任何未对齐访问（或半扭曲线程不按顺序访问内存的对齐访问）将导致 16 个独立的 32 字节事务。由于每个 32 字节事务只请求 4 个字节，因此可以预期有效带宽将减少 8 倍，这与上图（棕色线）中看到的偏移量（不是 16 个元素的倍数）大致相同，对应于线程的一半扭曲。
对于计算能力为 1 . 2 或 1 . 3 的 tesla c1060 或其他设备，未对准访问的问题较少。基本上，通过半个线程对连续数据的未对齐访问在几个“覆盖”请求的数据的事务中提供服务。由于未请求的数据正在传输，以及不同的半翘曲所请求的数据有些重叠，因此相对于对齐的情况仍然存在性能损失，但是这种损失远远小于 c870 。
计算能力为 2 . 0 的设备，如 tesla c250 ，在每个多处理器中都有一个 l1 缓存，其行大小为 128 字节。该设备将线程的访问合并到尽可能少的缓存线中，从而导致对齐对跨线程顺序内存访问吞吐量的影响可以忽略不计。
快速内存访问 步幅内核的结果如下图所示。
对于快速的全局内存访问，我们有不同的看法。对于大步进，无论架构版本如何，有效带宽都很差。这并不奇怪：当并发线程同时访问物理内存中相距很远的内存地址时，硬件就没有机会合并这些访问。从上图中可以看出，在 tesla c870 上，除 1 以外的任何步幅都会导致有效带宽大幅降低。这是因为 compute capability 1 . 0 和 1 . 1 硬件需要跨线程进行线性、对齐的访问以进行合并，因此我们在 offset 内核中看到了熟悉的 1 / 8 带宽。 compute capability 1 . 2 及更高版本的硬件可以将访问合并为对齐的段（ cc 1 . 2 / 1 . 3 上为 32 、 64 或 128 字节段，在 cc 2 . 0 及更高版本上为 128 字节缓存线），因此该硬件可以产生平滑的带宽曲线。
当访问多维数组时，线程通常需要索引数组的更高维，因此快速访问是不可避免的。我们可以使用一种名为 共享内存 的 cuda 内存来处理这些情况。共享内存是一个线程块中所有线程共享的片上内存。共享内存的一个用途是将多维数组的 2d 块以合并的方式从全局内存提取到共享内存中，然后让连续的线程跨过共享内存块。与全局内存不同，对共享内存的快速访问没有惩罚。我们将在下一篇文章中详细介绍共享内存。
概括 在这篇文章中，我们讨论了如何从 cuda 内核代码中有效地访问全局内存的一些方面。设备上的全局内存访问与主机上的数据访问具有相同的性能特征，即数据局部性非常重要。在早期的 cuda 硬件中，内存访问对齐和跨线程的局部性一样重要，但在最近的硬件上，对齐并不是什么大问题。另一方面，快速的内存访问会损害性能，使用片上共享内存可以减轻这种影响。在 下一篇文章 中，我们将详细探讨共享内存，之后的文章中，我们将展示如何使用共享内存来避免在 矩阵转置 过程中出现跨步全局内存访问。
关于作者
mark harris 是 nvidia 杰出的工程师，致力于 rapids 。 mark 拥有超过 20 年的 gpus 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： gpgpu （图形处理单元上的通用计算）。
VF转换电路计算三角波的高低阈值
ITO PET技术,ITO PET技术原理是什么
IHS最新预测：全球处理器出货量年底将达15亿片
三星李在镕被限制就业，会伤及其物联网布局吗？
Xlinx FPGA的DSP设计工具和设计流程
如何有效地从内核中访问设备的全局内存
高端黑科技华为nova 2 Plus魔镜版美得太不像话了！
苹果全新配色：渐变机身+5G网络+3860毫安电池
人工智能应该学什么专业
四种新型现代水产生态养殖模式详解
应用于低速应用所需的 DAC的详细分析
周志华：AI伦理研究将就加入南京大学人工智能学院
LG OLED 8k电视价值30多万 贵在这些地方你们知道吗
显像管灯丝保护电路图
华硕灵珑III笔记本重新定义轻薄本
智慧校园创造特色管理新模式！
Instagram将通过AI筛查检测欺凌性言语，以对抗网络欺凌行为
cpu发热如何解决 如何正确选择CPU
全面解读功能安全之暴露率的实例和解释
陶瓷耳机价格不菲但却备受年轻人追捧，这是为何
在 cuda 设备上有几种内存，每种内存的作用域、生存期和缓存行为都不同。到目前为止，在本系列中，我们已经使用了驻留在设备 dram 中的 全局内存 ，用于主机和设备之间的传输，以及内核的数据输入和输出。这里的名称 global 是指作用域，因为它可以从主机和设备访问和修改。全局内存可以像下面代码片段的第一行那样使用 __device__ de clara 说明符在全局（变量）范围内声明，或者使用 cudamalloc() 动态分配并分配给一个常规的 c 指针变量，如第 7 行所示。全局内存分配可以在应用程序的生命周期内保持。根据设备的 计算能力 ，全局内存可能被缓存在芯片上，也可能不在芯片上缓存。
__device__ int globalarray[256];void foo(){ ... int *mydevicememory = 0; cudaerror_t result = cudamalloc(&mydevicememory, 256 * sizeof(int)); ... }在讨论全局内存访问性能之前，我们需要改进对 cuda 执行模型的理解。我们已经讨论了如何将 线程被分组为线程块 分配给设备上的多处理器。在执行过程中，有一个更精细的线程分组到 warps 。 gpu 上的多处理器以 simd （ 单指令多数据 ）方式为每个扭曲执行指令。所有当前支持 cuda – 的 gpus 的翘曲尺寸（实际上是 simd 宽度）是 32 个线程。
全局内存合并 将线程分组为扭曲不仅与计算有关，而且与全局内存访问有关。设备 coalesces 全局内存加载并存储由一个 warp 线程发出的尽可能少的事务，以最小化 dram 带宽（在计算能力小于 2 . 0 的旧硬件上，事务合并在 16 个线程的一半扭曲内，而不是整个扭曲中）。为了弄清楚 cuda 设备架构中发生聚结的条件，我们在三个 tesla 卡上进行了一些简单的实验： a tesla c870 （计算能力 1 . 0 ）、 tesla c1060 （计算能力 1 . 3 ）和 tesla c2050 （计算能力 2 . 0 ）。
我们运行两个实验，使用如下代码（ github 上也有 ）中所示的增量内核的变体，一个具有数组偏移量，这可能导致对输入数组的未对齐访问，另一个是对输入数组的跨步访问。
#include#include// convenience function for checking cuda runtime api results// can be wrapped around any runtime api call. no-op in release builds.inlinecudaerror_t checkcuda(cudaerror_t result){#if defined(debug) || defined(_debug) if (result != cudasuccess) { fprintf(stderr, cuda runtime error: %sn, cudageterrorstring(result)); assert(result == cudasuccess); }#endif return result;}template__global__ void offset(t* a, int s){ int i = blockdim.x * blockidx.x + threadidx.x + s; a[i] = a[i] + 1;}template__global__ void stride(t* a, int s){ int i = (blockdim.x * blockidx.x + threadidx.x) * s; a[i] = a[i] + 1;}templatevoid runtest(int deviceid, int nmb){ int blocksize = 256; float ms; t *d_a; cudaevent_t startevent, stopevent; int n = nmb*1024*1024/sizeof(t); // nb: d_a(33*nmb) for stride case checkcuda( cudamalloc(&d_a, n * 33 * sizeof(t)) ); checkcuda( cudaeventcreate(&startevent) ); checkcuda( cudaeventcreate(&stopevent) ); printf(offset, bandwidth (gb/s):n); offset(d_a, 0); // warm up for (int i = 0; i <= 32; i++) { checkcuda( cudamemset(d_a, 0.0, n * sizeof(t)) ); checkcuda( cudaeventrecord(startevent,0) ); offset(d_a, i); checkcuda( cudaeventrecord(stopevent,0) ); checkcuda( cudaeventsynchronize(stopevent) ); checkcuda( cudaeventelapsedtime(&ms, startevent, stopevent) ); printf(%d, %fn, i, 2*nmb/ms); } printf(n); printf(stride, bandwidth (gb/s):n); stride(d_a, 1); // warm up for (int i = 1; i <= 32; i++) { checkcuda( cudamemset(d_a, 0.0, n * sizeof(t)) ); checkcuda( cudaeventrecord(startevent,0) ); stride(d_a, i); checkcuda( cudaeventrecord(stopevent,0) ); checkcuda( cudaeventsynchronize(stopevent) ); checkcuda( cudaeventelapsedtime(&ms, startevent, stopevent) ); printf(%d, %fn, i, 2*nmb/ms); } checkcuda( cudaeventdestroy(startevent) ); checkcuda( cudaeventdestroy(stopevent) ); cudafree(d_a);}int main(int argc, char **argv){ int nmb = 4; int deviceid = 0; bool bfp64 = false; for (int i = 1; i < argc; i++) { if (!strncmp(argv[i], dev=, 4)) deviceid = atoi((char*)(&argv[i][4])); else if (!strcmp(argv[i], fp64)) bfp64 = true; } cudadeviceprop prop; checkcuda( cudasetdevice(deviceid) ) ; checkcuda( cudagetdeviceproperties(&prop, deviceid) ); printf(device: %sn, prop.name); printf(transfer size (mb): %dn, nmb); printf(%s precisionn, bfp64 ? double : single); if (bfp64) runtest(deviceid, nmb); else runtest(deviceid, nmb);
}此代码可以通过传递“ fp64 ”命令行选项以单精度（默认值）或双精度运行偏移量内核和跨步内核。每个内核接受两个参数，一个输入数组和一个表示访问数组元素的偏移量或步长的整数。内核在一系列偏移和跨距的循环中被称为。
未对齐的数据访问 下图显示了 tesla c870 、 c1060 和 c2050 上的偏移内核的结果。
设备内存中分配的数组由 cuda 驱动程序与 256 字节内存段对齐。该设备可以通过 32 字节、 64 字节或 128 字节的事务来访问全局内存。对于 c870 或计算能力为 1 . 0 的任何其他设备，半线程的任何未对齐访问（或半扭曲线程不按顺序访问内存的对齐访问）将导致 16 个独立的 32 字节事务。由于每个 32 字节事务只请求 4 个字节，因此可以预期有效带宽将减少 8 倍，这与上图（棕色线）中看到的偏移量（不是 16 个元素的倍数）大致相同，对应于线程的一半扭曲。
对于计算能力为 1 . 2 或 1 . 3 的 tesla c1060 或其他设备，未对准访问的问题较少。基本上，通过半个线程对连续数据的未对齐访问在几个“覆盖”请求的数据的事务中提供服务。由于未请求的数据正在传输，以及不同的半翘曲所请求的数据有些重叠，因此相对于对齐的情况仍然存在性能损失，但是这种损失远远小于 c870 。
计算能力为 2 . 0 的设备，如 tesla c250 ，在每个多处理器中都有一个 l1 缓存，其行大小为 128 字节。该设备将线程的访问合并到尽可能少的缓存线中，从而导致对齐对跨线程顺序内存访问吞吐量的影响可以忽略不计。
快速内存访问 步幅内核的结果如下图所示。
对于快速的全局内存访问，我们有不同的看法。对于大步进，无论架构版本如何，有效带宽都很差。这并不奇怪：当并发线程同时访问物理内存中相距很远的内存地址时，硬件就没有机会合并这些访问。从上图中可以看出，在 tesla c870 上，除 1 以外的任何步幅都会导致有效带宽大幅降低。这是因为 compute capability 1 . 0 和 1 . 1 硬件需要跨线程进行线性、对齐的访问以进行合并，因此我们在 offset 内核中看到了熟悉的 1 / 8 带宽。 compute capability 1 . 2 及更高版本的硬件可以将访问合并为对齐的段（ cc 1 . 2 / 1 . 3 上为 32 、 64 或 128 字节段，在 cc 2 . 0 及更高版本上为 128 字节缓存线），因此该硬件可以产生平滑的带宽曲线。
当访问多维数组时，线程通常需要索引数组的更高维，因此快速访问是不可避免的。我们可以使用一种名为 共享内存 的 cuda 内存来处理这些情况。共享内存是一个线程块中所有线程共享的片上内存。共享内存的一个用途是将多维数组的 2d 块以合并的方式从全局内存提取到共享内存中，然后让连续的线程跨过共享内存块。与全局内存不同，对共享内存的快速访问没有惩罚。我们将在下一篇文章中详细介绍共享内存。
概括 在这篇文章中，我们讨论了如何从 cuda 内核代码中有效地访问全局内存的一些方面。设备上的全局内存访问与主机上的数据访问具有相同的性能特征，即数据局部性非常重要。在早期的 cuda 硬件中，内存访问对齐和跨线程的局部性一样重要，但在最近的硬件上，对齐并不是什么大问题。另一方面，快速的内存访问会损害性能，使用片上共享内存可以减轻这种影响。在 下一篇文章 中，我们将详细探讨共享内存，之后的文章中，我们将展示如何使用共享内存来避免在 矩阵转置 过程中出现跨步全局内存访问。
关于作者
mark harris 是 nvidia 杰出的工程师，致力于 rapids 。 mark 拥有超过 20 年的 gpus 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： gpgpu （图形处理单元上的通用计算）。
VF转换电路计算三角波的高低阈值
ITO PET技术,ITO PET技术原理是什么
IHS最新预测：全球处理器出货量年底将达15亿片
三星李在镕被限制就业，会伤及其物联网布局吗？
Xlinx FPGA的DSP设计工具和设计流程
如何有效地从内核中访问设备的全局内存
高端黑科技华为nova 2 Plus魔镜版美得太不像话了！
苹果全新配色：渐变机身+5G网络+3860毫安电池
人工智能应该学什么专业
四种新型现代水产生态养殖模式详解
应用于低速应用所需的 DAC的详细分析
周志华：AI伦理研究将就加入南京大学人工智能学院
LG OLED 8k电视价值30多万 贵在这些地方你们知道吗
显像管灯丝保护电路图
华硕灵珑III笔记本重新定义轻薄本
智慧校园创造特色管理新模式！
Instagram将通过AI筛查检测欺凌性言语，以对抗网络欺凌行为
cpu发热如何解决 如何正确选择CPU
全面解读功能安全之暴露率的实例和解释
陶瓷耳机价格不菲但却备受年轻人追捧，这是为何