6.6_映射锁页内存
6.6 映射锁页内存
映射锁页内存可用于PCIe传输和内核处理的重叠执行,尤其是对于设备端到主机端的复制(这里不需要隐藏太长时间主机内存的延迟)。映射锁页内存比原生GPU内存复制有更严格的对齐要求,因为它们必须被合并读取。使用非合并的内存事务处理会比使用映射锁页内存慢2~6倍。
我们对concurrencyMemcpyKernelMapped.cu程序作一简单移植会产生一个有趣的结果:在亚马逊EC2的cgl.4xlarge实例上,映射锁页内存在cycles低于64的情况下运行非常缓慢。
对于小的cycles,内核需要很长的时间来运行,就好像cycles大于200一样。虽然只有英伟达能发现这种表现异常的确切原因,但是它
并不难解决:通过展开内核的内部循环,为每个线程创造更多的工作,提高性能。
请注意,代码清单6-6中这个版本的AddKernel()与代码清单6-3中的那个功能相同[1]。它只是计算每个循环迭代的unrollFactor输出。既然展开因子是一个模板参数,编译器可以使用寄存器来保存value数组,并且最内层的循环可以完全展开。
代码清单6-6 循环展开的AddKernel()
template<const int unrollFactor>
device__void
AddKernel_helper( int *out, const int *in, size_t N, int increment, int cycles )
{
for ( size_t i = unrollFactor*blockIdx.x*blockDim.x+threadIdx.x;
i < N;
i += unrollFactor*blockDim.x*gridDim.x )
{
int values[unrollFactor];
for ( int iUnroll = 0; iUnroll < unrollFactor; iUnroll++) {
size_t index = i+ iUnroll*blockDim.x;
values[iUnroll] = in[index];
}
for ( int iUnroll = 0; iUnroll < unrollFactor; iUnroll++) {
for ( int k = 0; k < cycles; k++) {
values[iUnroll] += increment;
}
}
for ( int iUnroll - 0; iUnroll < unrollFactor; iUnroll++) {
size_t index = i+ iUnroll*blockDim.x;
out[index] = values[iUnroll];
}
}
device__void
AddKernel( int *out, const int *in, size_t N, int increment, int cycles, int unrollFactor )
{
switch ( unrollFactor ) {
case 1: return AddKernel_helper<1>(out, in, N, increment, cycles);
case 2: return AddKernel_helper<2>(out, in, N, increment, cycles);
case 4: return AddKernel_helper<4>(out, in, N, increment, cycles);
}
}对于unrollFactor ,其实现与代码清单6-3中相同。对于unrollFactor ,映射锁页方案比流方案显示出一些改进。转折点是从cycles=64下降到cycles=48。对于unrollFactor ,性能同样优于流版本。
这些值是针对32M的整数给出的,所以程序需要读取和写入128MB的数据。对于cycles==48,程序运行在26毫秒内。为了达到此等有效带宽(在PCIe 2.0上超过9GB/s),GPU在执行内核处理的同时,通过PCIe进行并发读和写。
[1] 除非N不可被unrollFactor整除。这当然可以对for循环做一微小改变,事后做一清理来进行修正。