CUDA C编程权威指南 第3章 CUDA执行模型

一、CUDA执行模型概述

CUDA编程模型中两个主要的抽象概念：内存层次结构和线程层次结构。它们能够控制大规模并行GPU。

1. GPU架构概述

GPU架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。

GPU中每一个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，所以在一个GPU上并发执行数千个线程是可能的。当启动一个内核网格时，它的线程块被分布在了可用的SM上来执行。多个线程块可能会被分配到同一个SM上，而且是根据SM资源的可用性进行调度的。

CUDA采用单指令多线程（SIMT）架构来管理和执行线程，每32个线程为一组，被称为线程束（warp）。线程束中的所有线程同时执行相同的指令。每个线程都有自己的指令地址计数器和寄存器状态，利用自身的数据执行当前的指令。每个SM都将分配给它的线程块划分到包含32个线程的线程束，然后在可用的硬件资源上调度执行。

SIMT与SIMD的区别：两者都是将相同的指令广播给多个执行单元来实现并行。一个关键的区别是SIMD要求同一个向量中的所有元素要在一个统一的同步组中一起执行，而SIMT允许属于同一线程束的多个线程独立执行。尽管一个线程束中的所有线程在相同的程序地址上同时开始执行，但是单独的线程仍有可能有不同的行为。

一个线程块只能在一个SM上被调度。一旦线程块在一个SM上被调度，就会保存在该SM上直到执行完成。在同一时间，一个SM可以容纳多个线程块。

共享内存被分配在SM上的常驻线程块中，寄存器在线程中被分配。

在并发的线程束间切换并没有开销，因为硬件资源已经被分配到了SM上的所有线程和块中，所以最新被调度的线程束的状态已经存储在SM上。

2. Fermi架构

Fermi是第一个完整的GPU计算架构。

Fermi的特征是多达512个加速核心，这被称为CUDA核心。每个CUDA核心都有一个全流水线的整数算术逻辑单元（ALU）和一个浮点运算单元（FPU），在这里每个时钟周期执行一个整数或是浮点数指令。CUDA核心被组织到16个SM中，每一个SM含有32个CUDA核心。Fermi架构有6个384位的GDDR5 DRAM存储器接口，支持多达6GB的全局机载内存，这是许多应用程序关键的计算资源。主机接口通过PCIe总线将GPU与CPU相连。

GigaThread引擎是一个全局调度器，用来分配线程块到SM线程束调度器上。

Fermi架构包含一个耦合的768KB的二级缓存，被16个SM所共享。在图中，一个垂直矩形条表示一个SM，包含了以下内容：

· 执行单元（CUDA核心）

· 调度线程束的调度器和调度单元

· 共享内存、寄存器文件和一级缓存

每一个多处理器有16个加载/存储单元，允许每个时钟周期内有16个线程（线程束的一半）计算源地址和目的地址。特殊功能单元（SFU）执行固有指令，如正弦、余弦、平方根和插值。每个SFU在每个时钟周期内的每个线程上执行一个固有指令。

每个SM有两个线程束调度器和两个指令调度单元。当一个线程块被指定给一个SM时，线程块中的所有线程就分成了线程束。两个线程束调度器选择两个线程束，再把一个指令从线程束中发送到一个组上，组里有16个CUDA核心，16个加载/存储单元或四个特殊功能单元。Fermi架构，计算性能2.x，可以在每个SM上同时处理48个线程束，即可在一个SM上同时常驻1536个线程。

Fermi架构的一个关键特征是有一个64KB的片内可配置存储器，它在共享内存与一级缓存之间进行分配。共享内存允许一个块上的线程相互合作，这有利于芯片内数据的广泛重用，并大大降低了片外的通信量。CUDA提供了一个运行时API，它可以用来调整共享内存和一级缓存的数量。根据给定的内核中共享内存或缓存的使用修改片内存储器的配置，可以提高性能。

Fermi架构也支持并发内核执行：在相同的GPU上执行相同应用程序的上下文中，同时启动多个内核。Fermi架构允许多达16个内核同时在设备上运行。

3. Kepler架构

Kepler架构的3个重要创新：

· 强化的SM

· 动态并行

· Hyper-Q技术

下图是Kepler K20X芯片框图，它包含了15个SM和6个64位的内存控制器。

Kepler K20X的关键部分是有一个新的SM单元，其中包括一些结构的创新，以提高编程效率和功率效率。每个Kepler SM单元包含192个单精度CUDA核心，64个双精度单元，32个特殊功能单元（SFU）以及32个加载/存储单元（LD/ST）。

