如何推出CUDA内核？

Question

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

我知道这是一个非常基本的概念,但我不知道这一点.关于流量我很困惑.

Answer 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-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个块,可以以三维方式排列.(65535*65535*65535).

网格中的每个块最多可以有1024个线程.这些线程也可以以三维方式排列(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书籍,不同的算法和纸笔当然!