使用gpu实现多线程

在配置GPU时一般都看重其的架构，流处理器数，以及显存数。

以英伟达的GPU为例架构一般以科学家的名字来命名，如Fermi(费米），Kepler(开普勒），现在主流的Maxwell（麦克斯韦），Pascal（帕斯卡），不同的架构主要体现在如纹理单元，流处理器，带宽等较为底层的东西不同，为线程与块中主要关心的是其流多处理器（streaming multiprocessor，SM）以及一个流多处理器包含的多个流处理器（scalar processor,SP） 或称为CUDA核（CUDA core）。当控制器将一个线程分配给一个流多处理器后，流多处理器的核协调工作，并行处理所有的线程。

在并行运算时，会出现网格（grid），线程块（block），线程（thread）三个常见的概念。一个网格下包含一个或多个线程块，一个线程块下包含多个线程。

在使用时 block和grid都可以用三维的向量表示，其中block向量的元素是thread，grid的元素是block。当然线程块内的线程数不是无限的，如先前的G80一个线程块内线程数最多为512个，Fermi架构下的线程块的线程数增加到1024个，Kepler架构下的线程数达到了2048个。

如图中grid block,可用dim3来表示三维的向量：

dim3 `gridSize(3×2×1）；`
dim3 blockSize（2，2，2）；
kernel<<<gridSize,blockSize>>>();

如需二维，则定义时将第三元缺省，如果需一维，则可以直接用整型量int来表示线程和块数；
在核函数kernel中线程是并行执行的，而当一个线程中需要用到另一个线程的量时就需要用到线程的索引：
grid内每个block的位置可以通过blockIdx变量来获得,block的大小可以由变量blockDim来获得,thread在block中位置可以由threadId来获得;

下面再来看先前的两个数组相加的程序，这次我们加上线程索引和块索引；并且将数组的大小扩大到1000：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<stdlib.h>
#include <stdio.h>

cudaError_t addWithCuda(int *c,  int *a,  int *b, unsigned int size);

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

int main()
{
     const int arraySize = 1000;
    int a[arraySize];
    int b[arraySize];
    int c[arraySize] = { 0 };

    for (int i = 0; i < arraySize; i++)
    {
        a[i] = rand() % 100;
        b[i] = rand() % 100;

    }
    // Add vectors in parallel.
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }
    for (int i = 0; i < 1000; i++)
    {
        printf("c[%d]=%d\n", i, c[i]);
    }

    getchar();
    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c,  int *a,  int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;


    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<4,256>>>(dev_c, dev_a, dev_b);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}

这里单个块内的线程数为256，块的个数为4。
块内线程数的选择有一个warp的概念，一个block内的线程同时也被打包成多个warp，一个warp由32各线程组成，所以为了让资源不浪费，即warp内的线程数被充分利用，线程数经常设置为32的整数倍。