每个Kepler SM包括4个线程束调度器和8个指令调度器，以确保在单一的SM上同时发送和执行4个线程束。Kepler K20X架构（计算能力3.5）可以同时在每个SM上调度64个线程束，即在一个SM上可同时常驻2048个线程。K20X架构中寄存器文件容量达到64KB，Fermi架构中只有32KB。同时，K20X还允许片内存储器在共享内存和一级缓存间有更多的分区。K20X能够提供超过1TFlops的峰值双精度计算能力，相较于Fermi的设计，功率效率提高了80%。每瓦的性能也提升了3倍。

动态并行是Kepler GPU的一个新特性，它允许GPU动态启动新的网格。有了这个特点，任一内核都能启动其他的内核，并且管理任何核间需要的依赖关系来正确地执行附加的工作。这一特点也让你更容易创建和优化递归及与数据相关的执行模式。

有了动态并行，GPU能够启动嵌套内核，消除了与CPU通信的需求。

Hyper-Q技术增加了更多的CPU和GPU之间的同步硬件连接，以确保CPU核心能够在GPU上同时运行更多的任务。Fermi GPU依赖一个单一的硬件工作队列来从CPU到GPU间传送任务。这可能会导致一个单独的任务阻塞队列中在该任务之后的所有其他任务。Kepler GPU在主机与GPU之间提供了32个硬件工作队列。

4. 配置文件驱动优化

CUDA提供了两个主要的性能分析工具：nvvp，独立的可视化分析器；nvprof，命令行分析器。

nvvp是可视化分析器，它可以可视化并优化CUDA程序的性能。这个工具会显示CPU与GPU上的程序活动的时间表，从而找到可以改善性能的机会。此外，nvvp可以分析应用程序潜在的性能瓶颈，并给出建议以消除或减少这些瓶颈。该工具既可以作为一个独立的应用程序，也可以作为Nsight Eclipse Edition（nsight）的一部分。

nvprof在命令行上收集和显示分析数据。跟nvvp一样，它可以获得CPU与GPU上CUDA关联活动的时间表，其中包括内核执行、内存传输和CUDA的API调用。它也可以获得硬件计数器和CUDA内核的性能指标。

事件和指标

在CUDA性能分析中，事件是可计算的活动，它对应一个在内核执行期间被收集的硬件计数器。指标是内核的特征，它由一个或多个事件计算得到。

· 大多数计数器通过流式多处理器来报告，而不是通过整个GPU。

· 一个单一的运行只能获得几个计数器。有些计数器的获得是相互排斥的。多个性能分析运行往往需要获取所有相关的计数器。

· 由于GPU执行中的变化（如线程块和线程束调度指令），经重复运行，计数器值可能不是完全相同的。

在本书中主要使用nvprof来提高内核性能。

有3种常见的限制内核性能的因素：

· 存储带宽

· 计算资源

· 指令和内存延迟

如果想改善内核的性能，必须对硬件资源有一定的了解。

二、理解线程束执行的本质

1. 线程束和线程块

线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程就会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。

从硬件的角度来看，所有的线程都被组织成了一维的，线程块可以被配置为一维、二维或三维的。用x维度作为最内层的维度，y维度作为第二个维度，z作为最外层的维度，则二维或三维线程块的逻辑布局可以转化为一维物理布局。例如，对于一个给定的二维线程块，在一个块中每个线程的独特标识符都可以用内置变量threadIdx和blockDim来计算：

        threadIdx.y * blockDim.x + threadIdx.x 

对于一个三维线程块，计算如下：

        threadIdx.z * blockDim.y *  blockDim.x + threadIdx.y * blockDim.x + threadIdx.x

线程束不会在不同的线程块之间分离。如果线程块的大小不是线程束的偶数倍，那么在最后的线程束里有些线程就不会活跃。

2. 线程束分化

一个线程束中的所有线程在同一周期中必须执行相同的指令。

如果一个线程束中的线程产生分化，线程束将连续执行每一个分支路径，而禁用不执行这一路径的线程。为了获得最佳的性能，应该避免在同一线程束中有不同的执行路径。

核函数在设备上第一次运行时可能会增加间接开销。

3. 资源分配

线程束的本地执行上下文主要由以下资源组成：

· 线程计数器

· 寄存器

· 共享内存

每个SM都有32位的寄存器组，它存储在寄存器文件中，并且可以在线程中进行分配，同时固定数量的共享内存用来在线程块中进行分配。对于一个给定的内核，同时存在于同一个SM中的线程块和线程束的数量取决于在SM中可用的且内核所需的寄存器和共享内存的数量。

