cuda之线程分配

线程是cuda编程的核心

原文请查看：http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model

为方便起见，threadidx是一个三维向量，使线程可以确定使用一个一维、二维、三维线或指标，形成一维，二维，或三维块的线程，称为一个线程块。这提供了一种自然方法来调用域中元素的计算，如向量、矩阵或卷。

索引的线程和线程ID以简单的方式相互关联：一个一维的块，它们是相同的；一个二维块大小（dx，dy），一个线程索引线程ID（x，y）是（x + y DX）；一尺寸三维块（dx，dy，dz），一个线程索引线程ID（x，y，z）是（X + Y + Z DX Dx Dy）。

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

每个块的线程数是有限制的，因为块的所有线程都预期位于相同的处理器核心上，并且必须共享该核心的有限内存资源。在目前的GPU的线程块可以包含多达1024个线程。

然而，一个内核可以由多个相同形状的线程块执行，因此线程的总数等于每个块的线程数和块的个数。

块被组织成一维、二维或三维的线程块网格。网格中线程块的数量通常取决于正在处理的数据的大小或系统中处理器的数量，这大大超过了它的数量。

线程的数量每块和多块网格中指定的每个<< <…> > >语法的类型可以是int或dim3。可以在上面的例子中指定二维块或网格。

网格内的每个块可以通过一维、二维或三维指数，在内核通过内置的blockidx变量访问。该线程块尺寸可在内核通过内置的blockdim变量。

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

一个16x16的线程块大小（256个线程），虽然在这种情况下，任意的，是一种常见的选择。网格的创建有足够的块，每个矩阵元素有一个线程和以前一样。为了简单起见，这个示例假设每个维度中每个网格的线程数可以被该维度中每个块的线程数均匀地整除，尽管情况并非如此。

线程块需要独立执行：必须以并行或串行的顺序执行它们。这个独立性要求允许线程块以任意顺序排列在任意数量的内核中，如图5所示，使程序员能够编写与内核数量相匹配的代码。

块中的线程可以通过共享内存共享数据，并通过同步执行来协调内存访问。更确切地说，人可以指定同步点的内核通过调用__syncthreads()固有功能；__syncthreads()作为一个屏障，块中的所有线程必须在任何继续等待。共享内存提供了使用共享内存的例子。

有效的合作，共享内存将在每个处理器核心的低延迟内存（很像一个L1缓存）和__syncthreads()有望轻。