CUDA 内核是如何启动的？

Posted 2023-04-15

【中文标题】CUDA 内核是如何启动的？

【英文标题】：How is a CUDA kernel launched?

【发布时间】：2012-08-23 17:30:31

【问题描述】：

我创建了一个简单的 CUDA 应用程序来添加两个矩阵。它编译得很好。我想知道所有线程将如何启动内核以及 CUDA 内部的流程是什么？我的意思是，每个线程将以何种方式执行矩阵的每个元素。

我知道这是一个非常基本的概念，但我不知道。我对流程感到困惑。

【参考方案1】：

你启动一个方块网格。

块被不可分割地分配给多处理器（其中多处理器上的块数决定了可用共享内存的数量）。

块被进一步分成经线。对于有 32 个线程的 Fermi GPU，它们要么执行相同的指令，要么处于非活动状态（因为它们分支了，例如，通过比同一个 warp 中的邻居更早退出循环或不采用 if 他们所做的）。在 Fermi GPU 上，一次最多在一个多处理器上运行两个 warp。

只要有延迟（即内存访问或数据依赖完成的执行停止），就会运行另一个warp（适合一个多处理器的warp数量 - 相同或不同的块 - 由寄存器的数量决定每个线程使用的共享内存量以及一个/块使用的共享内存量。

这种调度是透明地发生的。也就是说，您不必考虑太多。 但是，您可能想要使用预定义的整数向量 threadIdx（我的线程在块中的哪里？）、blockDim（一个块有多大？）、blockIdx（我的块在网格中的哪里？）和gridDim（网格有多大？）在线程之间拆分工作（读取：输入和输出）。您可能还想了解如何有效地访问不同类型的内存（因此可以在单个事务中为多个线程提供服务） - 但这是题外话。

NSight 提供了一个图形调试器，让您在穿越行话丛林后对设备上发生的事情有一个很好的了解。对于那些你在调试器中看不到的东西（例如停顿原因或内存压力），它的分析器也是如此。

您可以通过另一个内核启动来同步网格内的所有线程（所有线程）。 对于非重叠、顺序内核执行，不需要进一步同步。

一个网格中的线程（或一个内核运行 - 无论您想如何称呼它）可以使用原子操作（用于算术）或适当的内存栅栏（用于加载或存储访问）通过全局内存进行通信。

您可以使用内部指令__syncthreads() 同步一个块内的所有线程（之后所有线程都将处于活动状态 - 尽管与往常一样，最多两个 warp 可以在 Fermi GPU 上运行）。一个块中的线程可以使用原子操作（用于算术）或适当的内存栅栏（用于加载或存储访问）通过共享或全局内存进行通信。

如前所述，warp 中的所有线程总是“同步”的，尽管有些线程可能处于非活动状态。它们可以通过共享或全局内存（或即将推出的具有计算能力 3 的硬件上的“通道交换”）进行通信。您可以使用原子操作（用于算术）和 volatile 限定的共享或全局变量（加载或存储访问在同一个 warp 中顺序发生）。 volatile 限定符告诉编译器始终访问内存，而不是其他线程无法看到其状态的寄存器。

此外，warp 范围内的投票函数可以帮助您做出分支决策或计算整数（前缀）总和。

好的，基本上就是这样。希望有帮助。写得很好:-)。

【讨论】：

感谢您的回复..我帮了很多忙。你还能说出每个线程是如何启动内核的吗？

“每个线程”是什么意思？设备线程无法在计算能力 3 之前启动内核（还没有硬件输出）。否则，它们从一个或多个主机线程中启动。在高端显卡上，可以使用多个主机线程来控制并发主机设备数据传输。

【参考方案2】：

让我们举一个 4*4 矩阵相加的例子。你有两个矩阵 A 和 B，尺寸为 4*4..

int main()

 int *a, *b, *c;            //To store your matrix A & B in RAM. Result will be stored in matrix C
 int *ad, *bd, *cd;         // To store matrices into GPU's RAM. 
 int N =4;                 //No of rows and columns.

 size_t size=sizeof(float)* N * N;

 a=(float*)malloc(size);     //Allocate space of RAM for matrix A
 b=(float*)malloc(size);     //Allocate space of RAM for matrix B

//allocate memory on device
  cudaMalloc(&ad,size);
  cudaMalloc(&bd,size);
  cudaMalloc(&cd,size);