当计算资源（如寄存器和共享内存）已分配给线程块时，线程块被称为活跃的块。它所包含的线程束被称为活跃的线程束。活跃的线程束可以进一步被分为以下三种类型：

· 选定的线程束

· 阻塞的线程束

· 符合条件的线程束

一个SM上的线程束调度器在每个周期都选择活跃的线程束，然后把它们调度到执行单元。活跃执行的线程束被称为选定的线程束。如果一个活跃的线程束准备执行但尚未执行，它是一个符合条件的线程束。如果一个线程束没有做好执行的准备，它是一个阻塞的线程束。如果同时满足以下两个条件则线程束符合执行条件。

· 32个CUDA核心可用于执行

· 当前指令中所有的参数已就绪

4. 延迟隐藏

在指令发出和完成之间的时钟周期被定义为指令延迟。当每个时钟周期中所有的线程调度器都有一个符合条件的线程束时，可以达到计算资源的完全利用。

吞吐量和带宽

5. 占用率

使用maxThreadsPerMultiProcessor除以32，可以得到最大线程束数量。

6. 同步

栅栏同步是一个原语，它在许多并行编程语言中都很常见，在CUDA中，同步可以在两个级别执行：

线程块中的线程可以通过共享内存和寄存器来共享数据。

在不同的块间没有线程同步。块间同步，唯一安全的方法是在每个内核执行结束端使用全局同步点；也就是说，在全局同步之后，终止当前的核函数，开始执行新的核函数。

不同块中的线程不允许相互同步，因此GPU可以以任意顺序执行块。这使得CUDA程序在大规模并行GPU上是可扩展的。

7. 可扩展性

能够在可变数量的计算核心上执行相同的应用程序代码的能力被称为透明可扩展性。

CUDA内核启动时，线程块分布在多个SM中。网格中的线程块以并行或连续或任意的顺序被执行。这种独立性使得CUDA程序在任意数量的计算核心间可以扩展。

三、并行性的表现

一个内核的可实现占用率（achieved_occupancy）被定义为：每周期内活跃线程束的平均数量与一个SM支持的线程束最大数量的比值。

全局加载效率（gld_efficiency）：被请求的全局加载吞吐量占所需的全局加载吞吐量的比值。它衡量了应用程序的加载操作利用设备内存带宽的程度。

增大并行性

一个块的最内层维数（block.x）应该是线程束大小的倍数。

最后一次的执行配置块的大小为（256， 8），这是无效的。一个块中线程总数超过1024个（这是GPU的硬件限制）。

没有一个单独的指标能直接优化性能。我们需要在几个相关的指标间寻找一个恰当的平衡来达到最佳的总体性能。

四、避免分支分化

在向量中执行满足交换律和结合律的运算，被称为归约问题。并行归约问题是这种运算的并行执行。（加法、求最大值、最小值、平均值和乘积等都可以）

以并行归约为例，介绍避免分支分化的基本技术。

1. 并行归约问题

假如要对一个有N个元素的整数数组求和，使用如下串行代码很容易实现算法：

int sum = 0;
for (int i = 0; i < N; i++)
    sum += array[i];

大量数据元素可以用以下方法执行并行加法运算：

1. 将输入向量划分到更小的数据块中。

2. 用一个线程计算一个数据块的部分和。

3.对每个数据块的部分和再求和得出最终结果。

并行加法的一个常用方法是使用迭代成对实现。一个数据块只包含一对元素，并且一个线程对这两个元素求和产生一个局部结果。然后，这些局部结果在最初的输入向量中就地保存。这些新值被作为下一次迭代求和的输入值。因为输入值的数量会在每一次迭代后会减半，当输出向量的长度达到1时，最终的和就已经被计算出来了。

根据每次迭代后输出元素就地存储的位置，成对的并行和实现可以被进一步分为以下两种类型：

· 相邻配对：元素与它们直接相邻的元素配对

· 交错配对：根据给定的跨度配对元素

2. 并行归约中的分化

相邻配对

__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n)
{

    // set thread ID
    unsigned int tid =  threadIdx.x;
    
    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // boundary check
    if (idx >= n) return;

    // in-place reduction in global memory
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        if ((tid % (2 * stride)) == 0)
            idata[tid] += idata[tid + stride];
        //  synchronize within block
        __syncthreads();
    }
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

因为线程块间无法同步，所以每个线程块产生的部分和被复制回了主机，并且在哪儿进行串行求和。

