点击上方AI人工智能初学者，订阅我！此刻开始我们一起学习进步！
目录
1、CUDA线程
2、存储器架构
2.1、全局内存
2.2、本地内存和寄存器堆
2.3、告诉缓冲存储器
3、线程同步
3.1、共享内存
3.2、原子操作
笔记来源书籍推荐
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#include #include //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 kernelint 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 arraysint h_a[N], h_b[N], h_c[N];//Defining device pointersint *d_a, *d_b, *d_c;// allocate the memorycudaMalloc((void**)&d_a, N * sizeof(int));cudaMalloc((void**)&d_b, N * sizeof(int));cudaMalloc((void**)&d_c, N * sizeof(int));//Initializing Arraysfor (int i = 0; i < N; i++) {h_a[i] = 2 * i*i;h_b[i] = i;}// Copy input arrays from host to device memorycudaMemcpy(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 parametersgpuAdd << <512, 512 >> >(d_a, d_b, d_c);//Copy result back to host memory from device memorycudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);cudaDeviceSynchronize();int Correct = 1;printf("Vector addition on GPU \n");//Printing result on consolefor (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 memorycudaFree(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 #define N 5__global__ void gpu_global_memory(int *d_a){// "array" is a pointer into global memory on the deviced_a[threadIdx.x] = threadIdx.x;}int main(int argc, char **argv){// Define Host Arrayint h_a[N];//Define device pointerint *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 memorycudaMemcpy((void *)h_a, (void *)d_a, sizeof(int) *N, cudaMemcpyDeviceToHost);printf("Array in Global Memory is: \n");//Printing result on consolefor (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 #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缓存进行。纹理内存和常量内存也分别有它们独立的缓存。
3、线程同步
3.1、共享内存
共享内存位于芯片内部，因此它比全局内存快得多。(CUDA里面存储器的快慢有两方面，一个是延迟低，一个是带宽大。这里特指延迟低)，相比没有经过缓存的全局内存访问，共享内存大约在延迟上低100倍。同一个块中的线程可以访问相同的一段共享内存(注意：不同块中的线程所见到的共享内存中的内容是不相同的)，这在许多线程需要与其他线程共享它们的结果的应用程序中非常有用。但是如果不同步，也可能会造成混乱或错误的结果。如果某线程的计算结果在写入到共享内存完成之前被其他线程读取，那么将会导致错误。因此，应该正确地控制或管理内存访问。这是由__syncthreads()指令完成的，该指令确保在继续执行程序之前完成对内存的所有写入操作。这也被称为barrier。barrier的含义是块中的所有线程都将到达该代码行，然后在此等待其他线程完成。当所有线程都到达了这里之后，它们可以一起继续往下执行。
#include __global__ void gpu_shared_memory(float *d_a){// Defining local variables which are private to each threadint i, index = threadIdx.x;float average, sum = 0.0f;//Define shared memory__shared__ float sh_arr[10];sh_arr[index] = d_a[index];__syncthreads(); // This ensures all the writes to shared memory have completedfor (i = 0; i<= index; i++) { sum += sh_arr[i]; }average = sum / (index + 1.0f);d_a[index] = average; sh_arr[index] = average;}int main(int argc, char **argv){//Define Host Arrayfloat h_a[10]; //Define Device Pointerfloat *d_a; for (int i = 0; i < 10; i++) {h_a[i] = i;}// allocate global memory on the devicecudaMalloc((void **)&d_a, sizeof(float) * 10);// now copy data from host memory to device memory cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10, cudaMemcpyHostToDevice);gpu_shared_memory << <1, 10 >> >(d_a);// copy the modified array back to the host memorycudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);printf("Use of Shared Memory on GPU: \n");//Printing result on consolefor (int i = 0; i < 10; i++) {printf("The running average after %d element is %f \n", i, h_a[i]);}return 0;}
在main函数中，当分配好主机和设备上的数组后，用0.0到9.0填充主机上的数组，然后将这个数组复制到显存。内核将对显存中的数据进行读取，计算并保存结果。最后结果从显存中传输到内存，然后在控制台上输出。控制台上的输出结果如图所示：
这个程序还含有额外的一个CUDA函数调用：cudaDeviceSynchronize()。为何要加这句？这是因为启动内核是一个异步操作，只要发布了内核启动命令，不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。在上述的代码中，CPU线程返回，继续执行的下一句是printf()。而再之后，在内核完成之前，进程就会结束，终止控制台窗口。所以，如果不加上这句同步函数，你就看不到任何的内核执行结果输出。在程序退出后内核生成的输出结果，将没有地方可去，你没法看到它们，因此，如果我们不包含这个指令，你将不会看到任何内核执行的printf语句的输出结果。要能看到内核生成的输出结果，我们必须包含这句同步函数。这样，内核的结果将通过可用的标准输出显示，而应用程序则会在内核执行完成之后才退出。
3.2、原子操作
考虑当大量的线程需要试图修改一段较小的内存区域的情形，这是(在日常的算法实现中)常发生的现象。当我们试图进行“读取-修改-写入”操作序列的时候，这种情形经常会带来很多麻烦。
一个例子是代码d_out[i]++，这代码首先将d_out[i]的原值从存储器中读取出来，然后执行了+1操作，再将结果回写到存储器。然而，如果多个线程试图在同一个内存区域中进行这个操作，则可能会得到错误的结果。
假设某内存区域中有初始值6，两个线程p和q分别试图将这段区域中的内容+1，则最终的结果应当是8。但是在实际执行的时候，可能p和q两个线程同时读取了这个初始值，两者都得到了6，执行+1操作都得到了7，然后它们将7写回这个内存区域。这样，和正确的结果8不同，我们得到的最终结果是7，这是错误的。这种错误是如何的危险，我们通过ATM取现操作来演示。假设你的账户余额为5000卢比，你的账户下面开了两张银行卡，你和你的朋友同时去2个不同的ATM上取现4000卢比，你俩在同一瞬间刷卡取现。所以，当两个ATM检查余额的时候，都将显示5000卢比的余额。当你俩同时取现4000卢比的时候，两个ATM机都只根据初始值5000卢比判断，要取的现金4000卢比小于当前余额。所以两个机器将会给你们每人4000卢比。即使你之前只有5000卢比的余额，你们也能得到8000卢比，这很危险。为了示范一下这种情形，做了一个很多线程试图同时访问一个小数组的例子：
#include #define NUM_THREADS 10000#define SIZE 10#define BLOCK_WIDTH 100__global__ void gpu_increment_without_atomic(int *d_a){// Calculate thread id for current threadint tid = blockIdx.x * blockDim.x + threadIdx.x;// each thread increments elements wrapping at SIZE variabletid = tid % SIZE;d_a[tid] += 1;}int main(int argc, char **argv){printf("%d total threads in %d blocks writing into %d array elements\n",NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);// declare and allocate host memoryint h_a[SIZE];const int ARRAY_BYTES = SIZE * sizeof(int);// declare and allocate GPU memoryint * d_a;cudaMalloc((void **)&d_a, ARRAY_BYTES);//Initialize GPU memory to zerocudaMemset((void *)d_a, 0, ARRAY_BYTES);gpu_increment_without_atomic <> >(d_a);// copy back the array to host memorycudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);printf("Number of times a particular Array index has been incremented without atomic add is: \n");for (int i = 0; i < SIZE; i++){printf("index: %d --> %d times\n ", i, h_a[i]);}cudaFree(d_a);return 0;}
可能如同已经猜到的那样，每次运行你的程序，每个内存区域中的元素值都可能会不同。这是设备上不定顺序的多线程执行导致的。
为了解决这个问题，CUDA提供了atomicAdd这种原子操作函数。该函数会从逻辑上保证，每个调用它的线程对相同的内存区域上的“读取旧值-累加-回写新值”操作是不可被其他线程扰乱的原子性的整体完成的。使用atomicAdd进行原子累加的内核函数代码如下：
#include #define NUM_THREADS 10000#define SIZE 10#define BLOCK_WIDTH 100__global__ void gpu_increment_atomic(int *d_a){// Calculate thread id for current threadint tid = blockIdx.x * blockDim.x + threadIdx.x;// each thread increments elements wrapping at SIZE variabletid = tid % SIZE;atomicAdd(&d_a[tid], 1);}int main(int argc, char **argv){printf("%d total threads in %d blocks writing into %d array elements\n",NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);// declare and allocate host memoryint h_a[SIZE];const int ARRAY_BYTES = SIZE * sizeof(int);// declare and allocate GPU memoryint * d_a;cudaMalloc((void **)&d_a, ARRAY_BYTES);//Initialize GPU memory to zerocudaMemset((void *)d_a, 0, ARRAY_BYTES);gpu_increment_atomic << > >(d_a);// copy back the array to host memorycudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);printf("Number of times a particular Array index has been incremented is: \n");for (int i = 0; i < SIZE; i++) { printf("index: %d --> %d times\n ", i, h_a[i]); }cudaFree(d_a);return 0;}
在main函数中，具有10个元素的数组被初始化成0值，然后传递给了内核，但现在，内核中的代码将执行原子累加操作。所以，这个程序输出的结果将是对的，数组中的每个元素将被累加1000。运行结果显示如图：
如果你测量一下这个程序的运行时间，相比之前的那个简单地在全局内存上直接进行加法操作的程序它用的时间更长。这是因为使用原子操作后程序具有更大的执行代价。可以通过使用共享内存来加速这些原子累加操作。如果线程规模不变，但原子操作的元素数量扩大，则这些同样次数的原子操作会更快地完成。这是因为更广泛的分布范围上的原子操作有利于利用多个能执行原子操作的单元，以及每个原子操作单元上面的竞争性的原子事务也相应减少了。
笔记来源书籍：
《基于GPU加速的计算机视觉编程》——使用OpenCV和CUDA实时处理复杂图像数据
比较好的一本书，推荐给大家，希望可以帮助到你，后续小编也会持续摘录本书内容以及调试好代码并上传，并在最后以视觉项目的形式进行收尾。小编也是第一次接触CUDA编程，做笔记知识希望可以和大家一起交流学习。
关注【AI人工智能初学者】公众号，回复【CUDA4】建议长按复制，即可获得完整的项目代码文件。
希望您可以关注公众号，也非常期待您的打赏。
声明：转载请说明出处
下方为小生公众号，还望包容接纳和关注，非常期待与您的美好相遇，让我们以梦为马，砥砺前行。
希望技术与灵魂可以一路同行
长按识别二维码关注一下
更多精彩内容可回复关键词
每篇文章的主题即可
点“在看”给我一朵小黄花