当前位置： 首页 > article >正文

CUDA Dynamic Parallelism测试

article 2025/2/24 0:22:26

CUDA Dynamic Parallelism测试

  • 一.知识点
  • 二.测试内容
    • 1.查看动态并行生成的PTX
    • 2.性能对比测试(测试一个向量的N次累加)
    • 3.动态并行是如何调度SM的
  • 三.查看动态并行生成的PTX
  • 四.性能对比测试
  • 五.动态并行是如何调度SM的

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));
  }  
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -ptx -arch=sm_86 -rdc=true  \
    -o dynamic_parallelism.ptx dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
cat dynamic_parallelism.ptx
  • 输出
.version 8.1
.target sm_86
.address_size 64

        // .globl       _Z6kernelPfi
.extern .func  (.param .b64 func_retval0) __cudaCDP2GetParameterBufferV2
(
        .param .b64 __cudaCDP2GetParameterBufferV2_param_0,
        .param .align 4 .b8 __cudaCDP2GetParameterBufferV2_param_1[12],
        .param .align 4 .b8 __cudaCDP2GetParameterBufferV2_param_2[12],
        .param .b32 __cudaCDP2GetParameterBufferV2_param_3
);
.extern .func  (.param .b32 func_retval0) __cudaCDP2LaunchDeviceV2
(
        .param .b64 __cudaCDP2LaunchDeviceV2_param_0,
        .param .b64 __cudaCDP2LaunchDeviceV2_param_1
);

.visible .entry _Z6kernelPfi(
        .param .u64 _Z6kernelPfi_param_0,
        .param .u32 _Z6kernelPfi_param_1
)
{
  .reg .pred      %p<4>;
  .reg .f32       %f<3>;
  .reg .b32       %r<11>;
  .reg .b64       %rd<8>;
  
  ld.param.u64    %rd2, [_Z6kernelPfi_param_0];
  ld.param.u32    %r2, [_Z6kernelPfi_param_1];
  mov.u32         %r1, %ntid.x;
  setp.lt.s32     %p1, %r2, 1;
  @%p1 bra        $L__BB0_5;

  mov.u32         %r3, %ctaid.x;
  mov.u32         %r4, %tid.x;
  mad.lo.s32      %r5, %r3, %r1, %r4;
  cvta.to.global.u64      %rd3, %rd2;
  mul.wide.u32    %rd4, %r5, 4;
  add.s64         %rd5, %rd3, %rd4;
  ld.global.f32   %f1, [%rd5];
  add.f32         %f2, %f1, 0f3F800000;
  st.global.f32   [%rd5], %f2;
  setp.ne.s32     %p2, %r5, 0;
  @%p2 bra        $L__BB0_5;

  // begin inline asm
  pmevent         0;
  // end inline asm
  mov.u32         %r6, %nctaid.x;
  mov.u32         %r7, 1;
  mov.u64         %rd6, _Z6kernelPfi;
  mov.u32         %r8, 0;
  { // callseq 0, 0
    .reg .b32 temp_param_reg;
    .param .b64 param0;
    st.param.b64    [param0+0], %rd6;
    .param .align 4 .b8 param1[12];
    st.param.b32    [param1+0], %r6;
    st.param.b32    [param1+4], %r7;
    st.param.b32    [param1+8], %r7;
    .param .align 4 .b8 param2[12];
    st.param.b32    [param2+0], %r1;
    st.param.b32    [param2+4], %r7;
    st.param.b32    [param2+8], %r7;
    .param .b32 param3;
    st.param.b32    [param3+0], %r8;
    .param .b64 retval0;
    call.uni (retval0), __cudaCDP2GetParameterBufferV2,(param0,param1,param2,param3);
    ld.param.b64    %rd1, [retval0+0];
  } // callseq 0
  setp.eq.s64     %p3, %rd1, 0;
  @%p3 bra        $L__BB0_4;
  add.s32         %r9, %r2, -1;
  st.u64  [%rd1], %rd2;
  st.u32  [%rd1+8], %r9;
  mov.u64         %rd7, 4;
  { // callseq 1, 0
    .reg .b32 temp_param_reg;
    .param .b64 param0;
    st.param.b64    [param0+0], %rd1;
    .param .b64 param1;
    st.param.b64    [param1+0], %rd7;
    .param .b32 retval0;
    call.uni (retval0),__cudaCDP2LaunchDeviceV2,(param0,param1); 
    ld.param.b32    %r10, [retval0+0];
  } // callseq 1
$L__BB0_4:
  // begin inline asm
  pmevent         1;
  // end inline asm
$L__BB0_5:
  ret;
}