3. 改善并行归约的分化

条件表达式 

if ((tid % (2 * stride)) == 0)

只对偶数ID的线程为true，所以这会导致很高的线程束分化。在每次迭代中，所有的线程都必须得调度。通过重新组织每个线程的数组索引来强制ID相邻的线程执行求和操作，线程束分化就能被归约了。

__global__ void reduceNeighboredLess(int *g_idata, int *g_odata, unsigned int n) {
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // boundary check
    if (idx >= n) return;

    // in-place reduction in global memory
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        // convert tid into local array index
        int index = 2 * stride * tid;
        if (index < blockDim.x) {
            idata[index] += idata[index+ stride];    
        }
        // synchronize within threadblock
        __syncthreads();
    }
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

4. 交错配对的归约

// Interleaved Pair Implementation with less divergence
__global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) {
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;
    
    // boundary check
    if (idx >= n) return;

    // in-place reduction in global memory
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            idata[tid] += idata[tid + stride];
        }
        __syncthreads();
    }
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

五、展开循环

循环展开时一个尝试通过减少分支出现的频率和循环维护指令来优化循环的技术。循环体的复制数量被称为循环展开因子，迭代次数就变为了原始循环迭代次数除以循环展开因子。

1. 展开的归约

__global__ void reduceUnrolling2 (int *g_idata, int *g_odata, unsigned int n) {
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x * 2;
    
    // unrolling 2 data blocks
    if (idx + blockDim.x < n)
        idata[idx] += idata[idx + blockDim.x];
    __syncthreads();
    
    // in-place reduction in global memory
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride)
            idata[tid] += idata[tid + stride];
        // synchronize within threadblock
        __syncthreads();
    }
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

__syncthreads是用于块内同步的。在归约核函数中，它用来确保在线程进入下一轮之前，每一轮中所有线程已经将局部结果写入全局内存中了。

2. 展开线程的归约

因为线程束的执行是SIMT的，每条指令之后都有隐式地线程束内同步过程。

考虑一下只剩下32个或更少线程（即一个线程束）的情况，归约循环的最后6个迭代可以用下述语句来展开：

__global__ void reduceUnrollWarp8 (int *g_idata, int *g_odata, unsigned int n) {
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x * 8;

    // unrolling 8
    if (idx + 7 * blockDim.x < n) {
        int a1 = g_idata[idx];
        int a2 = g_idata[idx + blockDim.x];
        int a3 = g_idata[idx + blockDim.x * 2];
        int a4 = g_idata[idx + blockDim.x * 3];
        int b1 = g_idata[idx + blockDim.x * 4];
        int b2 = g_idata[idx + blockDim.x * 5];
        int b3 = g_idata[idx + blockDim.x * 6];
        int b4 = g_idata[idx + blockDim.x * 7];
        g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
    }    
    __syncthreads();

    // in-place reduction in global memory
    for (int stride = blockDim.x / 2; stride > 32; stride >>= 1) {
        if (tid < stride)
            idata[tid] += idata[tid + stride];
        // synchronize within threadblock
        __syncthreads();
    }
    // unrolling warp
    if (tid < 32) {
        volatile int *vmem = idata;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1]; 
    }
    // write result for this block to global mem
    if (tid == 0)
        g_odata[blockIdx.x] = idata[0];
}

4. 模板函数的归约

使用模板函数有助于进一步减少分支消耗。在设备函数上CUDA支持模板参数。

template<unsigned int iBlockSize>
__global__ void reduceCompleteUnroll (int *g_idata, int *g_odata, unsigned int n) {
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x * 8;

    // unrolling 8
    if (idx + 7 * blockDim.x < n) {
        int a1 = g_idata[idx];
        int a2 = g_idata[idx + blockDim.x];
        int a3 = g_idata[idx + blockDim.x * 2];
        int a4 = g_idata[idx + blockDim.x * 3];
        int b1 = g_idata[idx + blockDim.x * 4];
        int b2 = g_idata[idx + blockDim.x * 5];
        int b3 = g_idata[idx + blockDim.x * 6];
        int b4 = g_idata[idx + blockDim.x * 7];
        g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
    }    
    __syncthreads();

    // in-place reduction and complete
    if (iBlockSize >= 1024 && tid < 512) idata[tid] += idata[tid + 512];
    __syncthreads();
    
    if (iBlockSize >= 512 && tid < 256) idata[tid] += idata[tid + 256];
    __syncthreads();

    if (iBlockSize >= 256 && tid < 128) idata[tid] += idata[tid + 128];
    __syncthreads();

    if (iBlockSize >= 128 && tid < 64) idata[tid] += idata[tid + 64];
    __syncthreads();

    // unrolling warp
    if (tid < 32) {
        volatile int *vmem = idata;
        vmem[tid] += vmem[tid + 32];
        vmem[tid] += vmem[tid + 16];
        vmem[tid] += vmem[tid + 8];
        vmem[tid] += vmem[tid + 4];
        vmem[tid] += vmem[tid + 2];
        vmem[tid] += vmem[tid + 1]; 
    }
    // write result for this block to global mem
    if (tid == 0)
        g_odata[blockIdx.x] = idata[0];
}

