parallel-processing - CUDA 内核是如何启动的？

Question

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

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

Answer 1

您启动一个块网格。

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

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

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

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

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

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

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

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

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

此外，还有 warp-wide 投票函数可以帮助您做出分支决策或计算整数（前缀）总和。

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

Answer 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 书籍、不同的算法和纸笔！

Answer 3

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