1、CUDA线程
CUDA关于并行执行具有分层结构。每次内核启动时可以被切分成多个并行执行的块，而每个块又可以进一步地被切分成多个线程。
在上一推文我们已经知道，maxThreadPerBlock属性限制了每个块能启动的线程数量。这个值对于最新的GPU卡来说是1024。类似地，第二种方式能最大启动的块数量被限制成2^31-1个。
更加理想的则是，我们并不单独启动1个块，里面多个线程；也不启动多个块，每个里面1个线程。我们一次并行启动多个块，每个块里面多个线程（最多可以是maxThread-PerBlock那么多哦）。所以，假设上一章的那个向量加法例子你需要启动N=50000这么多的线程，我们可以这样调用内核：
最大的块能有1024个线程。不过我们这里举例，对于N个线程来说，每个块有512个线程，则需要有N/512个块。但是如果N不是512的整数倍，那么N除以512会计算得到错误的块数量，比实际的块数量少1个。所以为了计算得到下一个最小的能满足要求的整数结果，N需要加上511，然后再除以512。这基本上是一个除法的向上取整操作。
还是直接撸代码吧：
#include "stdio.h" #include<iostream> #include <cuda.h> #include <cuda_runtime.h> //Defining number of elements in Array #define N50000 //Defining Kernel function for vector addition __global__ void gpuAdd(int *d_a, int *d_b, int *d_c) { //Getting block index of current kernel int tid = threadIdx.x + blockIdx.x * blockDim.x; while (tid < N) { d_c[tid] = d_a[tid] + d_b[tid]; tid += blockDim.x * gridDim.x; } } int main(void) { //Defining host arrays int h_a[N], h_b[N], h_c[N]; //Defining device pointers int *d_a, *d_b, *d_c; // allocate the memory cudaMalloc((void**)&d_a, N * sizeof(int)); cudaMalloc((void**)&d_b, N * sizeof(int)); cudaMalloc((void**)&d_c, N * sizeof(int)); //Initializing Arrays for (int i = 0; i < N; i++) { h_a[i] = 2 * i*i; h_b[i] = i; } // Copy input arrays from host to device memory cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice); //Calling kernels with N blocks and one thread per block, passing device pointers as parameters gpuAdd << <512, 512 >> >(d_a, d_b, d_c); //Copy result back to host memory from device memory cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); int Correct = 1; printf("Vector addition on GPU \n"); //Printing result on console for (int i = 0; i < N; i++) { if ((h_a[i] + h_b[i] != h_c[i])) { Correct = 0; } } if (Correct == 1) { printf("GPU has computed Sum Correctly\n"); } else { printf("There is an Error in GPU Computation\n"); } //Free up memory cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; }
本内核的代码和上一文写过的那个很相似。但是有两处不同：
（1）计算初始的tid的时候；
（2）是添加了while循环部分。
计算初始的tid的变化，是因为我们现在是启动多个块，每个里面有多个线程，直接看成ID的结构，多个块横排排列，每个块里面有N个线程，那么自然计算tid的时候是用：
当前块的ID*当前块里面的线程数量+当前线程在块中的ID
即tid=blockIdx.x（当前块的ID）*blockDim.x（当前块里面的线程数量）+threadIdx.x（当前线程在块中的ID）。
而while部分每次增加现有的线程数量（因为你没有启动到N），直到达到N。这就如同你有一个卡，一次最多只能启动100个块，每个块里有7个线程，也就是一次最多能启动700个线程。但N的规模是8000，远远超过700怎么办？答案是直接启动K个（K≥700），这样就能安全启动。然后里面添加一个while循环，这700个线程第一次处理[0,699)，第二次处理[700,1400)，第三次处理[1400,2100)……直到这8000个元素都被处理完。这就是我们本例中看到的代码。初始化时候的tid=threadIdx.x+blockDim.x*blockIdx.x，每次while循环的时候tid+=blockDim.x*gridDim.x（注意一个是=，一个是+=，后者是增加的由来）。下面的2D表格用来辅助理解。
对于任意一个线程，使用blockIdx.x命令可以得到当前的块的ID，而使用threadIdx.x命令可以得到本线程在该块中的ID。例如，对于表格中绿色标记的线程，它的块ID是2，线程ID是1，如果想将这两个数字进行ID化，得到每个线程唯一的总ID，可以用块的ID乘以块中的线程总数，然后加上线程在这个块中的ID。数学表达式如下：
这次的main函数，和我们上次写过的那个非常类似。唯一的不同点在于内核的启动方式。现在我们用512个块，每个块里面有512个线程启动该内核。这样N非常大的问题就得到了解决。此外，我们不再将很长的结果数组中的每个值都打印出来，只打印结果是否正确。
2、存储器架构
在GPU上的代码执行被划分为流多处理器、块和线程。GPU有几个不同的存储器空间，每个存储器空间都有特定的特征和用途以及不同的速度和范围。这个存储空间按层次结构划分为不同的组块，比如全局内存、共享内存、本地内存、常量内存和纹理内存，每个组块都可以从程序中的不同点访问。此存储器架构如图所示：
如图所示，每个线程都有自己的本地存储器和寄存器堆。与处理器不同的是，GPU核心有很多寄存器来存储本地数据。当线程使用的数据不适合存储在寄存器堆中或者寄存器堆中装不下的时候，将会使用本地内存。寄存器堆和本地内存对每个线程都是唯一的。寄存器堆是最快的一种存储器。同一个块中的线程具有可由该块中的所有线程访问的共享内存。全局内存可被所有的块和其中的所有线程访问。它具有相当大的访问延迟，但存在缓存这种东西来给它提速。如下表，GPU有一级和二级缓存（即L1缓存和L2缓存）。常量内存则是用于存储常量和内核参数之类的只读数据。最后，存在纹理内存，这种内存可以利用各种2D和3D的访问模式。
所有存储器特征总结如下。
上表表述了各种存储器的各种特性。作用范围栏定义了程序的哪个部分能使用该存储器。而生存期定义了该存储器中的数据对程序可见的时间。除此之外，L1和L2缓存也可以用于GPU程序以便更快地访问存储器。
总之，所有线程都有一个寄存器堆，它是最快的。共享内存只能被块中的线程访问，但比全局内存块。全局内存是最慢的，但可以被所有的块访问。常量和纹理内存用于特殊用途。存储器访问是程序快速执行的最大瓶颈。
2.1、全局内存
所有的块都可以对全局内存进行读写。该存储器较慢，但是可以从你的代码的任何地方进行读写。缓存可加速对全局内存的访问。所有通过cudaMalloc分配的存储器都是全局内存。下面的简单代码演示了如何从程序中使用全局内存：
#include <stdio.h> #define N 5 __global__ void gpu_global_memory(int *d_a) { // "array" is a pointer into global memory on the device d_a[threadIdx.x] = threadIdx.x; } int main(int argc, char **argv) { // Define Host Array int h_a[N]; //Define device pointer int *d_a; cudaMalloc((void **)&d_a, sizeof(int) *N); // now copy data from host memory to device memory cudaMemcpy((void *)d_a, (void *)h_a, sizeof(int) *N, cudaMemcpyHostToDevice); // launch the kernel gpu_global_memory << <1, N >> > (d_a); // copy the modified array back to the host memory cudaMemcpy((void *)h_a, (void *)d_a, sizeof(int) *N, cudaMemcpyDeviceToHost); printf("Array in Global Memory is: \n"); //Printing result on console for (int i = 0; i < N; i++) { printf("At Index: %d --> %d \n", i, h_a[i]); } return 0; }
这段代码演示了如何从设备代码中进行全局内存的写入，以及如何从主机代码中用cudaMalloc进行分配，如何将指向该段全局内存的指针作为参数传递给内核函数。内核函数用不同的线程ID的值来填充这段全局内存。然后（用cudaMemcpy）复制到内存以便显示内容。最终结果如图所示：
2.2、本地内存和寄存器堆
本地内存和寄存器堆对每个线程都是唯一的。寄存器是每个线程可用的最快存储器。当内核中使用的变量在寄存器堆中装不下的时候，将会使用本地内存存储它们，这叫寄存器溢出。
请注意使用本地内存有两种情况：
（1）、寄存器不够了
（2）、某些情况根本就不能放在寄存器中
例如对一个局部数组的下标进行不定索引的时候。基本上可以将本地内存看成是每个线程的唯一的全局内存部分。相比寄存器堆，本地内存要慢很多。虽然本地内存通过L1缓存和L2缓存进行了缓冲，但寄存器溢出可能会影响你的程序的性能。
下面演示一个简单的程序：
#include <stdio.h> #define N 5 __global__ void gpu_local_memory(int d_in) { int t_local; t_local = d_in * threadIdx.x; printf("Value of Local variable in current thread is: %d \n", t_local); } int main(int argc, char **argv) { printf("Use of Local Memory on GPU:\n"); gpu_local_memory << <1, N >> > (5); cudaDeviceSynchronize(); return 0; }
代码中的t_local变量是每个线程局部唯一的，将被存储在寄存器堆中。用这种变量计算的时候，计算速度将是最快速的。以上代码的输出如图所示：
2.3、高速缓冲存储器
在较新的GPU上，每个流多处理器都含有自己独立的L1缓存，以及GPU有L2缓存。L2缓存是被所有的GPU中的流多处理器都共有的。所有的全局内存访问和本地内存访问都使用这些缓存，因为L1缓存在流多处理器内部独有，接近线程执行所需要的硬件单位，所以它的速度非常快。一般来说，L1缓存和共享内存共用同样的存储硬件，一共是64KB（注意：这是和计算能力有关，不一定共用相同的存储硬件，也不一定可以配置互相占用的比例，例如计算能力5.X和6.X的GPU卡就不能。同时L1缓存和共享内存在这两个计算能力上也不是共用的，但旧的计算能力和7.X GPU卡是如此），你可以配置L1缓存和共享内存分别在这64KB中的比例。所有的全局内存访问通过L2缓存进行。纹理内存和常量内存也分别有它们独立的缓存。
