1、CUDA程序结构

　　CUDA程序是在主机或者设备上执行的函数的组合。不显示并行性的函数在CPU上执行，显示数据并行性的函数在GPU上执行，GPU在编译期间要隔离这些函数。CUDA代码基本上与C代码相同，只是添加了一些开发数据并行性所需的关键字。

２、CUDA　C中的双变量加法程序

　　编写一个将两个变量相加的内核（kernel）函数

#include<iostream>
#include<cuda.h>
#include<cuda_runtime.h>
__global__ void gpuAdd(int d_a, int d_b, int* d_c) {
    *d_c = d_a + d_b;
}

　　main函数代码

int main(void) {
    //定义主机变量，统一以h_开头
    int h_c;
    //定义设备变量，统一以d_开头，作为参数传递给内核的指针应该仅指向设备显存
    int* d_c;
    //在设备上分配d_c变量的内存，cudaMalloc类似于C中的malloc函数
    //分配了一个整数变量大小的内存块，并返回指向该内存位置的指针d_c
    cudaMalloc((void**)&d_c, sizeof(int));
    //调用gpuAdd，其中1和4是两个输入变量，d_c是一个作为输出指针变量的设备显存指针
    //
    gpuAdd << <1, 1 >> > (1, 4, d_c);
    //将结果从设备内存复制到主机内存中
    cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
    //释放设备上通过cudaMalloc使用的所有内存
    cudaFree(d_c);
    return 0;
}

　　需要注意的是，作为参数传递给内核的指针应该仅指向设备显存，如果它指向主机内存，会导致程序崩溃。其次，当通过对调用的内核传递参数时i，所有的指针仅指向设备显存。

３、内核调用

　　使用ASNSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码（Host Code）通过内核调用的方式来启动的。简单地说，内核调用的含义是我们从主机代码启动设备代码。内核调用通常会产生大量的块（Block）和线程（Thread）来在GPU上并行地处理数据。内核使用__global__关键字定义，使用kernel<< <希望执行的块数， 每个块将具有的线程数， 内核使用的共享内存大小>> >来启动内核。在变量相加程序中，内核启动的语法如下：

gpuAdd << <1, 1 >> > (1, 4, d_c)

4、CUDA中的向量运算

       上面的双变量加法没有利用到GPU设备的并行处理能力，下面利用GPU的并行结构修改CPU上的向量加法程序，首先在CPU上编写向量加法程序

4.1、CPU向量加法

#include "stdio.h"
#include<iostream>
//定义向量元素个数
#define N    5
//定义CPU向量加法
void cpuAdd(int *h_a, int *h_b, int *h_c) {
    int tid = 0;    
    while (tid < N)
    {
        h_c[tid] = h_a[tid] + h_b[tid];
        tid += 1;
    }
}

int main(void) {
    int h_a[N], h_b[N], h_c[N];
        //向量初始化
    for (int i = 0; i < N; i++) {
        h_a[i] = 2 * i*i;
        h_b[i] = i;
    }
    //调用CPU向量加法函数
    cpuAdd (h_a, h_b, h_c);
    //输出结果
    printf("Vector addition on CPU\n");
    for (int i = 0; i < N; i++) {
        printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    return 0;
}

4.2、GPU向量加法

       在GPU上添加修改后的内核函数：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//定义向量元素个数
#define N 5
//定义向量加法内核函数
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
    //获取当前内核的块ID
    int tid = blockIdx.x;    //在当前块ID来初始化tid
    //因为是同时启动的，每个块必须直到自己要处理的数据是什么
    if (tid < N)
        d_c[tid] = d_a[tid] + d_b[tid];
}

int main(void) {
    //定义主机变量
    int h_a[N], h_b[N], h_c[N];
    //定义设备指针，指向cudaMalloc分配的显存
    int *d_a, *d_b, *d_c;
    //给设备变量分配内存
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));
    //主机向量初始化
    for (int i = 0; i < N; i++) {
        h_a[i] = 2*i*i;
        h_b[i] = i ;
    }
    //将主机向量复制到设备向量中
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
    //使用内核启动符号<<<>>>调用内核函数，启动N个块（和向量长度一样，使用不同的块并行处理向量N个位置的加法），每个块1个线程，传递参数为设备指针
    gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
    //将结果从设备内存复制到主机内存中
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    printf("Vector addition on GPU \n");
    //在控制台中输出结果
    for (int i = 0; i < N; i++) {
        printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    //释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    return 0;
}

       在main函数中调用内核函数时，并行启动了N个块，这意味着我们同时启动了N个执行该内核代码的线程副本。因为N个块是同时启动的，只有每个块分工合作效率才会高，main函数调用gpuAdd内核函数时，启动的块是随机的，但可以通过blockIdx.x查看启动的块的ID，用该ID作为块处理的数据对应的ID，这样就可以让N个块分别处理向量中N个数据的加法。GPU的并行计算的方式比CPU的串行计算提高了吞吐量。

4.3、GPU对向量每个元素进行平方

       上面是并行启动N个块，每个块只有一个线程。也可以在1个块里面并行启动N个线程。

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//定义向量元素个数
#define N    5
//平方内核函数
__global__ void gpuSquare(float *d_in, float *d_out) {
    //获取当前调用内核函数启动的线程ID
    int tid = threadIdx.x;    //使用内核函数计算线程ID对应的向量中的元素
    float temp = d_in[tid];
    d_out[tid] = temp*temp;
}

int main(void) {
    //定义主机向量
    float h_in[N], h_out[N];
    //定义设备指针
    float *d_in, *d_out;

    //在CPU上分配内存
    cudaMalloc((void**)&d_in, N * sizeof(float));
    cudaMalloc((void**)&d_out, N * sizeof(float));
    //Initializing Array
    for (int i = 0; i < N; i++) {
        h_in[i] = i;
    }
    //将主机上的数组复制到设备上
    cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
    //调用1个块，N个线程的内核函数
    gpuSquare << <1, N >> >(d_in, d_out);
    //将设备上的结果复制到主机内存上
    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    //控制台输出结果
    printf("Square of Number on GPU \n");
    for (int i = 0; i < N; i++) {
        printf("The square of %f is %f\n", h_in[i], h_out[i]);
    }
    //释放内存
    cudaFree(d_in);
    cudaFree(d_out);
    return 0;
}

       当使用这种方式启动N个线程时，每个块的最大线程数不超过512或1024，如果N=2000，而GPU的线程最大数量为512，那么不能写成<<<1, 2000>>>，而应该用<<<4, 500>>>。应当选择合适的块和线程的数量。

5、并行通信模式

       当多个线程并行执行时，他们遵循一定的通信模式，直到它们在显存哪里输入，哪里输出，通讯模式包括映射、收集、分散式、蒙版、转置。

5.1、映射（Map）

       单一输入，单一输出。Map的代码模式看起来如下：

d_out[i] = d_in[i] * 2;

5.2、收集（Gather）

      多个输入，单个输出。Gather模式的代码看起来如下：

out[i] = (in[i - 1] + in[i] + in[i + 1]) / 3;

5.3、分散式（Scatter）

      单一输入，多个输出。Scatter模式的代码看起来如下：

out[i - 1] += 2 * in[i] and out[i + 1] += 3 * in[i];

5.4、蒙版（Stencil）

      从数组中读取固定形状的相邻元素时，称为stencil模式。比如3X3或者5X5卷积就是Stencil，它是Gather操作的一种模式，代码和Gather相似。

5.5、转置（Transpose）

      对矩阵进行转换。