CUDA Pro：通过向量化内存访问提高性能

CUDA Pro：通过向量化内存访问提高性能
许多CUDA内核受带宽限制，而新硬件中触发器与带宽的比率不断提高，导致带宽受限制的内核更多。这使得采取措施减轻代码中的带宽瓶颈非常重要。本文将展示如何在CUDA C / C ++中使用向量加载和存储，以帮助提高带宽利用率，同时减少已执行指令的数量。
从以下简单的内存复制内核开始。
global void device_copy_scalar_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
d_out[i] = d_in[i];
}
}

void device_copy_scalar(int* d_in, int* d_out, int N)
{
int threads = 128;
int blocks = min((N + threads-1) / threads, MAX_BLOCKS);
device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
代码使用的是网格跨度循环。图1显示了内核吞吐量（GB / s）与副本大小的关系。

图1：复制带宽与复制大小的关系。
可以使用CUDA Toolkit 附带的cuobjdump工具检查该内核的程序集。
％> cuobjdump -sass可执行文件
标量复制内核主体的SASS如下：
/ * 0058 * / IMAD R6.CC，R0，R9，c [0x0] [0x140]
/ * 0060 * / IMAD.HI.X R7，R0，R9，c [0x0] [0x144]
/ * 0068 * / IMAD R4.CC，R0，R9，c [0x0] [0x148]
/ * 0070 * / LD.E R2，[R6]
/ * 0078 * / IMAD.HI.X R5，R0，R9，c [0x0] [0x14c]
/ * 0090 * / ST.E [R4]，R2
可以看到总共六个与复制操作相关的指令。四个IMAD指令计算加载和存储地址和LD.E与ST.E负载位和32位来自这些地址存储。
可以使用向量化的加载和存储指令LD.E.{64,128}和来提高此操作的性能ST.E.{64,128}。这些操作也可以加载和存储数据，但可以64位或128位宽度进行加载和存储。使用矢量化负载减少了指令总数，减少了等待时间，并提高了带宽利用率。
使用矢量载荷的最简单的方法是使用在CUDA C / C ++标准头中定义的向量的数据类型，如int2，int4，或 float2。可以通过C / C ++中的类型转换轻松地使用这些类型。例如，在C ++可以重铸int指针d_in到一个int2使用指针reinterpret_cast<int2*>(d_in)。在C99中，可以使用强制转换运算符做相同的事情：(int2*(d_in))。
取消引用那些指针将导致编译器生成矢量化指令。但是，有一个重要警告：这些指令需要对齐的数据。设备分配的内存会自动对齐到数据类型大小的倍数，但是如果偏移指针，则偏移也必须对齐。例如reinterpret_cast<int2*>(d_in+1)，无效是因为d_in+1未与对齐sizeof(int2)。
如果使用“对齐”偏移量，则可以安全地偏移数组，如 reinterpret_cast<int2*>(d_in+2)中所示。也可以使用结构生成矢量化载荷，只要该结构的大小为2个字节即可。
struct Foo {int a，int b，double c}; // 16个字节
Foo * x，* y;
…
x [i] = y [i];
既然已经看到了如何生成向量化指令，那么让修改内存复制内核以使用向量加载。
global void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
}

// in only one thread, process final element (if there is one)
if (idxN/2 && N%21)
d_out[N-1] = d_in[N-1];
}

void device_copy_vector2(int* d_in, int* d_out, int n) {
threads = 128;
blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS);

device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
该内核只有几处更改。首先，循环现在仅执行N/ 2次，因为每次迭代处理两个元素。其次，在副本中使用上述技术。第三，处理所有可能N被2整除的剩余元素。最后，启动的线程数量是标量内核中数量的一半。
检查SASS，看到以下内容。
/ * 0088 * / IMAD R10.CC，R3，R5，c [0x0] [0x140]
/ * 0090 * / IMAD.HI.X R11，R3，R5，c [0x0] [0x144]
/ * 0098 * / IMAD R8.CC，R3，R5，c [0x0] [0x148]
/ * 00a0 * / LD.E.64 R6，[R10]
/ * 00a8 * / IMAD.HI.X R9，R3，R5，c [0x0] [0x14c]
/ * 00c8 * / ST.E.64 [R8]，R6
编译器生成LD.E.64和ST.E.64。其他所有指令均相同。由于循环仅执行N / 2次，因此将执行一半的指令。在指令绑定或延迟绑定的内核中，指令数量的2倍改进非常重要。
还可以编写复制内核的vector4版本。
_global void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
}