四.性能对比测试

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>
#include <chrono>
#include <thread>

#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 case_0(float *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  iodata[tid]+=1;
}

__global__ void case_1(float *iodata,int count=100)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  for(int i=0;i<count;i++)
  {
      iodata[tid]+=1;
  }
}

__global__ void case_2(float *iodata,int count,bool is_block)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(count>0)
  {
    iodata[tid]+=1;
    if(is_block) __syncthreads(); // 同步所有线程
    if(tid == 0)
      case_2<<<gridDim.x, blockDim.x,0>>>(iodata, count - 1,is_block);
  }
}

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

template <typename F>
void TIMEIT(F const &f,cudaStream_t &stream,cudaEvent_t &start_ev,cudaEvent_t&stop_ev)
{ 
    f(stream); 
    CHECK_CUDA(cudaDeviceSynchronize());
    auto start = std::chrono::high_resolution_clock::now();
    cudaEventRecord(start_ev, stream); 
    f(stream); 
    cudaEventRecord(stop_ev, stream); 
    CHECK_CUDA(cudaEventSynchronize(stop_ev)); 
    auto end = std::chrono::high_resolution_clock::now(); 
    std::chrono::duration<double> diff = end - start; 
    float milliseconds = 0; 
    cudaEventElapsedTime(&milliseconds, start_ev, stop_ev); 
    printf("E2E:%7.2fms Kernel:%7.2fms\n",diff.count()*1000,milliseconds);
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int block_count=100000;
  int block_size=1024;
  int count=1000;
  
  cudaStream_t stream;
  cudaStreamCreate(&stream);

  cudaEvent_t start_ev, stop_ev;
  cudaEventCreate(&start_ev);
  cudaEventCreate(&stop_ev);
  
  size_t value;
  CHECK_CUDA(cudaDeviceGetLimit(&value, cudaLimitDevRuntimePendingLaunchCount));
  printf("cudaLimitDevRuntimePendingLaunchCount:%ld\n",value);

  {//Host循环Lanuch count次
      printf(" ----------------- case 0 ----------------- \n");
      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;
      TIMEIT([&](cudaStream_t &stream)-> void {for(int i=0;i<count;i++){case_0<<<block_count, block_size,0,stream>>>(iodata);}},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel内循环加count次
      printf(" ----------------- case 1 ----------------- \n");
      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;
      TIMEIT([&](cudaStream_t &stream)-> void {case_1<<<block_count, block_size>>>(iodata,count);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel通过动态并行,递归count次,每次同步线程块
      printf(" ----------------- case 2 ----------------- \n");
      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;
      TIMEIT([&](cudaStream_t &stream)-> void {case_2<<<block_count, block_size>>>(iodata,count,true);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }
  {//Kernel通过动态并行,递归count次,每次不需要同步线程块
      printf(" ----------------- case 3 ----------------- \n");
      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;
      TIMEIT([&](cudaStream_t &stream)-> void {case_2<<<block_count, block_size>>>(iodata,count,false);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
  {//Kernel通过动态并行,递归count次,每次不需要同步线程块,cudaStreamFireAndForget
      printf(" ----------------- case 4 ----------------- \n");
      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;
      TIMEIT([&](cudaStream_t &stream)-> void {case_3<<<block_count, block_size>>>(iodata,count);},stream,start_ev,stop_ev);
      printf("%f %f\n",iodata[0],iodata[thread_size-1]);
      CHECK_CUDA(cudaFreeHost(iodata));
  }  
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -rdc=true -lineinfo \
    -o dynamic_parallelism dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./dynamic_parallelism
  • 输出
cudaLimitDevRuntimePendingLaunchCount:2048
 ----------------- case 0 -----------------
E2E:84655.40ms Kernel:84654.82ms   #Host循环Lanuch count次
2000.000000 2000.000000
 ----------------- case 1 -----------------
E2E:  85.26ms Kernel:  82.59ms     #Kernel内循环加count次
2000.000000 2000.000000
 ----------------- case 2 -----------------
E2E:85089.90ms Kernel:85088.64ms   #Kernel通过动态并行,递归count次,每次同步线程块
2000.000000 2000.000000
 ----------------- case 3 -----------------
E2E:84906.45ms Kernel:84904.73ms   #Kernel通过动态并行,递归count次,每次不需要同步线程块
2000.000000 2000.000000
 ----------------- case 4 -----------------
E2E:84757.05ms Kernel:84755.00ms   #Kernel通过动态并行,递归count次,每次不需要同步线程块,cudaStreamFireAndForget
2000.000000 2000.000000

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

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>
#include <chrono>
#include <thread>

#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 child_kernel(unsigned int *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int smid;
  asm volatile("mov.u32 %0, %smid;" : "=r"(smid));  
  iodata[tid]=smid;
}

__global__ void main_kernel(unsigned int *iodata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if(tid == 0) child_kernel<<<28,1024,0>>>(iodata);
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int count=28*1024;
  
  unsigned int last_smid=-1;
  {
      unsigned int *iodata;
      CHECK_CUDA(cudaHostAlloc(&iodata,count*sizeof(unsigned int),cudaHostAllocDefault));
      for(int i=0;i<count;i++) iodata[i]=0;
      main_kernel<<<1, 1>>>(iodata);
      CHECK_CUDA(cudaDeviceSynchronize());
      for(int i=0;i<count;i++) 
      {
        if(iodata[i]!=last_smid)
        {
          printf("tid:%06d smid:%04d\n",i,iodata[i]);
          last_smid=iodata[i];
        }
      }
      CHECK_CUDA(cudaFreeHost(iodata));
  }
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -rdc=true -lineinfo \
    -o dynamic_parallelism dynamic_parallelism.cu \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./dynamic_parallelism
  • 输出
tid:000000 smid:0000
tid:001024 smid:0002
tid:002048 smid:0004
tid:003072 smid:0006
tid:004096 smid:0008
tid:005120 smid:0010
tid:006144 smid:0012
tid:007168 smid:0014
tid:008192 smid:0016
tid:009216 smid:0018
tid:010240 smid:0020
tid:011264 smid:0022
tid:012288 smid:0024
tid:013312 smid:0026
tid:014336 smid:0001
tid:015360 smid:0003
tid:016384 smid:0005
tid:017408 smid:0007
tid:018432 smid:0009
tid:019456 smid:0011
tid:020480 smid:0013
tid:021504 smid:0015
tid:022528 smid:0017
tid:023552 smid:0019
tid:024576 smid:0021
tid:025600 smid:0023
tid:026624 smid:0025
tid:027648 smid:0027

http://www.kler.cn/a/327642.html

相关文章：

  • Django-form表单
  • 【JAVA高级】如何使用Redis加锁和解锁（二）做分布式锁案例（防误删锁）
  • https访问报错：net::ERR_CERT_DATE_INVALLD
  • 【Simulink仿真】两级式三相光伏并网发电系统
  • Linux启动mysql报错
  • 【QT Quick】C++交互：QML对象操作
  • Linux应急响应技巧整理
  • 如何确定光纤用几芯 用光纤与网线区别在哪里
  • TensorRT-LLM高级用法
  • 五子棋双人对战项目（1）——WebSocket介绍
  • ubuntu 开启root
  • 太速科技-FMCJ457-基于JESD204B的2路2Gsps AD 2路2Gsps DA FMC子卡
  • LED驱动电路
  • SQL | 查询带有单引号的异常数据和replace与insert插入的区别
  • YOLO11震撼发布！
  • 《Python编程:从入门到实践》数据可视化
  • Linux操作系统中Redis
  • 从源码中学习动态代理模式
  • SpringBoot自定义异常
  • Webpack 4 优化指南：提升构建性能与加载速度