CUDA学习之CUDA执行模型--part5

3.6 动态并行

到目前为止，我们所有的内核都是在主机线程中调用的，是否我们可以在内核中调用内核，这个内核可以是别的内核，也可以是自己，那么我们就需要动态并行了，这个功能在早期的设备上是不支持的。

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

动态并行的好处之一就是能让复杂的内核变得有层次，坏处就是写出来的程序更复杂，因为并行行为本来就不好控制.

有了动态并行，可以推迟到运行时决定需要在GPU上创建多少个块和网格，可以动态地利用GPU硬件调度器和加载平衡器,并进行调整以适应数据驱动或工作负载。

在GPU端直接创建工作的能力可以减少在主机和设备之间传输执行控制和数据的需求,因为在设备上执行的线程可以在运行时决定启动配置。

3.6.1 嵌套执行

内核中启动内核，和cpu并行中有一个相似的概念，就是父线程和子线程。子线程由父线程启动，但是到了GPU，这类名词相对多了些，比如父网格，父线程块，父线程，对应的子网格，子线程块，子线程。子网格被父线程启动，且必须在对应的父线程，父线程块，父网格结束之前结束。
所有的子网格结束后，父线程，父线程块，父网格才会结束。

线程可能与由该线程启动的或由相同线程块中其他线程启动的子网格同步。在线程块中，只有当所有线程创建的所有子网格完成之后，线程块的执行才会完成。如果块中所有线程在所有的子网格完成之前退出，那么在那些子网格上隐式同步会被触发。

父线程块启动子网格需要显式的同步，也就是说不同的线程束需要都执行到子网格调用那一句，这个线程块内的所有子网格才能依据所在线程束的执行，一次执行。

1. 父网格和子网格共享相同的全局和常量内存。
2. 父网格子网格有不同的局部内存
3. 有了子网格和父网格间的弱一致性作为保证，父网格和子网格可以对全局内存并发存取。
4. 有两个时刻父网格和子网格所见到的内存一致：子网格启动的时候，子网格结束的时候
5. 当父线程优于子网格调用时,所有的全局内存操作要保证对子网格是可见的。当父母在子网格完成时进行同步操作后,子网格所有的内存操作应保证对父母是可见的。
6. 共享内存和局部内存分别对于线程块和线程来说是私有的
7. 局部内存对线程私有，对外不可见。

3.6.2 在GPU上嵌套Hello World

#include <cuda_runtime.h>
#include <stdio.h>
__global__ void nesthelloworld(int iSize,int iDepth)
{
    unsigned int tid=threadIdx.x;
    printf("depth : %d blockIdx: %d,threadIdx: %d\n",iDepth,blockIdx.x,threadIdx.x);
    if (iSize==1)
        return;
    int nthread=(iSize>>1);
    if (tid==0 && nthread>0)
    {
        nesthelloworld<<<1,nthread>>>(nthread,++iDepth);
        printf("-----------> nested execution depth: %d\n",iDepth);
    }

}

int main(int argc,char* argv[])
{
    int size=64;
    int block_x=2;
    dim3 block(block_x,1);
    dim3 grid((size-1)/block.x+1,1);
    nesthelloworld<<<grid,block>>>(size,0);
    cudaGetLastError();
    cudaDeviceReset();
    return 0;
}

编译

nvcc -arch=sm_35  nested_Hello_World.cu -o nested_Hello_World -lcudadevrt --relocatable-device-code true

-lcudadevrt –relocatable-device-code true 是前面没有的，这两个指令是动态并行需要的一个库，
relocatable-device-code表示生成可重新定位的代码，第十章将会讲解更多重新定位设备代码的内容。

主机应用程序调用父网格，该父网格在一个线程块中有8个线程。然后，该父网格中的线程0调用一个子网格，该子网格中有一半线程,即4个线程。之后,第一个子网格中的线程0再调用一个新的子网格,这个新的子网格中也只有一半线程,即2个线程,以此类推,直到最后的嵌套中只剩下一个线程。

动态并行的限制条件

1. 动态并行只有在计算能力为3.5或更高的设备上才能被支持。
2. 通过动态并行调用的内核不能在物理方面独立的设备上启动。然而,在系统中允许查询任一个带CUDA功能的设备性能。
3. 动态并行的最大嵌套深度限制为24,但是实际上,在每一个新的级别中大多数内核受限于设备运行时系统需要的内存数量。因为为了对每个嵌套层中的父网格和子网格之间进行同步管理,设备运行时要保留额外的内存。

3.6.3 嵌套归约

嵌套归约并不能降低代码复杂度，其次，其运行效率也没有提高，动态并行，相当于串行编程的中的递归调用，递归调用如果能转换成迭代循环，一般为了效率的时候是要转换成循环的，只有当效率不是那么重要，而更注重代码的简洁性的时候，我们才会使用。

不看了，马德，脑瓜子嗡嗡的，以后再说