CUDA Dynamic Parallelism测试

CUDA 动态并行（CUDA Dynamic Parallelism）是 NVIDIA 在其 CUDA 编程模型中引入的一个强大特性。它允许 GPU 上运行的内核（kernel）直接在设备端启动新的内核，而无需返回主机（CPU）进行控制。这一特性使得我们可以在 GPU 上实现更复杂、更动态的算法，提高程序的并行度和执行效率。

一.知识点

1. CUDA 动态并行性（Dynamic Parallelism）：这是对 CUDA 编程模型的扩展，允许 CUDA 内核直接在 GPU 上创建和同步新工作。它使得在程序的任何需要的地方动态创建并行性成为可能。
2. 减少主机和设备之间的数据传输：通过在设备上运行的线程在运行时决定启动配置，动态并行性可以减少在主机和设备之间传输执行控制和数据的需要。
3. 数据驱动的并行工作生成：在运行时，内核可以根据数据驱动的决策或工作负载，在内核内生成依赖于数据的并行工作，动态利用 GPU 的硬件调度器和负载均衡器。
4. 表达复杂的算法和编程模式：以前需要修改以消除递归、不规则循环结构或其他不适合单级并行性的算法，现在可以更透明地表达。
5. 支持的计算能力：动态并行性仅支持计算能力为 3.5 及以上的设备。
6. CUDA 执行模型的扩展：支持动态并行性的 CUDA 执行模型现在允许设备线程配置、启动新网格（grids），并在设备上对其进行隐式同步。
7. 父子网格的关系：
   • 父线程、线程块、网格：启动新网格的实体，被称为父级。
   • 子网格：由父级启动的新网格。
   • 嵌套执行：子网格的启动和完成是正确嵌套的，父网格在其所有子网格完成之前不会被视为完成。
8. 设备运行时的作用：提供使内核函数能够使用动态并行性的运行时系统和 API。
9. 网格范围内的资源共享：在设备上，所有线程在网格内共享已启动的内核和 CUDA 对象。这意味着一个线程创建的流可以被网格内的任何其他线程使用。
10. 流和事件的使用：
    • 设备上创建的流：仅在创建它们的网格范围内存在，超出该范围的行为是未定义的。
    • NULL 流的特殊性：在设备上，隐式的 NULL 流只在线程块内共享，不同线程块中的线程对 NULL 流的启动可能会并发执行。
11. 并发执行的保证：CUDA 编程模型不保证设备上不同线程块之间的并发执行，包括父网格和子网格之间。
12. 多 GPU 支持的限制：设备运行时不支持多 GPU；它只能在当前执行的设备上操作。
13. 内存一致性和可见性：
    • 全局和常量内存：父子网格共享相同的全局和常量内存，但具有独立的本地和共享内存。
    • 内存一致性保证：子网格只有在启动时保证与父线程的内存视图是一致的。由于 cudaDeviceSynchronize() 的移除，父网格无法在退出前保证看到子网格的内存修改。
14. 零拷贝内存：与全局内存具有相同的一致性保证，但内核不能在设备上分配或释放零拷贝内存。
15. 设备运行时 API：
    • 类似于主机运行时 API：设备运行时的语法和语义与主机运行时 API 基本相同，便于代码重用。
    • 内核启动的异步性：与主机端启动相同，设备端内核启动相对于启动线程是异步的。
    • 不支持的功能：设备运行时不支持像 cudaStreamSynchronize() 和 cudaStreamQuery() 这样的 API，也不支持从设备上创建或销毁纹理和表面对象。
16. 特殊流的使用：
    • Fire-and-Forget 流（cudaStreamFireAndForget）：用于立即调度启动，无需依赖之前的启动，无法与事件配合使用。
    • Tail Launch 流（cudaStreamTailLaunch）：用于在父网格完成后调度新的网格启动，同样无法与事件配合使用。
17. 事件的限制：仅支持用于流间同步的 CUDA 事件，不支持 cudaEventSynchronize()、cudaEventElapsedTime() 和 cudaEventQuery() 等功能。
18. 设备属性查询的限制：只能查询当前设备的属性，不支持在设备运行时切换设备。
19. 全局和常量内存变量的行为：设备上的所有内核都可以读取或写入全局变量，但不能修改常量内存中的数据。
20. 错误处理：
    • 错误代码的获取：每个线程可以通过 cudaGetLastError() 获取其生成的最后一个错误代码。
    • 错误传播：子网格中的错误（例如访问无效地址）将返回到主机。
21. PTX 支持：CUDA 提供了底层的 PTX API，如 cudaLaunchDevice() 和 cudaGetParameterBuffer()，供需要在 PTX 级别支持动态并行性的编程语言和编译器实现者使用。
22. 编译和链接：
    • 不需要显式包含头文件：在编译 CUDA 程序时，会自动包含设备运行时 API 的原型。
    • 设备运行时库：使用动态并行性的 CUDA 程序需要链接设备运行时静态库 libcudadevrt。
23. 系统资源的限制和配置：
    • 启动池的大小：受限于系统资源，可以使用 cudaDeviceSetLimit() 配置启动池的大小。
    • 堆栈大小的控制：可以通过 cudaDeviceSetLimit() 设置每个 GPU 线程的堆栈大小。
24. 内存分配注意事项：
    • 设备上的 cudaMalloc() 和 cudaFree()：在设备上调用时，与主机上的行为不同，映射到设备端的 malloc() 和 free()，受限于设备 malloc 堆大小。
    • 指针的限制：在设备上分配的内存指针不能在主机上释放，反之亦然。
25. 线程重调度的注意事项：设备运行时可能会将线程块重新调度到不同的 SM，以更有效地管理资源，因此依赖 %smid 或 %warpid 保持不变是不安全的。
26. ECC 错误处理：CUDA 内核中无法通知 ECC 错误，所有 ECC 错误将在整个启动树完成后在主机端报告。

二.测试内容

1.查看动态并行生成的PTX

2.性能对比测试(测试一个向量的N次累加)

• A.调用N次Kernel
• B.在Kernel里循环N次
• C.使用动态并行,递归N次

3.动态并行是如何调度SM的

三.查看动态并行生成的PTX

tee dynamic_parallelism.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>

#define CHECK_CUDA(call)                      \
  do {
                                   \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {
                      \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)


__global__ void kernel(float *iodata,int count)
{
   
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(count>0)
  {
   
    iodata[tid]+=1;
    if(tid == 0)
    {
   
      __prof_trigger(0);
      kernel<<<gridDim.x, blockDim.x,0,cudaStreamFireAndForget >>>(iodata, count - 1);
      __prof_trigger(1);
    }
  }
}

int main(int argc,char *argv[])
{
   
  int deviceid=0;cudaSetDevice(deviceid); 
  int block_count=100000;
  int block_size=1024;
  int count=1000;
    
  size_t value;
  CHECK_CUDA(cudaDeviceGetLimit(&value, cudaLimitDevRuntimePendingLaunchCount));
  printf("cudaLimitDevRuntimePendingLaunchCount:%ld\n",value);

  {
   
      int thread_size=block_count*block_size;
      float *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,thread_size*sizeof(float),cudaHostAllocDefault));  
      for(int i=0;i<thread_size;i++) iodata[i]=0;
      kernel<<<block_count, block_size>>>(iodata,count);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }