线程层次的概念：
简单说，就是一个grid有多个block,一个block有多个thread.
grid有多大，用gridDim表示它有多少个block，具体分为gridDim.x, gridDim.y，gridDim.z。
block有多大，用blockDim表示它有多少个thread，具体分为blockDim.x，blockDim.y，blockDim.z。
怎么表示thread在block中的相对位置呢？用 threadIdx.x,threadIdx.y,threadIdx.z表示。
怎么表示block在grid中的相对位置呢？用blockIdx.x,blockIdx.y,blockIdx.z表示。
顺便解释下 华为云论坛_云计算论坛_开发者论坛_技术论坛-华为云 中hello_from_gpu<<<x,y>>>(); 中的x和y是什么意思？它们分别表示 gridDim和blockDim。
对于下面这个函数：
表示gridDim是1，表示grid有1个block，blockDim是4。表示block有4个thread。
所以对于上面的核函数，相当于有4个thread分别执行了 c[n]=a[n]+b[n]的操作，n=threadIdx.x
在调用的时候，所有的CUDA核都是执行同一个函数。这与CPU多线程可能会执行不同的任务不同。
如上图所示，Thread在CUDA core中执行，Block在 SM中执行，Grid在Device中执行。
那么，CUDA是如何执行的呢？看下面这张图：
如果没有block的概念，要同时进行同步、通信、协作时，整体的核心都要产生等待的行为，如要进行扩展时，扩展的越多等待也越多。所以性能会受影响。
但是有block的概念后，可以实现可扩展性。用block或warp就可以很容易实现扩展了。
如何找到线程该处理的数据在哪里呢？这就要提到线程索引的概念。
以上：假定每8个thread时一个block。
具体的公式如下：
具体的索引位置 index = blockDim.x * blockIdx.x + threadIdx.x
那么一个CUDA程序到底应该怎么写呢？
以将一个CPU实现的代码转换为GPU为例：
CPU的实现过程大致如下：
（1）主程序main：
先分配 源地址空间a,b，目的地址空间c，并生成a，b的随机数。然后调用 一维矩阵加的CPU函数。
（2）一维矩阵加的CPU函数：
遍历a,b地址空间，分别将 a[i] 与 b[i]相加，写入 c[i]地址。
这个时候，请注意是要显式地进行for循环遍历。
那么，GPU该如何实现呢？
（1）主程序main：
因为GPU存在Host和Device内存，所以先申请host内存h_a,h_b，存放a，b的一维矩阵的内容（也可以生成随机数），并申请host内存h_c存放c的计算结果。
然后申请device内存，这个时候，需要申请 d_a,d_b两个源device内存(cudaMalloc)，以及d_c这个目的device内存(cudaMalloc)。将h_a和h_b的内容拷贝到d_a和d_b (显然需要使用 cudaMemcpyHostToDevice)；
然后调用核函数完成GPU的并行计算，结果写入h_c；
最后将d_c的device内存写回到h_c(cudaMemcpyDeviceToHost)，并释放所有的host内存(使用free)和device内存（使用cudaFree)。
（2）核函数
这里就是重点了。核函数只需要去掉最外层的循环，并且根据前面 的index写法，将i替换成index的写法即可。
如何设置Gridsize和blocksize呢？
对于一维的情况：
block_size=128;
grid_size = (N+ block_size-1)/block_size;
（没有设成什么值是最好的）
每个block可以申请多少个线程呢？
总数也是1024。如（1024，1，1）或者（512，2，1）
grid大小没有限制。
底层是以warp为单位申请。 如果blockDim为160，则正好申请5个warp。如果blockDim为161，则不得不申请6个warp。
如果数据过大，线程不够用怎么办？
这样子，每个线程需要处理多个数据。
比如对于上图，线程0，需要处理 0,8,16,24 四个数据。核函数需要将每一个大块都跑一遍。代码如下：
这里引入了一个stride的概念，它的大小为blockDim.x X gridDim.x 。核函数需要完成每个满足 index = index + stride * count对应的相关地址的计算。
范例1：体验index
Index_of_thread.cu
#include <stdio.h>__global__ void hello_from_gpu()
{//仅仅是在原先代码的基础上打印 blockIdx.x 和 threadIdx.xconst int bid = blockIdx.x;const int tid = threadIdx.x;printf("Hello World from block %d and thread %d!\n", bid, tid);
}int main(void)
{hello_from_gpu<<<5, 5>>>();//记得加上同步，不然结果会出不来。cudaDeviceSynchronize();return 0;
}
Makefile：
TEST_SOURCE = Index_of_thread.cuTARGETBIN := ./Index_of_threadCC = /usr/local/cuda/bin/nvcc$(TARGETBIN):$(TEST_SOURCE)$(CC) $(TEST_SOURCE) -o $(TARGETBIN).PHONY:clean
clean:-rm -rf $(TARGETBIN)
编译并执行：
范例2：完成一维向量计算：add
vectorAdd.cu
#include <math.h>
#include <stdio.h>void __global__ add(const double *x, const double *y, double *z, int count)
{const int n = blockDim.x * blockIdx.x + threadIdx.x;//这里判断是防止溢出if( n < count){z[n] = x[n] + y[n];}}
void check(const double *z, const int N)
{bool error = false;for (int n = 0; n < N; ++n){//检查两个值是否相等，如不等则error=true.if (fabs(z[n] - 3) > (1.0e-10)){error = true;}}printf("%s\n", error ? "Errors" : "Pass");
}int main(void)
{const int N = 1000;const int M = sizeof(double) * N;//分配host内存double *h_x = (double*) malloc(M);double *h_y = (double*) malloc(M);double *h_z = (double*) malloc(M);//初始化一维向量的值for (int n = 0; n < N; ++n){h_x[n] = 1;h_y[n] = 2;}double *d_x, *d_y, *d_z;//分配device内存cudaMalloc((void **)&d_x, M);cudaMalloc((void **)&d_y, M);cudaMalloc((void **)&d_z, M);//host->devicecudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);//这个是公式。记住就可以了。const int block_size = 128;const int grid_size = (N + block_size - 1) / block_size;//核函数计算add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);//device->hostcudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);//检查结果check(h_z, N);//释放host内存free(h_x);free(h_y);free(h_z);//释放device内存cudaFree(d_x);cudaFree(d_y);cudaFree(d_z);return 0;
}Makefile-add
TEST_SOURCE = vectorAdd.cuTARGETBIN := ./vectorAddCC = /usr/local/cuda/bin/nvcc$(TARGETBIN):$(TEST_SOURCE)$(CC) $(TEST_SOURCE) -o $(TARGETBIN).PHONY:clean
clean:-rm -rf $(TARGETBIN)编译后执行：
