CUDA 线程如何工作答案

【问题标题】：How do CUDA threads workCUDA 线程如何工作
【发布时间】：2016-02-27 21:57:44
【问题描述】：

我对线程的形成和执行方式有很多疑问。

首先，文档将 GPU 线程描述为轻量级线程。假设我希望将两个 100*100 矩阵相乘。如果每个元素由不同的线程计算，这将需要100*100 线程。但是，我的 GPU（NVIDIA GT 640M LE）规格显示了两个 SM，每个 SM 只能支持 2048 个线程。考虑到我的 GPU 无法支持这么多线程，如何计算其余两个元素并行 y。

还要考虑基本的向量添加代码。假设我调用一个具有 1 个块和 64 个线程的内核来添加两个数组，每个数组包含 100 个元素，如下所示：

    __global__ void add(int* a,int* b,int* c)
    {
        int i = threadIdx.x;
        for(i<100)
        {
            c[i] = a[i] + b[i];
        {    
     }

由于只初始化了 64 个线程，我假设并行添加了 64 个元素。

其余元素是如何添加的？
warp 调度程序如何决定分配哪些线程来添加最后 36 个元素？

我的主要问题是：

我不明白线程如何知道要操作哪些元素。

【问题讨论】：

标签： cuda

【解决方案1】：

您的卡具有计算能力3.0，请参阅here。

根据 CUDA C 编程指南的表 12，您提到的计算能力的 2048 线程数是指每个多处理器的最大驻留线程数。这并不意味着您不能启动超过2048 的线程。例如，从该表上方的几行中，您可以看到线程块网格的最大 x-dimension 为 2^31-1。这意味着启动例如1d 线程网格是完全合法的，例如8192 线程。原因是该卡将执行线程扭曲之间的上下文切换，如本文所示：What is the context switching mechanism in GPU?。

关于您问题的第二部分，您对 add 函数的实现在概念上是错误的。您将索引i 用作线程索引和for 循环索引。更正确的实现如下

__global__ void add(int* a,int* b,int* c)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

上面写法的意思是：每个线程都会执行两个赋值，即

    int i = threadIdx.x;
    c[i] = a[i] + b[i];

现在，例如，对于线程 #3，threadIdx.x 变量的值将是 3。因此，线程#3 将处理其内存空间私有的局部变量i，其值将分配给3。此外，它将从全局内存中加载a[3]和b[3]，将它们相加，将结果分配给c[3]，然后将最终结果存储到全局内存中。因此，当您启动网格时，您当然不能仅通过 64 线程填充整个 100 元素数组，您将需要 100 线程。

请注意，上述解释过于简单。我建议您阅读一些基础教科书，例如著名的 CUDA By Example。

【讨论】：

感谢您的回答。关于启动大量线程的部分我明白了。
进入第二部分，我仍然不明白其余元素是如何添加的。由于仅启动了 64 个线程，因此块上仅存在两个 warp。添加 64 个元素后，是否会添加剩余的元素？如果是，那么由哪些线程添加？
@mastercheif141 我试图更好地解释你的第二个问题。请参阅编辑后的答案。

【解决方案2】：

将为您提供 CUDA 中 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 =16;   

          //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 书籍... :-)

【讨论】：

【解决方案3】：

for 在这里是非常错误的 - 一些 threadid

int i = threadIdx.x;
c[i] = a[i] + b[i];

如果您的数组大小为 loke 100 与 64x 的块大小不匹配，为了某些线程不会越界读/写，请执行以下操作：

int i = threadIdx.x;
    if(i < 100){

        c[i] = a[i] + b[i];
    }

您只会在最后一个区块出现分歧。可能你想要那个

【讨论】：