logo标签列表

如何启动CUDA内核？

parallel-processingcudagpgpunvidia
8

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

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

3个回答
12

您启动了一个块的网格。

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

块进一步分为线程束。对于Fermi GPU，它是32个线程，要么执行相同的指令，要么处于非活动状态（因为它们通过提前退出循环或不执行其他线程束中的邻居执行的if等方式进行了分支）。在Fermi GPU上，最多有两个线程束在一个多处理器上运行。

每当存在延迟（即执行由于内存访问或数据依赖性需等待完成而导致的停顿）时，就会运行另一个线程束（相同或不同块的线程束数由每个线程使用的寄存器数和一个/多个块使用的共享内存量确定）。

这种调度是透明的。也就是说，您不必过多考虑它。但是，您可能希望使用预定义的整数向量threadIdx（我的线程在块中的位置？），blockDim（一个块的大小是多少？），blockIdx（我的块在网格中的位置？）以及gridDim（网格的大小是多少？）将工作（读取：输入和输出）分配给线程。您可能还希望了解如何有效地访问不同类型的内存（以便可以在单个事务中服务多个线程）。但这已经偏离了主题。

NSight提供了一个图形化调试器，一旦您克服了术语障碍，就能很好地了解设备上正在发生的情况。对于那些您在调试器中看不到的事情（例如停顿原因或内存压力），其分析器也是如此。

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

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

您可以使用内置指令__syncthreads()同步一个block内的所有线程（所有线程之后都将处于活动状态——但是，在Fermi GPU上最多只能运行两个warp）。一个block中的线程可以通过原子操作（用于算术运算）或适当的内存栅栏（用于加载或存储访问）来使用共享或全局内存进行通信。
正如早先提到的，一个warp内的所有线程始终是“同步”的，尽管有些可能是非活动的。它们可以通过共享或全局内存进行通信（或在具有计算能力3的新硬件上进行“lane swapping”）。您可以使用原子操作（用于算术运算）和具有volatile限定符的共享或全局变量(在同一warp内顺序发生的加载或存储访问)。volatile限定符告诉编译器始终访问内存而不是状态不能被其他线程看到的寄存器。
此外，还有整个warp范围的投票功能，可帮助您做出分支决策或计算整数（前缀）总和。
好的，基本就是这样了。希望对您有所帮助。写作时感觉很畅快哦 :-)。
谢谢您的回复，对我帮助很大。 不过，您能否告诉我每个线程如何启动内核？ - ATG
“每个线程”是什么意思？在计算能力低于3的情况下，设备线程无法启动内核（尚未有硬件输出）。否则，它们将从一个或多个主机线程中启动。在高端图形卡上，可以使用多个主机线程来控制并发的主机<->设备数据传输。 - Dude
8

让我们以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个线程。
以下程序将使您清楚地了解其执行过程。 理解索引部分（即线程ID、blockDim、blockID）是重要的部分。您需要阅读CUDA文献。一旦您对索引有清晰的理解，您就会赢得一半的胜利！因此，请花时间阅读CUDA书籍、不同的算法和纸笔材料！
你正在展示4 * 4矩阵，而不是16 * 16矩阵。 - Robert Crovella
@RobertCrovella：已更正！谢谢 - sandeep.ganage
你错过了几个引用。 - Robert Crovella
0
这个回答如何回答问题？ - talonmies
在Cuda-gdb中，您可以看到内核的执行方式。 - chaohuang
NVIDIA NSIGHT也做同样的事情吗？ - ATG
是的，NSight是一款集成于IDE的图形化调试器。虽然它尚未正式支持Linux，但基于CUDA 5的Eclipse预览版已经发布。 - Dude
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接