// in only one thread, process final elements (if there are any)
int remainder = N%4;
if (idx==N/4 && remainder!=0) {
while(remainder) {
int idx = N - remainder–;
d_out[idx] = d_in[idx];
}
}
}

void device_copy_vector4(int* d_in, int* d_out, int N) {
int threads = 128;
int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);

device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
相应的SASS是：
/0090/ IMAD R10.CC, R3, R13, c[0x0][0x140]
/0098/ IMAD.HI.X R11, R3, R13, c[0x0][0x144]
/00a0/ IMAD R8.CC, R3, R13, c[0x0][0x148]
/00a8/ LD.E.128 R4, [R10]
/00b0/ IMAD.HI.X R9, R3, R13, c[0x0][0x14c]
/00d0/ ST.E.128 [R8], R4
在这里可以看到生成的LD.E.128和ST.E.128。此版本的代码将指令数减少了4倍。可以在图2中看到所有3个内核的整体性能。

图2：矢量化内核的复制带宽与复制大小的关系。
在几乎所有情况下，矢量化载荷都优于标量载荷。但是请注意，使用矢量化负载会增加寄存器压力并降低总体并行度。因此，如果的内核已经受到寄存器限制或并行度很低，则可能需要坚持标量加载。同样，如前所述，如果指针未对齐或以字节为单位的数据类型大小不是2的幂，则不能使用矢量化加载。
向量化加载是应该尽可能使用的基本CUDA优化，因为它们会增加带宽，减少指令数量并减少延迟。本文展示了如何通过较少的更改就可以轻松地将向量化负载合并到现有内核中。

CUDA Pro：通过向量化内存访问提高性能相关推荐

深度解读Facebook刚开源的beringei时序数据库——数据压缩delta of delta+充分利用内存以提高性能...
转自:https://yq.aliyun.com/topic/58?spm=5176.100239.blogcont69354.9.MLtp4T 摘要: Facebook最近开源了beringei时序 ...
python3 性能提升_整理下Python性能语法,非常有效的提高性能的tips
阅读 Zen of Python,在Python解析器中输入 import this. 一个犀利的Python新手可能会注意到"解析"一词, 认为Python不过是另一门脚本语言. ...
面试官问：Integer 如何实现节约内存和提升性能的？
点击上方"方志朋",选择"设为星标" 回复"666"获取新整理的面试资料作者:Byte_Liu 来源:https://urlify.cn/ ...
7. CUDA内存访问（一）提高篇------按部就班 ------GPU的革命
序言:从上一篇< CUDA编程接口(二)------一十八般武器>到现在,差不多有三个月了,不知道大家在"暑假"里面过得怎么样,又经历了什么?花了两个星期的睡觉前的时间 ...
cuda合并访问的要求_【CUDA 基础】4.3 内存访问模式
Abstract: 本文介绍内存的访问过程,也就是从应用发起请求到硬件实现的完整操作过程,这里是优化内存瓶颈的关键之处,也是CUDA程序优化的基础.Keywords: 内存访问模式,对齐,合并,缓存, ...
CUDA的global内存访问的问题
http://blog.csdn.net/OpenHero/article/details/3520578 关于CUDA的global内存访问的问题,怎么是访问的冲突,怎样才能更好的访问内存,达到更高 ...
Java内存访问重排序的研究
什么是重排序请先看这样一段代码1 public class PossibleReordering { static int x = 0, y = 0; static int a = 0, b = 0 ...
cuda二维数组内存分配和数据拷贝
uda二维数组内存分配和数据拷贝 2016-04-20 10:54 138人阅读评论(0) 收藏举报分类: 机器学习(11) 人工智能(9) 版权声明:本文为博主原创文章,允许转载. 因为cud ...
java项目怎样提高性能_Java程序员成长之路(如何提高Java程序性能？)
1.尽量在合适的场合使用单例使用单例可以减轻加载的负担,缩短加载的时间,提高加载的效率,但并不是所有地方都适用于单例,简单来说,单例主要适用于以下三个方面第一,控制资源的使用,通过线程同步来控制资 ...

CUDA Pro：通过向量化内存访问提高性能

CUDA Pro：通过向量化内存访问提高性能相关推荐

最新文章

热门文章