5 线程网格、线程块以及线程（2）

5.3 线程块

即使有512个线程，也不能让你在GPU上斩获颇丰。对很多在CPU上编程的编程人员来说，这似乎是一个很大的数量，但其实在GPU上编程的时候，512个线程并不一定会让你获得很高的效益，对于GPU而言，通常我们可能会创建成千上万个并发线程来实现设备上的高吞吐量。   

与前面线程一节所讲的一样，num_blocks 是内核调用中<<<和>>>的第一个参数:

kernel_function<<<num blocks，num_threads>>>(paraml,param2,…)

如果我们将这个参数从1修改成2，那么就是告诉GPU硬件，我们将启动两倍于之前线程数量的线程，例如:

some_kernel_func<<<2，128>>>(a，b，c);

这将会调用名为some_kernel_func的GPU函数共2x128次，每次都是不同的线程。然而，这样做通常会使 thread_idx参数的计算变得更加复杂，而thread_idx通常又用来表示数组的位置下标。因此，我们之前简单的内核就要稍作调整:  

__global__ void some_kernel_func(int* const a, constint* const b, const int * const c)

{

const unsigned int thread idx = (blockIdx.x * blockDim.x)+ threadIdx.x;

a[thread_idx] = b[thread_idx] * c[thread_idx];

    }

为了计算 thread_idx这个参数，我们必须考虑线程块的数量。对第一个线程块而言blockIdx.x是0，因此 thread_idx直接就等于之前使用过的 threadIdx.x，然而，对于第二个线程块而言，它的blockIdx.x的值是1，blockDim.x表示本例中所要求的每个线程块启动的线程数量，它的值是128，那么对第二个线程块而言，在计算thread_idx时，要在threadIdx.x的基础上加上一个1x128 线程的基地址。

不知道你有没有注意到，在介绍线程块加一个基地址的时候有一个错误?现在我们一共启动了256个线程，数组的下标是0~255，如果不更改数组的大小，那么第128个元素~256个元素，将会出现元素访问和写人越界的问题。这种数组越界错误是不会被编译器发现的，程序代码也会根据数组a边界之外的内容来正常执行，因此在调用内核函数的时候要尽量小心，避免这种内存越界访问错误。

对这个例子而言，我们使用一个128字节大小的数组，并将启动的两个线程块中的每个线程块的线程数量改成 64:

some_kernel_func<<<2，64 >>>(a，b，c);

你可以从图 5-8中看到它的表示。

图5-8 线程块映射的地址空间

注意，尽管我们启动了两个线程块，但thread_idx这个参数依然同之前一样等于数组的下标。那么，我们使用线程块的意义究竟在哪?在这个简单的例子中，很明显，没有什么意义。但是在很多现实问题中，我们将不仅仅只处理512个元素，很可能更多。事实上，查看线程块的数量限制，你会发现可以使用65536个线程块。

如果使用65536个线程块，每个线程块启动512个线程，那么我们一共可以调度33 554432(大约3350万)个线程。如果每个线程块启动512个线程，那么每个SM最多可以处理3个线程块。事实上，这个限制是基于每个SM最多能处理的线程数量。在最新的费米架构的硬件上，每个SM每次最多能执行1536个线程，而在G80的硬件上，只能执行768 个线程。

如果你打算在费米架构的硬件上每个线程块调度1024个线程，那么65536个线程块-共就能调度接近6400万个线程，但很不幸的是，如果每个线程块是1024个线程，那每个SM 每次最多运行一个线程块。造成的结果是，除非每个SM 分配执行的线程块数量是一个以上，否则在单个GPU上你将需要65536个SM来执行你的程序。目前，在任何GPU上SM 的最大数目都是30。因此，在线程块需要的SM 数量超出硬件支持的 SM 数量之前CUDA 提供了一定的处理机制。这正是CUDA的迷人之处--它能扩展为上千个执行单元。并行的极限仅受限于应用程序可以分解的并行程度。

假定现在有6400万个线程，每个线程处理数组的一个元素，那么一共可以处理6400万个数组元素。假定数组的每个元素都是一个单精度的浮点数，那么每个元素占4个字节，总共大约需要2亿5600万个字节，即约 256MB 数据存储空间。而几乎所有的GPU都至少支持这个大小的内存空间。因此，仅使用线程和线程块就可以达到相当大量的并行性和数据覆盖。

很多人会担心大规模数据集问题。它们可能是 GB 级、TB 级，甚至 PB级的大规模数据。对于这类问题，这里提供多个解决方案。我们通常会选择一个线程处理多个元素或者使用线程块的其他维度来处理。接下来的小节我们将会进行详细介绍。

