CUDA学习笔记中，如何高效掌握的优化技巧？

CUDA [!NOTE] 本笔记推荐使用Typora的Mdmdt主题渲染 本笔记旨在简单记录一些《CUDA by Example. An Introduction to General-Purpose GPU Programming》书籍上的重要概念。 目录CUDA核函数核函数的声明和调用核函数的内部变量核函数内置变量的使用1D2D设备内存共享内存和同步常量内存和事件事件 核函数 我们通常把CPU和系统的内存称为主机（host），把GPU以及其内存称为设备（device） 在设备上执行的函数称为核函数（kernal） 核函数的声明和调用 标识符__global__告知编译器，该函数应该被编译成运行在设备上而不是主机上 标识符__device__告知编译器，该函数运行在设备上，而且仅能由__global__或__device__标识的函数所调用 CUDA的编译器和运行时会处理好从主机上调用核函数的一些琐碎事务，包括将核函数参数传递给设备 语法<<<Blocks, ThreadsPerBlock>>>告知CUDA运行时如何启动该核函数；第一个参数的含义是我们想让设备分配多少并行块来运行该核函数；第二个参数表示我们希望CUDA运行时为我们创建的每个块的线程数。 上面提到的，运行在GPU上的并行块的集合称为网格（grid） [!NOTE] 分配并行块的数量和每个块的线程数量都有限制，若超出限制可能会导致核函数调用失败，因此任务的计算量最好不能与这两个参数有关联，最好的情况是任务的计算量限制仅与设备内存相关 针对上面的问题常见的做法是提前固定好三重括号（<<<_, _>>>）内参数的数量，在核函数内进行多轮循环以覆盖全部任务。（实际上这种做法可能也存在一定的坏处，后面会介绍） 核函数的内部变量 我们前面已经解释了三重括号的作用，我们利用GPU的原因就是其强大的并行处理能力，使用三重括号可以告知CUDA运行时以什么方式并行启动我们的核函数，此时会有多个线程（这些线程共享同一份核函数代码，执行相同的代码逻辑）同时执行，这里有一个问题，既然所有线程都执行相同的逻辑，那么不就是在做重复的工作吗？ 为了解决这个问题，CUDA运行时为我们提供了可以在核函数内使用的一些变量，用来标识当前是执行的是哪个线程。也就是说我们在核函数内要做的就是利用这个线程标识来准确的分配我们的计算任务给每一个线程。 线程索引（threadIdx）和块索引（blockIdx）分别用来标识当前线程位于块中的第几个，以及位于网格的哪个块中 块维度（blockDim）和网格维度（girdDim）分别用来表示块的线程数量好网格的块数量 上面的三个变量的数据类型都是dim3，也就是有三个维度用来适配在1、2、3维做计算的计算任务 向三重括号传递时既可以直接传递整数（运行时自动识别为1维的配置），或者传入像这样初始化的变量（推测是用了宏处理）🚩 dim3 grid(dimX, dimY); 核函数内置变量的使用 1D 在一维的情况下（也就是直接向三重括号传递整数），Grid的情况可以理解为下图： 因此该线程的索引便是： int idx = blockIdx.x * blockDim.x + threadIdx.x; 2D 2维的情况下（也就是向三重括号传递上面的dim3 ），Grid如下所示： 其中方括号内的分别代表着blockIdx.x、blockIdx.y、threadIdx.x和threadIdx.y，由图gridDim为(3, 2)，blockDim为(16, 16)。