CUDA内核是如何启动的？

Question

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

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

Answer 1

您启动一个网格块。

块被分地分配给多处理器（其中在多处理器块的数目确定可用共享存储器的量）。

块进一步分裂成经纱。对于32个线程的Fermi GPU来说，这些线程执行相同的指令或者是不活动的（因为它们分支离开，例如通过在同一个warp内退出循环之前的循环，或者没有使用它们的if）。在Fermi GPU上，最多两个经线一次在一个多处理器上运行。

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

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

NSight提供了图形化调试器，让你发生了什么设备上，一旦你通过行话丛林得到了一个好主意。对于那些在调试器中看不到的东西（例如失速原因或内存压力），它的分析器也是如此。

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

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

您可以使用内部指令__syncthreads()同步所有线程（所有线程之后都会激活 - 尽管一如既往，最多只能在Fermi GPU上运行两个线程）。一个块内的线程可以使用原子操作（用于算术）或适当的内存屏蔽（用于加载或存储访问）通过共享或全局内存进行通信。

正如前面提到的，经中的所有线程都始终“同步”，虽然有些可能是无效的。他们可以通过共享或全局存储器进行通信（或即将推出的具有计算能力的硬件上的“通道交换”3）。您可以使用原子操作（用于算术）和volatile限定的共享或全局变量（加载或存储访问在相同的warp内顺序进行）。volatile限定符告诉编译器总是访问内存，并且永远不会注册其他线程无法看到的状态。

此外，还有全范围的投票功能，可以帮助您制定分支决策或计算整数（前缀）总和。

好的，基本上就是这样。希望有所帮助。有一个很好的流动写作:-)。

Answer 2

尝试'Cuda-gdb'，这是CUDA调试器。

Answer 3

让我们除了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个块，它们可以以三维方式排列。 （65535 * 65535 * 65535）。

在网格中的每个块可以有最高1024不threads.Those线程也可以安排在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书籍，不同的算法和纸笔当然！