线程块的分配

为了确保能够真正地了解线程块的分配，接下来我们写一个简短的内核程序来输出线程块、线程、线程束和线程全局标号到屏幕上。现在，除非你使用的是3.2版本以上的SDK,否则内核中是不支持printf的。因此，我们可以将数据传送回CPU端然后输出到控制台窗口，内核的代码如下:

__g1obal__ void what_is_my_id(unsigned int*const block.
unsigned int*const thread.
unsigned int*const warp,
unsigned int*const calc thread)
{
    /* Thread id is block index * block size + thread offset into the block */
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;

    block[thread_idx] = blockIdx.x;
    thread[thread_idx] = threadIdx.x;

    /* Calculate warp using built in variable warpSize */
    warp[thread_idx] = threadIdx.x/warpSize;
    calc_thread[thread_idx] = thread_idx;
}

在CPU端，我们需要执行下面的一部分代码来为数组在 GPU上分配内存以及将算好的数组数据从 GPU 端复制回来并在 CPU 端显示。
#include <stdio.h>
#include <stdlib.h>
#include <conio.h>

__global__ void what_is_my_id(unsigned int * const block,
    unsigned int*const thread,
    unsigned int*const warp,
    unsigned int*const calc thread)
{

    /* Thread id is block index * block size + thread offset into the block */
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;

    block[thread_idx] = blockIdx.x;
    thread[thread_idx] = threadIdx.x;

    /* Calculate warp using built in variable warpSize */
    warp[thread_idx] = threadIdx.x/warpSize:

    calc_thread[thread_idx] = thread_idx;
}

#define ARRAY_SIZE 128
#define ARRAY_SIZE_IN_BYTES(sizeof(unsigned int)*(ARRAY_SIZE))

/* Declare statically four arrays of ARRAY_SIZE each */
unsigned int cpu_blocK[ARRAY_SIZE];
unsigned int cpu_thread[ARRAY_SIZE]:
unsigned int cpu_warp[ARRAY_SIZE];
unsigned int cpu_calc_thread[ARRAY_SIZE];

int main(void)
{
    /*Total thread count = 2 * 64 = 128 */
    const unsigned int num_blocks = 2:
    const unsigned int num_threads =64:
    char ch;

    /* Declare pointers for GPU based params */
    unsigned int *gpu_block;
    unsigned int *gpu_thread;
    unsigned int *gpu_warp;// 线程束
    unsigned int *gpu_calc_thread;

    /* Declare loop counter for use later */
    unsigned int i:

    /* Allocate four arrays on the GPu */
    cudaMalloc((void **)&gpu_block, ARRAY_SIZE_IN_BYTES);
    cudaMa11oc((void **)&gpu_thread, ARRAY_SIZE_IN_BYTES);
    cudaMa11oc((void **)&gpu_warp, ARRAY_SIZE_IN_BYTES);
    cudaMa11oc((void **)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES);

    /*Execute our kernel*/
    what_is_my_id<<<num_blocks, num_threads>>>(gpu_block, gpu_thread,             gpu_warpgpu_calc_thread);

    /*Copy back the gpu results to the CPU */
    cudaMemcpy(cpu_block,gpu_block,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost);
    cudaMemcpy(cpu_thread,gpu_thread,ARRAY_SIZE_IN_BYTES,cudaMemcpyDeviceToHost);
    cudaMemcpy(cpu_warp,gpu_warp,ARRAY_SIZE_IN_BYTEScudaMemcpyDeviceToHost);
    cudaMemcpy(cpu_calc_thread, gpu_ calc_thread,         ARRAY_SIZE_IN_BYTEScudaMemcpyDeviceToHost);

    /* Free the arrays on the GPU as now we're done with them */
    cudaFree(gpu_block);
    cudaFree(gpu_thread);
    cudaFree(gpu_warp);
    cudaFree(gpu_calc_thread);


    /*Iterate through the arrays and print */
    for(i = 0;i < ARRAY_SIZE; ++i)
    {
        printf("Calculated Thread:%3u-Block:%2u-Warp %2u- Thread %3u\n"
        cpu_calc_thread[i],cpu_block[i], cpu_warp[i], cpu_thread[i]);


        ch = getch();
    }
}

在这个例子中，我们可以看到线程块按照线程块的编号紧密相连。由于处理的是一维数组，所以我们对线程块采用相同的布局便可简单解决问题。以下是此程序的输出结果:

正如我们计算的那样，线程索引是0~127。一共有两个线程块，每个线程块包含64个线程，每个线程块内部线程的索引为0~63。一个线程块包含两个线程束。