//initialize host memory with its own indices
    for(i=0;i<N;i++)
      
    for(j=0;j<N;j++)
         
            a[i * N + j]=(float)(i * N + j);
            b[i * N + j]= -(float)(i * N + j);
         
      

//copy data from host memory to device memory
     cudaMemcpy(ad, a, size, cudaMemcpyHostToDevice);
     cudaMemcpy(bd, b, size, cudaMemcpyHostToDevice);

//calculate execution configuration 
   dim3 grid (1, 1, 1); 
   dim3 block (16, 1, 1);

//each block contains N * N threads, each thread calculates 1 data element

    add_matrices<<<grid, block>>>(ad, bd, cd, N);

   cudaMemcpy(c,cd,size,cudaMemcpyDeviceToHost);  
   printf("Matrix A was---\n");
    for(i=0;i<N;i++)
    
        for(j=0;j<N;j++)
            printf("%f ",a[i*N+j]);
        printf("\n");
    

   printf("\nMatrix B was---\n");
   for(i=0;i<N;i++)
    
        for(j=0;j<N;j++)
            printf("%f ",b[i*N+j]);
        printf("\n");
    

    printf("\nAddition of A and B gives C----\n");
    for(i=0;i<N;i++)
    
        for(j=0;j<N;j++)
            printf("%f ",c[i*N+j]);   //if correctly evaluated, all values will be 0
        printf("\n");
    



    //deallocate host and device memories
    cudaFree(ad); 
    cudaFree(bd); 
    cudaFree (cd);

    free(a);
    free(b);
    free(c);

    getch();
    return 1;


/////Kernel Part

__global__ void add_matrices(float *ad,float *bd,float *cd,int N)

  int index;
  index = blockIDx.x * blockDim.x + threadIDx.x            

  cd[index] = ad[index] + bd[index];

让我们举一个 16*16 矩阵相加的例子.. 你有两个矩阵 A 和 B，尺寸为 16*16..

首先你必须决定你的线程配置。 您假设启动一个内核函数，它将执行矩阵加法的并行计算，这将在您的 GPU 设备上执行。

现在，一个网格与一个内核函数一起启动.. 一个网格最多可以有 65,535 个可以以 3 维方式排列的块。 (65535 * 65535 * 65535)。

网格中的每个块最多可以有 1024 个线程。这些线程也可以以 3 维方式排列 (1024 * 1024 * 64)

现在我们的问题是添加 16 * 16 矩阵..

A | 1  2  3  4 |        B | 1  2  3  4 |      C| 1  2  3  4 |
  | 5  6  7  8 |   +      | 5  6  7  8 |   =   | 5  6  7  8 | 
  | 9 10 11 12 |          | 9 10 11 12 |       | 9 10 11 12 |  
  | 13 14 15 16|          | 13 14 15 16|       | 13 14 15 16|

我们需要 16 个线程来执行计算。

i.e. A(1,1) + B (1,1) = C(1,1)
     A(1,2) + B (1,2) = C(1,2) 
     .        .          .
     .        .          . 
     A(4,4) + B (4,4) = C(4,4) 

所有这些线程将同时执行。 所以我们需要一个有 16 个线程的块。 为方便起见，我们将在一个块中以 (16 * 1 * 1) 方式排列线程 由于没有 16 个线程，所以我们只需要一个块来存储这 16 个线程。

所以，网格配置将是dim3 Grid(1,1,1)，即网格将只有一个块 并且块配置将是dim3 block(16,1,1)，即块将有 16 个线程按列排列。

以下程序将使您清楚地了解其执行情况。 了解索引部分（即 threadIDs、blockDim、blockID）是重要的部分。您需要阅读 CUDA 文献。一旦你对索引有清晰的想法，你将赢得半战！因此，当然要花一些时间阅读 cuda 书籍、不同的算法和纸笔！

【讨论】：

您显示的是 4 * 4 矩阵，而不是 16 * 16 矩阵。

@RobertCrovella：已更正！谢谢

您遗漏了一些参考资料。

【参考方案3】：

试试'Cuda-gdb'，它是 CUDA 调试器。

【讨论】：

这如何回答这个问题？

在Cuda-gdb中，可以看到内核是如何执行的。

NVIDIA NSIGHT 是否也这样做？

是的。 NSight 是一个集成了 IDE 的图形调试器。它尚未正式在 Linux 上可用，但有一个基于 Eclipse 的 CUDA 5 预览版。