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 $= = 1$ ，其实现与代码清单6-3中相同。对于unrollFactor $= = 2$ ，映射锁页方案比流方案显示出一些改进。转折点是从cycles=64下降到cycles=48。对于unrollFactor $= = 4$ ，性能同样优于流版本。

这些值是针对32M的整数给出的，所以程序需要读取和写入128MB的数据。对于cycles==48，程序运行在26毫秒内。为了达到此等有效带宽（在PCIe 2.0上超过9GB/s），GPU在执行内核处理的同时，通过PCIe进行并发读和写。

[1] 除非N不可被unrollFactor整除。这当然可以对for循环做一微小改变，事后做一清理来进行修正。