六、动态并行

CUDA的动态并行允许在GPU端直接创建和同步新的GPU内核。

1. 嵌套执行

在动态并行中，内核执行分为两种类型：父母和孩子。子网格的调用和完成必须进行适当地嵌套，这意味着在线程创建的所有子网格都完成之后，父网格才会完成。如果调用的线程没有显式地同步启动子网格，那么运行时保证父母和孩子之间的隐式同步。下图中，在父线程中设置了栅栏，从而可以与其子网格显式地同步。

如果块中所有线程在所有的子网格完成之前退出，那么在那些子网格上隐式同步会被触发。

当父母启动一个子网格，父线程块与孩子显式同步之后，孩子才能开始执行。

父网格和子网格共享相同的全局和常量内存存储，但它们有不同的局部和共享内存。当父线程优于子网格调用时，所有的全局内存操作要保证对子网格是可见的。当父母在子网格完成时进行同步操作后，子网格所有的内存操作应保证对父母是可见的。

共享内存和局部内存分别对于线程块或线程来说是私有的。

2. 在GPU上嵌套Hello World

__global__ void nestedHelloWorld(const int iSize, int iDepth) {
    int tid = threadIdx.x;
    printf("Recursion = %d: Hello World from thread %d block %d\n", iDepth, tid, blockIdx.x);

    // condition to stop recursive execution
    if (iSize == 1) return;

    // reduce block size to half
    int nthreads = iSize >> 1;

    // thread 0 launches child grid recursively
    if (tid == 0 && nthreads > 0) {
        nestedHelloWorld<<<1, nthreads>>>(nthreads, ++iDepth);
        printf("------> nested execution depth: %d\n", iDepth);
    } 

}

使用两个线程块调用父网格：

注意：子网格线程块ID都是0，因为父网格嵌套的子网格中只包含一个线程块。

动态并行的限制条件

动态并行只有在计算能力为3.5或更高的设备上才能被支持。

通过动态并行调用的内核不能在物理方面独立的设备上启动。

动态并行的最大嵌套深度限制为24，但实际上，在每一个新的级别中大多数内核受限于设备运行时系统需要的内存数量。因为为了对每个嵌套层中的父网格和子网格之间进行同步管理，设备运行时要保留额外的内存。

3. 嵌套归约

在一个子网格被调用后，它看到的内存与父线程是完全一样的。因为每一个子线程只需要父线程的数值来指导部分归约，所以在每个子网格启动前执行线程块内部的同步是没有必要的。去除所有同步操作会产生如下的核函数。

__global__ void gpuRecursiveReduceNosync(int *g_idata, int *g_odata, unsigned int isize) {
    // set thread ID
    unsigned int tid = threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;
    int *odata = &g_odata[blockIdx.x];

    // stop condition
    if (iSize == 2 && tid == 0) {
        g_odata[blockIdx.x] = idata[0] + idata[1];
        return;
    }

    // nested invoke
    int istride = iSize >> 1;
    if (istride > 1 && tid < istride) {
        idata[tid] += idata[tid + istride];
        if (tid == 0)
        gpuRecursiveReduceNosync<<<1, istride>>>(idata, odata, istride);
    }
}

大量的内核调用造成内核效率很低。

__global__ void gpuRecursiveReduce2(int *g_idata, int *g_odata, int iStride, int const iDim) {
    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * iDim;
    
   // stop condition
    if (iStride == 1 && threadIdx.x == 0) {
        g_odata[blockIdx.x] = idata[0] + idata[1];
        return;
    }

    // in place reduction
    idata[threadIdx.x] += idata[threadIdx.x + iStride];
    
    // nested invocation to generate child grids
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        gpuRecursiveReduce2 <<<gridDim.x, iStride / 2>>> (g_idata, g_odata, iStride / 2, iDim);
    }

}

gpuRecursiveReduce2内核中创建的子网格数量大大减少，提升了性能。