参考：
优达学城：https://classroom.udacity.com/courses/cs344/lessons/55120467/concepts/670611900923
一、典型GPU程序构成
一个典型GPU程序有如下几个部分：
①CPU在GPU上分配内存
②CPU将CPU中的数据copy到GPU中
③调用内核函数来处理数据
④CPU将GPU中的数据copy到CPU中
*可以看出，四个步骤中有两个是数据的copy，因此如果你的程序需要不断地进行copy，那么运行效率会比较低，不适合利用GPU运算。
一般情况下，最好的方式是，让GPU进行大量运算，同时保证计算量与通信的比值很高。
二、运行程序时GPU是如何运作的？
前面提到，我们将运算看作是一个由一系列或多个内核组成的结构。GPU有很多并行的计算单元，那么当我们写kernels的时候，应该充分利用这些硬件并行。
★Big Idea（one of the very core concept of CUDA）：
当你在写程序的时候，你把它写的像是一个串行的程序，就好像是在一个单独的线程里运行一样。当你从CPU调用核的时候，就给核规定好要启动几个线程，每一个线程都会运行这个程序（的一部分？，存疑）。
小结：
GPU擅长：①启动很多的线程 ②并行地运行很多线程
因此，如果你的程序没有利用它进行多线程计算，那么就没有很好的利用GPU。
三、一个典型的程序
程序要求：对一个数组当中的所有元素进行平方计算。
先看一下C语言版本的程序：
for (int i = 0; i < 64; i++){
out[i] = in[i] * in[i];
}C语言的程序有以下两个特点：
①只有一个线程在执行计算，该线程遍历输入数组内所有元素。（将该线程定义为“执行代码的一个独立路径”，这一定义也将适用于GPU代码中。）
②代码中没有显式并行性，是串行代码。这段代码是一个循环64次的线程，每一个循环进行一次平方计算。
对于GPU来说，程序要分为两个部分：①在CPU上运行 ②在GPU上运行的
GPU上运行的程序很简单，可以看作是一个普通的串行运算，在这个例子里就是 OUT = IN * IN，即计算平方。那么这个核函数就可以看作跟“并行”无关。
CPU上的程序则要负责分配内存、从GPU复制数值、将数值传递给GPU、启动Kernel函数。
*CPU启动kernels的过程：用64个线程，启动一个执行平方计算的kernel，即square_kernel。每一个内核的实例（instance）将运算一个平方计算。
*线程索引：分辨每一个独立的内核做的工作。
CPU code：
square_kernel<<<64>>>(input_array, output_array)完整的代码如下：
#include <stdio.h>
//内核函数，就跟一个串行程序一样
//__global__ : 定义说明符，用来表明这是一个内核，而非普通函数
//kernel的参数都必须是分配在GPU上的，不然会程序崩溃
__global__ void square(float * d_out, float * d_in){
int idx = threadIdx.x; //用来得到线程的索引值,threadIdx是一个CUDA内置变量，其本质是一个C结构，有三个成员，x、y、z。
float f = d_in[idx];
d_out[idx] = f * f;
}
//下面是CPU代码
int main(int argc, char ** argv) {
const int ARRAY_SIZE = 64;
const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
//generate input array on the host
float h_in[ARRAY_SIZE];
for (int i = 0; i < ARRAY_SIZE; i++){
h_in[i] = float(i);
}
float h_out[ARRAY_SIZE];
//declare GPU memory pointers,看起来跟c program一样，是一个普通的数组
float * d_in;
float * d_out;
//allocate GPU memory，在这里给它分配GPU内存
cudaMalloc((void **) &d_in, ARRAY_BYTES);
cudaMalloc((void **) &d_out, ARRAY_BYTES);
//transfer the array to the GPU，前三个参数跟C Memcpy一样，目标地址、源地址、字节数，最后一个参数表明转移方向cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,cudaMemcpyDeviceToDevice
cudaMencpy(d_in, h_in, ARRAY_BYTES,cudaMemcpyHostToDevice);
//launch the kernel, cuda启动运算符，在一个具有64个元素的block上启动一个核，两个参数d_in,d_out
//这句话告诉CPU，在GPU上启动64个线程，运行64个内核副本
//注意，我们只能在调用内核处理GPU数据，而不能处理CPU数据
square<<<1,ARRAY_SIZE>>>(d_out, d_in);
//copy back the result array to the CPU，完成内核运行以后，结果在d_out中
cudaMemcpy(h_out, d_out, ATTAY_BYTES, cudaMemcpyDeviceToHost);
//print out the resulting array
for (int i = 0; i < ARRAY_SIZE; i++){
printf("%f",h_out[i]);
printf(((i % 4) != 3) ? "\t" : "\n" );
}
//free GPU memory allocation
cudaFree(d_in);
cudaFree(d_out);
return 0;
}
需要注意的是，在CUDA程序中，我们以h_开头表示host上的数据，d_开头表示GPU上的数据。
针对代码，有值得注意的几点：
①square<<<1,64>>>(d_out, d_in)
参数说明：
1：number of blocks
64：Threads Per Block
从硬件角度说，①硬件可以同时运行多个块（blocks），②每个块有线程数限制，较新的GPU一个块可以支持1024个线程，旧的512.
②另外，这是一个一维diagram，只在一个维度发展。如果想要处理2/3等维度问题，则需要如下代码：
Kernel<<<GRID OF BLOCKS, BLOCK OF THREADS，shemem>>>(...)
说明：
GRID OF BLOCKS：线程块网格维数，取值1-3
BLOCK OF THREADS：每个块中线程维数，取值1-3
dim3(x,y,z)：线程/blocks维数 = x*y*z，xyz分别代表三个方向的值
shemem：每一个块共享的内存，字节数
③线程索引
threadIdx：线程索引，有threadIdx.x，threadIdx.y，threadIdx.z来找三个方向的索引。
blockIdx:块索引，同上。
*CUDA擅长处理多块、多线程、多维度问题。