随着多核CPU和众核GPU的到来，并行编程已经得到了业界越来越多的重视，CPU-GPU异构程序能够极大提高现有计算机系统的运算性能，对于科学计算等运算密集型程序有着非常重要的意义。这一系列文章是根据《CUDA C语言编程指南》来整理的，该指南是NVIDIA公司提供的CUDA学习资料，介绍了CUDA编程最基本最核心的概念，是学习CUDA必不可少的阅读材料。

初学CUDA，笔记错误之处在所难免，还请

随着多核CPU和众核GPU的到来，并行编程已经得到了业界越来越多的重视，CPU-GPU异构程序能够极大提高现有计算机系统的运算性能，对于科学计算等运算密集型程序有着非常重要的意义。这一系列文章是根据《CUDA C语言编程指南》来整理的，该指南是NVIDIA公司提供的CUDA学习资料，介绍了CUDA编程最基本最核心的概念，是学习CUDA必不可少的阅读材料。

初学CUDA，笔记错误之处在所难免，还请发现问题的诸位读者不吝赐教。

1. 什么是CUDA？

2. CUDA编程模型如何扩展？

3. CUDA基本概念

3.1 内核（Kernels）

CUDA C是C语言的一个扩展，它允许程序员定义一种被称为内核函数（Kernel Functions）的C函数，内核函数运行在GPU上，一旦启动，CUDA中的每一个线程都将会同时并行地执行内核函数中的代码。

内核函数使用关键字__global__来声明，运行该函数的CUDA线程数则通过<<<...>>>执行配置语法来设置。（参见章节"C语言扩展"），每一个执行内核函数的线程都由一个唯一的线程ID，这一ID可以通过在内核函数中访问threadIdx变量来得到。

下面通过一些示例代码来展示刚刚提到的这些概念该如何应用在编程中：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void VecAdd(float* A, float* B, float* C) {
  3. int i = threadIdx.x;
  4. C[i] = A[i] + B[i];
  5. }
  6. int main() {
  7. ...
  8. // Kernel invocation with N threads
  9. VecAdd<<<1, N>>>(A, B, C);
  10. ...
  11. }

在上面的代码中，N个线程将会并行地同时执行加法运算。

3.2 线程层次（Thread Hierarchy）

下面的例子展示了两个NxN矩阵相加的CUDA实现：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
  3. int i = threadIdx.x;
  4. int j = threadIdx.y;
  5. C[i][j] = A[i][j] + B[i][j];
  6. }
  7. int main() {
  8. ...
  9. // Kernel invocation with one block of N * N * 1 threads
  10. int numBlocks = 1;
  11. dim3 threadsPerBlock(N, N);
  12. MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
  13. ...
  14. }
每个线程块（block）中的线程数量是有限制的，因为依据前面所说，同一线程块（block）中的所有线程都会被分配到同一个处理器核上运行，共享有限的存储资源，因此对于当前的GPU，线程块所能包含的最大线程数目为1024。
上面的例子中numBlocks代表线程块的数量，这里的值为1。在一般的CUDA程序中，这个值通常大于1，也就是说将会有多个线程块被分配到多个处理器核中同时进行处理，这样就大大提高了程序的并行性。
在CUDA中，线程块包含在线程格（grid）当中，线程格可以是一维、二维或者三维的，线程格的尺寸一般根据待处理数据的规模或者处理器的数量来指定。线程格中所包含的线程块数目通常远远大于GPU处理器核心的数目。下图展示了线程格（grid）、线程块（block）以及线程（thread）之间的关系：
内核函数的调用可以简化为kernel<<<A，B>>>(parameters)，在尖括号中，A代表线程格（grid）的尺寸，它可以是三维的，用类型dim3表示，也可以是一维的，用int类型表示。B代表线程块（block）的尺寸，它与A类似，也可分别用dim3或int类型表示。
在内核函数内部，CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：
1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。
2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。
3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3。
4. threadIdx：线程索引，前面章节已经详细探讨过了，这里不再赘述。
明白了这些变量的含义，那么下面的矩阵加法程序便不难理解了：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
  3. int i = blockIdx.x * blockDim.x + threadIdx.x;
  4. int j = blockIdx.y * blockDim.y + threadIdx.y;
  5. if (i < N && j < N)
  6. C[i][j] = A[i][j] + B[i][j];
  7. }
  8. int main() {
  9. ...
  10. // Kernel invocation
  11. dim3 threadsPerBlock(16, 16);
  12. dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
  13. MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
  14. ...
  15. }
在上面的程序中，线程块（block）的尺寸是16x16，这是CUDA编程中一个非常普遍的选择。线程格（grid）包含了足够多的线程块（block）来进行计算。
线程块（block）是独立执行的，在执行的过程中线程块之间互不干扰，因此它们的执行顺序是随机的。
同一线程块中的线程可以通过访问共享内存（shared memory）或者通过同步函数__syncthreads()来协调合作。这些概念将在以后的章节中详细解释。

3.3 内存层次（Memory Hierarchy）

在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存（global memory）的区域对所有的CUDA线程都是可访问的。

除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。

一个程序启动内核函数以后，全局内存、常量内存以及纹理内存将会一直存在直到该程序结束。下面是CUDA的内存层次图：

3.4 异构编程（Heterogeneous Programming）

4. CUDA C语言编程接口

异构程序设计跟传统的串行程序设计差别是很大的，学习起来也是非常不容易的。NVIDIA非常够意思，为了简化CUDA的学习曲线，它采用了绝大多数程序员都熟悉的C语言作为其根基，CUDA C是NVIDIA为程序员提供的一类编程接口，它实际上是一个C语言的扩展，在C的基础上增加了一些新的语法和变量，并且提供了功能丰富的库函数，方便程序员使用GPU进行异构计算。
除了前面章节提到的CUDA最基本、最核心的概念以外，CUDA C呈现给程序员的接口主要由两大类API构成，它们分别是CUDA Runtime API和CUDA Driver API，Runtime API实际上是对于Driver API的封装，其目的自然是方便程序员的代码编写工作。Driver API为用户提供了更细一层的控制手段，通过它可以控制诸如CUDA Contexts（一种类似主机进程的概念）以及CUDA Modules（类似主机动态加载库的概念）等更加底层的CUDA模块。

4.1 NVCC编译器

任何一种程序设计语言都需要相应的编译器将其编译为二进制代码，进而在目标机器上得到执行。对于异构计算而言，这一过程与传统程序设计语言是有一些区别的。为什么？因为CUDA它本质上不是一种语言，而是一种异构计算的编程模型，使用CUDA C写出的代码需要在两种体系结构完全不同的设备上执行：1、CPU；2、GPU。因此，CUDA C的编译器所做的工作就有点略多了。一方面，它需要将源代码中运行在GPU端的代码编译得到能在CUDA设备上运行的二进制程序。另一方面，它也需要将源代码中运行在CPU端的程序编译得到能在主机CPU上运行的二进制程序。最后，它需要把这两部分有机地结合起来，使得两部分代码能够协调运行。

4.2 兼容性

      1、二进制兼容性
      二进制代码是设备相关的，使用NVCC编译器编译时，若指定-code选项，则会编译产生目标设备的二进制cubin对象。例如，编译时使用-code=sm_13会产生适用于计算能力1.3的二进制代码。二进制代码在CUDA计算设备上具有小版本的向前兼容性，但是在大版本上不具备兼容性。也就是说，对于计算能力X.y的硬件，使用-code=sm_Xy编译后，程序能够运行于计算能力X.z（其中z>=y）的硬件上，但不能运行在计算能力M.n（M!=X）的硬件上。
      2、PTX代码兼容性
      不同计算能力的设备所支持的PTX指令条数是不同的，一些PTX指令只在拥有较高计算能力的设备上被支持。例如，全局内存（global Memory）的原子操作指令只能用于计算能力不小于1.1的设备；双精度浮点运算指令只能用于计算能力不小于1.3的设备。在将C语言编译为PTX代码时，NVCC使用-arch编译选项指定PTX代码目标设备的计算能力。因此，要想使用双精度运算，编译时必须使用选项-arch=sm_13（或使用更高的计算能力），否则NVCC会自动将双精度操作降级为单精度操作。
      为某一特定设备产生的PTX代码，在运行时总是能够被具有更高计算能力的设备JIT编译为可执行的二进制代码。
      3、应用程序兼容性
      执行CUDA程序有两种方式，一种是直接加载编译好的CUDA二进制代码运行，另一种是首先加载程序中的PTX代码，再执行JIT编译得到二进制的设备可执行文件，然后运行。特别需要注意的是，为了让程序运行具有更高计算能力的未来设备上，必须让程序加载PTX代码。
      事实上，在一个CUDA C程序中可以嵌入不止一个版本的PTX/二进制代码。那么，具体执行时哪一个版本的PTX或者二进制代码会得到执行呢？答案是：最兼容的那个版本。例如编译一个名为x.cu的CUDA源代码：
将会产生兼容计算能力1.1硬件的二进制代码（第一排的-gencode选项）以及兼容计算能力1.1设备的PTX和二进制代码，这些代码都将会嵌入到编译后的目标文件中。
      主机端将会产生一些额外的代码，在程序运行时，这些代码会自动决定装载哪一个版本的代码来执行。对于上面的例子：
  • 计算能力1.0的设备运行该程序将会装载1.0版本的二进制代码
  • 计算能力1.1、1.2或者1.3的设备运行该程序将会装载1.1版本的二进制代码
  • 计算能力2.0或者更高的设备运行该程序将会装载1.1版本的PTX代码进而对其进行JIT编译得到相应设备的二进制代码
      同时，x.cu还可以在程序中使用一些特殊的宏来改变不同设备的代码执行路径。例如，对于计算能力1.1的设备而言，宏__CUDA_ARCH__等于110，在程序中可以对该宏的值进行判断，然后分支执行程序。
      NVCC用户手册列出了很多-arch，-code和-gencode等编译选项的简化书写形式。例如，-arch=sm_13就是-arch=compute_13 -code=compute13, sm_13的简化形式。更多详尽的内容请参阅该手册。
      4、C/C++兼容性
      NVCC编译器前端使用C++语法啊规则来处理CUDA源文件。在主机端，CUDA支持完整的C++语法；而在设备端，只有部分C++语法是被支持的。这方面更为详尽的讨论请参见《CUDA C程序设计指南》的C/C++语言支持章节。
      5、64位兼容性
      64位版本的nvcc编译器将设备代码编译为64位模式，即指针是64位的。运行64位设备代码的先决条件是主机端代码必须也使用64位模式进行编译。同样，32位版本的nvcc将设备代码编译为32位模式，这些代码也必须与相应的32位主机端代码相配合方能运行。
      32位nvcc编译器可以使用-m64编译选项将设备代码编译为64位模式。同时64位nvcc编译器也可使用-m32编译选项将设备代码编译为32位模式。

4.3 CUDA C Runtime

        CUDA C Runtime使用cudart动态链接库实现（cudart.dll或者cudart.so），运行时中所有的入口函数都以cuda为前缀。

4.3.1 初始化

4.3.2 设备内存

正如前面异构计算章节所讲，CUDA编程模型假定系统是由主机和设备构成的，它们分别具有自己独立的内存空间。Runtime负责设备内存的分配，回收，拷贝以及在主机和设备间传输数据的工作。

设备内存可以有两种分配方式：线性内存或者CUDA数组

CUDA数组是一块不透明的内存空间，它主要被优化用于纹理存取。

线性内存空间与平时我们访问的内存类似，对于计算能力1.x的设备来说，它存在于一个32位的地址空间。对于更高计算能力的设备而言，它存在于一个40位的地址空间中。因此，单独分配的实体可以使用指针来相互应用。

我们通常使用cudaMalloc()函数分配线性内存空间，使用cudaFree()函数释放线性内存空间，使用cudaMemcpy()函数在主机和设备之间传输数据。下面是CUDA Vector Add代码示例的一些片段：

[cpp] view plaincopy
  1. // Device code
  2. __global__ void VecAdd(float *A, float *B, float *C, int N) {
  3. int i = blockDim.x * blockIdx.x + threadIdx.x;
  4. if (i < N)
  5. C[i] = A[i] + B[i];
  6. }
  7. // Host code
  8. int main() {
  9. int N = ...;
  10. size_t size = N * sizeof(float);
  11. // Allocate input vectors h_A and h_B in host memory
  12. float *h_A = (float*)malloc(size);
  13. float *h_B = (float*)malloc(size);
  14. // Initialize input vectors
  15. ...
  16. // Allocate vectors in device memory
  17. float *d_A, *d_B, *d_C;
  18. cudaMalloc(&d_A, size);
  19. cudaMalloc(&d_B, size);
  20. cudaMalloc(&d_C, size);
  21. // Copy vectors from host memory to device memory
  22. cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
  23. cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
  24. // Invoke kernel
  25. int threadsPerBlock = 256;
  26. int blocksPerGrid = (N +threadsPerBlock - 1) / threadsPerBlock;
  27. VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
  28. // Copy result from device memory to host Memory
  29. cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
  30. // Free device memory
  31. cudaFree(d_A);
  32. cudaFree(d_B);
  33. cudaFree(d_C);
  34. // Free host memory
  35. ...
  36. }

片段展示了设备内存的分配，传输以及回收过程。

除了上面展示的方法，我们还可以使用cudaMallocPitch()和cudaMalloc3D()函数来分配线性内存。这些函数能够确保分配的内存满足设备内存访问的对齐要求，对于行地址的访问以及多维数组间的数据传输提供高性能保证，因此非常适合对于二维和三维数组内存空间的分配。下面的代码片段展示了分配和使用尺寸为width x height的二维数组的技术：

[cpp] view plaincopy
  1. // Host code
  2. int width = 64, height = 64;
  3. float *devPtr;
  4. size_t pitch;
  5. cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);
  6. MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
  7. // Device code
  8. __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) {
  9. for (int r = 0; r < height; ++r) {
  10. float* row = (float*)((char*)devPtr + r * pitch);
  11. for (int c = 0; c < width; ++c) {
  12. float element = row[c];
  13. }
  14. }
  15. }

下面的代码片段展示了一个尺寸为width x height x depth的三维数组的分配和使用方法：

[cpp] view plaincopy
  1. // Host code
  2. int width = 64, height = 64, depth = 64;
  3. cudaExtent extent = make_cudaExtent(width * sizeof(float), height, depth);
  4. cudaPitchedPtr devPitchedPtr;
  5. cudaMalloc3D(&devPitchedPtr, extent);
  6. MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);
  7. // Device code
  8. __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, int width, int height, int depth) {
  9. char* devPtr = devPitchedPtr.ptr;
  10. size_t pitch = devPitchedPtr.pitch;
  11. size_t slicePitch = pitch * height;
  12. for (int z = 0; z < depth; ++z) {
  13. char* slice = devPtr + z * slicePitch;
  14. for (int y = 0; y < height; ++y) {
  15. float* row = (float*)(slice + y * pitch);
  16. for (int x = 0; x < width; ++x)
  17. float element = row[x];
  18. }
  19. }
  20. }

更多详细的内容请查阅参考手册。

下面的代码示例展示了多种使用Runtime API访问全局变量的技术：

[cpp] view plaincopy
  1. __constant__ float constData[256];
  2. float data[256];
  3. cudaMemcpyToSymbol(constData, data, sizeof(data));
  4. cudaMemcpyFromSymbol(data, constData, sizeof(data));
  5. __device__ float devData;
  6. float value = 3.14f;
  7. cudaMemcpyToSymbol(devData, &value, sizeof(float));
  8. __device__ float* devPointer;
  9. float* ptr;
  10. cudaMalloc(&ptr, 256 * sizeof(float));
  11. cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

使用cudaGetSymbolAddress()函数可以获得被声明存储在全局内存中的变量地址。为了获得分配内存的大小，可以使用cudaGetSymbolSize()函数。

4.3 CUDA C Runtime

4.3.3 共享内存（Shared Memory）

下面是矩阵乘法的CUDA C主要实现代码：

[cpp] view plaincopy
  1. // Matrices are stored in row-major order:
  2. // M(row, col) = *(M.elements + row * M.width + col)
  3. typedef struct {
  4. int width;
  5. int height;
  6. float *elements;
  7. } Matrix;
  8. // Thread block size
  9. #define BLOCK_SIZE 16
  10. // Forward declaration of the matrix multiplication kernel
  11. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
  12. // Matrix multiplication - Host code
  13. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
  14. void MatMul(const Matrix A, const Matrix B, Matrix C) {
  15. // Load A and B to device memory
  16. Matrix d_A;
  17. d_A.width = A.width; d_A.height = A.height;
  18. size_t size = A.width * A.height * sizeof(float);
  19. cudaMalloc(&d_A.elements, size);
  20. cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
  21. Matrix d_B;
  22. d_B.width = B.width; d_B.height = B.height;
  23. size = B.width * B.height * sizeof(float);
  24. cudaMalloc(&d_B.elements, size);
  25. cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
  26. // Allocate C in device memory
  27. Matrix d_C;
  28. d_C.width = C.width; d_C.height = C.height;
  29. size = C.width * C.height * sizeof(float);
  30. cudaMalloc(&d_C.elements, size);
  31. // Invoke kernel
  32. dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  33. dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
  34. MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
  35. // Read C from device memory
  36. cudaMemcpy(C.elements, d_c.elements, size, cudaMemcpyDeviceToHost);
  37. // Free device memory
  38. cudaFree(d_A.elements);
  39. cudaFree(d_B.elements);
  40. cudaFree(d_C.elements);
  41. }
  42. // Matrix multiplication kernel called by MatMul()
  43. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
  44. // Each thread computes one element of C
  45. // by accumulating results into Cvalue
  46. float Cvalue = 0;
  47. int row  = blockIdx.y * blockDim.y + threadIdx.y;
  48. int col = blockIdx.x * blockDim.x + threadIdx.xl
  49. for (int e = 0; e < A.width; ++e)
  50. Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
  51. C.elements[row * C.width + col] = Cvalue;
  52. }
可以看出，为了计算矩阵C的任何一个元素，程序都需要从全局内存（global memory）中获得矩阵A的一行和矩阵B的一列。因此，完成这一计算矩阵A被读取了B.width次，矩阵B被读取了A.height次。
现在我们来使用共享内存（shared memory）实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub，我们使用一个线程块（thread block）来负责某一子方阵的计算，线程块中的每一个线程（thread）正好负责子方阵Csub中一个元素的计算。这样划分后，任何一个结果子方阵Csub'（尺寸为block_size * block_size）都是与该方阵具有相同行索引的尺寸为A.width * block_size的A的子矩阵Asub和与该方阵具有相同列索引的尺寸为block_size * B.height的B的子矩阵Bsub相乘所得到。
        为了匹配设备的计算资源，两个子矩阵Asub和Bsub被划分为尽可能多的分离的维度为block_size的子方阵，Csub的值便是这些子矩阵相乘后相加所得到的结果。子矩阵乘法的执行顺序都是首先将它们从全局内存（global memory）拷贝到共享内存（shared memory）（线程块中的每一个线程正好负责方阵一个元素的拷贝），然后由线程自己完成相应元素的计算任务，利用寄存器存储局部结果，最后将寄存器的内容与新得到的计算结果依此累加起来得到最终运算结果并将其传输到全局内存（global memory）中。
        通过使用这种分治的计算策略，共享内存得到了很好的利用，采用这种方案计算完成时全局内存中矩阵A被访问的次数为B.width / block_size，矩阵B被访问的次数为A.height / block_size，很明显，这为我们节省了非常多的全局内存带宽。优化后的矩阵计算示意图如下所示：

为了提升计算效率，我们为类型Matrix增加了一个成员变量stride。__device__函数用来获得和设置子矩阵的元素。下面是优化后的代码：

[cpp] view plaincopy
  1. // Matrices are stored in row-major order;
  2. // M(row, col) = *(M.elements + row * M.stride + col)
  3. typedef struct {
  4. int width;
  5. int height;
  6. int stride;
  7. float* elements;
  8. } Matrix;
  9. // Get a matrix element
  10. __device__ float GetElement(const Matrix A, int row, int col) {
  11. return A.elements[row * A.stride + col];
  12. }
  13. // Set a matrix element
  14. __device__ void SetElement(Matrix A, int row, int col, float value) {
  15. A.elements[row * A.stride + col] = value;
  16. }
  17. // Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
  18. // located col sub-matrices to the right and row sub-matrices down
  19. // from the upper-left corner of A
  20. __device__ Matrix GetSubMatrix(Matrix A, int row, int col) {
  21. Matrix Asub;
  22. Asub.width = BLOCK_SIZE;
  23. Asub.height = BLOCK_SIZE;
  24. Asub.stride = A.stride;
  25. Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];
  26. return Asub;
  27. }
  28. // Thread block size
  29. #define BLOCK_SIZE 16
  30. // Forward declaration of the matrix multiplication kernel
  31. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
  32. // Matrix multiplication - Host code
  33. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
  34. void MatMul(const Matrix A, const Matrix B, Matrix C) {
  35. // Load A and B to device memory
  36. Matrix d_A;
  37. d_A.width = d_A.stride = A.width;
  38. d_A.height = A.height;
  39. size_t size = A.width * A.height * sizeof(float);
  40. cudaMalloc(&d_A.elements, size);
  41. cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
  42. Matrix d_B;
  43. d_B.width = d_B.stride = B.width;
  44. d_B.height = B.height;
  45. size = B.width * B.height * sizeof(float);
  46. cudaMalloc(&d_B.elements, size);
  47. cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
  48. // Allocate C in device memory
  49. Matrix d_C;
  50. d_C.width = d_C.stride = C.width;
  51. d_C.height = C.height;
  52. size = C.width * C.height * sizeof(float);
  53. cudaMalloc(&d_C.elements, size);
  54. // Invoke kernel
  55. dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  56. dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
  57. MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
  58. // Read C from device memory
  59. cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);
  60. // Free device memory
  61. cudaFree(d_A.elements);
  62. cudaFree(d_B.elements);
  63. cudaFree(d_C.elements);
  64. }
  65. // Matrix multiplication kernel called by MatMul()
  66. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
  67. // Block row and column
  68. int blockRow = blockIdx.y;
  69. int blockCol = blockIdx.x;
  70. // Each thread block computes one sub-matrix Csub of C
  71. Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
  72. // Each thread computes one element of Csub
  73. // by accumulating results into Cvalue
  74. float Cvalue = 0;
  75. // Thread row and column within Csub
  76. int row = threadIdx.y;
  77. int col = threadIdx.x;
  78. // Look over all the sub-matrices of A and B that are required to compute Csub
  79. // Multiply each pair of sub-matrices together and accumulate the results
  80. for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
  81. // Get sub-matrix Asub of A
  82. Matrix Asub = GetSubMatrix(A, blockRow, m);
  83. // Get sub-matrix Bsub of B
  84. Matrix Bsub = GetSubMatrix(B, m, blockCol);
  85. // Shared memory used to store Asub and Bsub respectively
  86. __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  87. __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
  88. // Load Asub and Bsub from device memory to shared memory
  89. // Each thread loads one element of each sub-matrix
  90. As[row][col] = GetElement(Asub, row, col);
  91. Bs[row][col] = GetElement(Bsub, row, col);
  92. // Synchronize to make sure the sub-matrices are loaded
  93. // before starting the computation
  94. __syncthreads();
  95. // Multiply Asub and Bsub together
  96. for (int e = 0; e < BLOCK_SIZE; ++e)
  97. Cvalue += As[row][e] * Bs[e][col];
  98. // Synchronize to make sure that the preceding computation is done before
  99. // loading two new sub-matrices of A and B in the next iteration
  100. __syncthreads();
  101. }
  102. // Write Csub to device memory
  103. // Each thread writes one element
  104. SetElement(Csub, row, col, Cvalue);
  105. }

异步并行执行

主机和设备间并行执行

将内核启动与数据传输重叠起来

对于一些计算能力等于或高于1.1的设备，它们可以将内核启动任务和锁页内存到设备内存的数据传输任务并行执行。应用程序可以检查设备属性中的asyncEngineCount项来确定设备是否支持这一功能。当该项值大于0时代表设备支持这一层次的并行。对于计算能力1.x的设备，该功能不支持通过cudaMallocPitch()函数分配的CUDA数组或2D数组。

并行内核执行

一些计算能力2.x或更高的设备可以同时并行执行多个内核函数。应用程序可以检查设备属性中的concurrentKernels项来确定设备是否支持这一功能，值为1代表支持。运算能力3.5的设备在同一时刻能够并行执行的最大内核函数数量为32，运算能力小于3.5的硬件则最多支持同时启动16个内核函数的执行。同时需要注意的是，在一个CUDA上下文中的内核函数不能与另一个CUDA上下文中的内核函数同时执行。使用很多纹理内存或者大量本地内存的内核函数也很可能无法与其它内核函数并行执行。

并行数据传输

一些计算能力为2.x或更高的设备可以将锁页内存到设备内存的数据传输和设备内存到锁页内存的数据传输并行执行。应用程序可检查设备属性中的asyncEngineCount项来确定这一功能的支持程度，等于2时表示支持。

流（Streams）

应用程序通过流来管理并行。一个流是一个顺次执行的命令序列。不同的流之间并行执行，没有固定的执行顺序。

1、流的创建与销毁

定义一个流的过程通常包括：创建一个流对象，然后指定它为内核启动或者主机设备间数据传输的流参数。下面的一段代码创建了两个流并且在锁页内存中分配了一块float类型的数组hostPtr：

[cpp] view plaincopy
  1. cudaStream_t stream[2];
  2. for (int i = 0; i < 2; ++i)
  3. cudaStreamCreate(&stream[i]);
  4. float *hostPtr;
  5. cudaMallocHost(&hostPtr, 2 * size);

下面的代码定义了每一个流的行为：从主机端拷贝数据到设备端，内核启动，从设备端拷贝数据到主机端：

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i) {
  2. cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
  3. MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
  4. cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
  5. }

这部分代码中有一点需要注意：为了并行化数据拷贝和内核执行，主机端内存必须分配为锁页（page-locked）内存。

要销毁一个流需要调用函数cudaStreamDestroy()

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i)
  2. cudaStreamDestroy(stream[i]);

cudaStreamDestroy() 函数等待之前流中的指令序列运行完成，然后销毁指定流，将控制权返还给主机端。

2、默认流（Default stream）

在内核启动或者数据拷贝过程中如果不指定流，或者设置流参数为0，则相应的指令将会运行在默认流上，它们也因此而顺次执行。

3、明同步（Explicit Synchronization）

在CUDA中有很多种方式可以用来同步流的执行：

cudaDeviceSynchronize()函数使得主机端线程阻塞直到所有流中的指令执行完成。

cudaStreamSynchronize()函数将一个流对象作为输入参数，用以等待指定流中的所有指令执行完成。

cudaStreamWaitEvent()函数将一个流对象和一个事件作为输入参数，它将延迟该函数调用后在指定流中所有新加入的命令的执行直到指定的事件完成为止。流参数可以为0，在该情形下所有流中的任何新加入的指令都必须等待指定事件的发生，然后才可以执行。

cudaStreamQuery()函数为应用程序提供了一个检测指定流中之前指令是否执行完成的方法。

为了避免同步带来的性能下降，所有上述同步函数最好用于计时目的或者分离错误的内核执行或数据拷贝。

4、暗同步（Implicit Synchronization）

如果任何一个流中正在执行以下操作，那么其它流是不能与其并行运行的：

a. 分配锁页内存空间

b. 设备内存分配

c. 设备内存置位

d. 同一设备两个不同地址间正在进行数据拷贝

e. 默认流中有指令正在执行

f. L1/shared内存配置的转换

对于支持并行内核执行并且计算能力3.0或以下的设备来说，任何一个需要检查依赖性以确定流内核启动是否完成的操作：

a. 只有当前CUDA上下文中所有流中所有之前的内核启动之后才能够启动执行。

b. 将会阻塞所有当前CUDA上下文中的任意流中新加入的内核调用直到内核检查完成。

需要进行依赖性检查的操作包括执行检查的内核启动所在流中的其它指令以及任何在该流上对cudaStreamQuery()函数的调用。因此，应用程序可以遵照以下指导原则来提升潜在并行性：

（1）所有非依赖操作应当比依赖性操作提前进行

（2）任何类型的同步越迟越好

5、重叠行为（Overlapping Behavior）

两个流间重叠行为的数量取决于以下几个因素：

（1）每个流中命令发出的次序

（2）设备是否支持内核启动与数据传输并行

（3）设备是否支持多内核并行启动

（4）设备是否支持多数据传输并行

例如，在不支持并行数据传输的设备上，“流的创建与销毁”章节中代码样例中的操作就不能并行，因为在stream[0]中发出设备端到主机端的数据拷贝后，stream[1]又发出主机端到设备端的数据拷贝命令，这两个命令式不能重叠执行的。假设设备支持数据传输与内核启动并行，那么如下代码：

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i)
  2. cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
  3. for (int i = 0; i < 2; ++i)
  4. MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
  5. for (int i = 0; i < 2; ++i)
  6. cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);

可将stream[0]的内核启动和stream[1]从主机端到设备端的数据拷贝重叠起来并行执行。

6、回调函数

CUDA运行时提供了cudaStreamAddCallback()函数以在流中的任意位置插入一个回调函数点。回调函数运行于主机端，如果在默认流中插入回调函数，那么它将等待所有其它流中的命令执行完成之后才会开始执行。

下面的代码展示了回调函数技术的应用：

[cpp] view plaincopy
  1. void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void **data) {
  2. printf("Inside callback %d\n", (int)data);
  3. }
  4. ...
  5. for (int i = 0; i < 2; ++i) {
  6. cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
  7. MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
  8. cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
  9. cudaStreamAddCallback(stream[i], MyCallback, (void**)i, 0);
  10. }

上面的代码定义了两个流的操作，每个流都完成一次主机端到设备端的数据拷贝，一次内核启动，一次设备端到主机端的数据拷贝，最后增加了一个加入回调函数的操作。当设备端代码运行到回调函数点的时候，设备将控制权交还给主机端，主机端运行完成以后再将控制权返还给设备端，然后设备端继续运行。

值得注意的是，在一个回调函数中，一定不能进行任何CUDA API的调用，直接的或者间接的都是不可以的。

发现问题的诸位读者不吝赐教。

1. 什么是CUDA？

2. CUDA编程模型如何扩展？

3. CUDA基本概念

3.1 内核（Kernels）

CUDA C是C语言的一个扩展，它允许程序员定义一种被称为内核函数（Kernel Functions）的C函数，内核函数运行在GPU上，一旦启动，CUDA中的每一个线程都将会同时并行地执行内核函数中的代码。

内核函数使用关键字__global__来声明，运行该函数的CUDA线程数则通过<<<...>>>执行配置语法来设置。（参见章节"C语言扩展"），每一个执行内核函数的线程都由一个唯一的线程ID，这一ID可以通过在内核函数中访问threadIdx变量来得到。

下面通过一些示例代码来展示刚刚提到的这些概念该如何应用在编程中：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void VecAdd(float* A, float* B, float* C) {
  3. int i = threadIdx.x;
  4. C[i] = A[i] + B[i];
  5. }
  6. int main() {
  7. ...
  8. // Kernel invocation with N threads
  9. VecAdd<<<1, N>>>(A, B, C);
  10. ...
  11. }

在上面的代码中，N个线程将会并行地同时执行加法运算。

3.2 线程层次（Thread Hierarchy）

下面的例子展示了两个NxN矩阵相加的CUDA实现：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
  3. int i = threadIdx.x;
  4. int j = threadIdx.y;
  5. C[i][j] = A[i][j] + B[i][j];
  6. }
  7. int main() {
  8. ...
  9. // Kernel invocation with one block of N * N * 1 threads
  10. int numBlocks = 1;
  11. dim3 threadsPerBlock(N, N);
  12. MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
  13. ...
  14. }
每个线程块（block）中的线程数量是有限制的，因为依据前面所说，同一线程块（block）中的所有线程都会被分配到同一个处理器核上运行，共享有限的存储资源，因此对于当前的GPU，线程块所能包含的最大线程数目为1024。
上面的例子中numBlocks代表线程块的数量，这里的值为1。在一般的CUDA程序中，这个值通常大于1，也就是说将会有多个线程块被分配到多个处理器核中同时进行处理，这样就大大提高了程序的并行性。
在CUDA中，线程块包含在线程格（grid）当中，线程格可以是一维、二维或者三维的，线程格的尺寸一般根据待处理数据的规模或者处理器的数量来指定。线程格中所包含的线程块数目通常远远大于GPU处理器核心的数目。下图展示了线程格（grid）、线程块（block）以及线程（thread）之间的关系：
内核函数的调用可以简化为kernel<<<A，B>>>(parameters)，在尖括号中，A代表线程格（grid）的尺寸，它可以是三维的，用类型dim3表示，也可以是一维的，用int类型表示。B代表线程块（block）的尺寸，它与A类似，也可分别用dim3或int类型表示。
在内核函数内部，CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：
1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。
2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。
3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3。
4. threadIdx：线程索引，前面章节已经详细探讨过了，这里不再赘述。
明白了这些变量的含义，那么下面的矩阵加法程序便不难理解了：

[cpp] view plaincopy
  1. // Kernel definition
  2. __global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
  3. int i = blockIdx.x * blockDim.x + threadIdx.x;
  4. int j = blockIdx.y * blockDim.y + threadIdx.y;
  5. if (i < N && j < N)
  6. C[i][j] = A[i][j] + B[i][j];
  7. }
  8. int main() {
  9. ...
  10. // Kernel invocation
  11. dim3 threadsPerBlock(16, 16);
  12. dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
  13. MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
  14. ...
  15. }
在上面的程序中，线程块（block）的尺寸是16x16，这是CUDA编程中一个非常普遍的选择。线程格（grid）包含了足够多的线程块（block）来进行计算。
线程块（block）是独立执行的，在执行的过程中线程块之间互不干扰，因此它们的执行顺序是随机的。
同一线程块中的线程可以通过访问共享内存（shared memory）或者通过同步函数__syncthreads()来协调合作。这些概念将在以后的章节中详细解释。

3.3 内存层次（Memory Hierarchy）

在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存（global memory）的区域对所有的CUDA线程都是可访问的。

除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。

一个程序启动内核函数以后，全局内存、常量内存以及纹理内存将会一直存在直到该程序结束。下面是CUDA的内存层次图：

3.4 异构编程（Heterogeneous Programming）

4. CUDA C语言编程接口

异构程序设计跟传统的串行程序设计差别是很大的，学习起来也是非常不容易的。NVIDIA非常够意思，为了简化CUDA的学习曲线，它采用了绝大多数程序员都熟悉的C语言作为其根基，CUDA C是NVIDIA为程序员提供的一类编程接口，它实际上是一个C语言的扩展，在C的基础上增加了一些新的语法和变量，并且提供了功能丰富的库函数，方便程序员使用GPU进行异构计算。
除了前面章节提到的CUDA最基本、最核心的概念以外，CUDA C呈现给程序员的接口主要由两大类API构成，它们分别是CUDA Runtime API和CUDA Driver API，Runtime API实际上是对于Driver API的封装，其目的自然是方便程序员的代码编写工作。Driver API为用户提供了更细一层的控制手段，通过它可以控制诸如CUDA Contexts（一种类似主机进程的概念）以及CUDA Modules（类似主机动态加载库的概念）等更加底层的CUDA模块。

4.1 NVCC编译器

任何一种程序设计语言都需要相应的编译器将其编译为二进制代码，进而在目标机器上得到执行。对于异构计算而言，这一过程与传统程序设计语言是有一些区别的。为什么？因为CUDA它本质上不是一种语言，而是一种异构计算的编程模型，使用CUDA C写出的代码需要在两种体系结构完全不同的设备上执行：1、CPU；2、GPU。因此，CUDA C的编译器所做的工作就有点略多了。一方面，它需要将源代码中运行在GPU端的代码编译得到能在CUDA设备上运行的二进制程序。另一方面，它也需要将源代码中运行在CPU端的程序编译得到能在主机CPU上运行的二进制程序。最后，它需要把这两部分有机地结合起来，使得两部分代码能够协调运行。

4.2 兼容性

      1、二进制兼容性
      二进制代码是设备相关的，使用NVCC编译器编译时，若指定-code选项，则会编译产生目标设备的二进制cubin对象。例如，编译时使用-code=sm_13会产生适用于计算能力1.3的二进制代码。二进制代码在CUDA计算设备上具有小版本的向前兼容性，但是在大版本上不具备兼容性。也就是说，对于计算能力X.y的硬件，使用-code=sm_Xy编译后，程序能够运行于计算能力X.z（其中z>=y）的硬件上，但不能运行在计算能力M.n（M!=X）的硬件上。
      2、PTX代码兼容性
      不同计算能力的设备所支持的PTX指令条数是不同的，一些PTX指令只在拥有较高计算能力的设备上被支持。例如，全局内存（global Memory）的原子操作指令只能用于计算能力不小于1.1的设备；双精度浮点运算指令只能用于计算能力不小于1.3的设备。在将C语言编译为PTX代码时，NVCC使用-arch编译选项指定PTX代码目标设备的计算能力。因此，要想使用双精度运算，编译时必须使用选项-arch=sm_13（或使用更高的计算能力），否则NVCC会自动将双精度操作降级为单精度操作。
      为某一特定设备产生的PTX代码，在运行时总是能够被具有更高计算能力的设备JIT编译为可执行的二进制代码。
      3、应用程序兼容性
      执行CUDA程序有两种方式，一种是直接加载编译好的CUDA二进制代码运行，另一种是首先加载程序中的PTX代码，再执行JIT编译得到二进制的设备可执行文件，然后运行。特别需要注意的是，为了让程序运行具有更高计算能力的未来设备上，必须让程序加载PTX代码。
      事实上，在一个CUDA C程序中可以嵌入不止一个版本的PTX/二进制代码。那么，具体执行时哪一个版本的PTX或者二进制代码会得到执行呢？答案是：最兼容的那个版本。例如编译一个名为x.cu的CUDA源代码：
将会产生兼容计算能力1.1硬件的二进制代码（第一排的-gencode选项）以及兼容计算能力1.1设备的PTX和二进制代码，这些代码都将会嵌入到编译后的目标文件中。
      主机端将会产生一些额外的代码，在程序运行时，这些代码会自动决定装载哪一个版本的代码来执行。对于上面的例子：
  • 计算能力1.0的设备运行该程序将会装载1.0版本的二进制代码
  • 计算能力1.1、1.2或者1.3的设备运行该程序将会装载1.1版本的二进制代码
  • 计算能力2.0或者更高的设备运行该程序将会装载1.1版本的PTX代码进而对其进行JIT编译得到相应设备的二进制代码
      同时，x.cu还可以在程序中使用一些特殊的宏来改变不同设备的代码执行路径。例如，对于计算能力1.1的设备而言，宏__CUDA_ARCH__等于110，在程序中可以对该宏的值进行判断，然后分支执行程序。
      NVCC用户手册列出了很多-arch，-code和-gencode等编译选项的简化书写形式。例如，-arch=sm_13就是-arch=compute_13 -code=compute13, sm_13的简化形式。更多详尽的内容请参阅该手册。
      4、C/C++兼容性
      NVCC编译器前端使用C++语法啊规则来处理CUDA源文件。在主机端，CUDA支持完整的C++语法；而在设备端，只有部分C++语法是被支持的。这方面更为详尽的讨论请参见《CUDA C程序设计指南》的C/C++语言支持章节。
      5、64位兼容性
      64位版本的nvcc编译器将设备代码编译为64位模式，即指针是64位的。运行64位设备代码的先决条件是主机端代码必须也使用64位模式进行编译。同样，32位版本的nvcc将设备代码编译为32位模式，这些代码也必须与相应的32位主机端代码相配合方能运行。
      32位nvcc编译器可以使用-m64编译选项将设备代码编译为64位模式。同时64位nvcc编译器也可使用-m32编译选项将设备代码编译为32位模式。

4.3 CUDA C Runtime

        CUDA C Runtime使用cudart动态链接库实现（cudart.dll或者cudart.so），运行时中所有的入口函数都以cuda为前缀。

4.3.1 初始化

4.3.2 设备内存

正如前面异构计算章节所讲，CUDA编程模型假定系统是由主机和设备构成的，它们分别具有自己独立的内存空间。Runtime负责设备内存的分配，回收，拷贝以及在主机和设备间传输数据的工作。

设备内存可以有两种分配方式：线性内存或者CUDA数组

CUDA数组是一块不透明的内存空间，它主要被优化用于纹理存取。

线性内存空间与平时我们访问的内存类似，对于计算能力1.x的设备来说，它存在于一个32位的地址空间。对于更高计算能力的设备而言，它存在于一个40位的地址空间中。因此，单独分配的实体可以使用指针来相互应用。

我们通常使用cudaMalloc()函数分配线性内存空间，使用cudaFree()函数释放线性内存空间，使用cudaMemcpy()函数在主机和设备之间传输数据。下面是CUDA Vector Add代码示例的一些片段：

[cpp] view plaincopy
  1. // Device code
  2. __global__ void VecAdd(float *A, float *B, float *C, int N) {
  3. int i = blockDim.x * blockIdx.x + threadIdx.x;
  4. if (i < N)
  5. C[i] = A[i] + B[i];
  6. }
  7. // Host code
  8. int main() {
  9. int N = ...;
  10. size_t size = N * sizeof(float);
  11. // Allocate input vectors h_A and h_B in host memory
  12. float *h_A = (float*)malloc(size);
  13. float *h_B = (float*)malloc(size);
  14. // Initialize input vectors
  15. ...
  16. // Allocate vectors in device memory
  17. float *d_A, *d_B, *d_C;
  18. cudaMalloc(&d_A, size);
  19. cudaMalloc(&d_B, size);
  20. cudaMalloc(&d_C, size);
  21. // Copy vectors from host memory to device memory
  22. cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
  23. cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
  24. // Invoke kernel
  25. int threadsPerBlock = 256;
  26. int blocksPerGrid = (N +threadsPerBlock - 1) / threadsPerBlock;
  27. VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
  28. // Copy result from device memory to host Memory
  29. cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
  30. // Free device memory
  31. cudaFree(d_A);
  32. cudaFree(d_B);
  33. cudaFree(d_C);
  34. // Free host memory
  35. ...
  36. }

片段展示了设备内存的分配，传输以及回收过程。

除了上面展示的方法，我们还可以使用cudaMallocPitch()和cudaMalloc3D()函数来分配线性内存。这些函数能够确保分配的内存满足设备内存访问的对齐要求，对于行地址的访问以及多维数组间的数据传输提供高性能保证，因此非常适合对于二维和三维数组内存空间的分配。下面的代码片段展示了分配和使用尺寸为width x height的二维数组的技术：

[cpp] view plaincopy
  1. // Host code
  2. int width = 64, height = 64;
  3. float *devPtr;
  4. size_t pitch;
  5. cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);
  6. MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
  7. // Device code
  8. __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) {
  9. for (int r = 0; r < height; ++r) {
  10. float* row = (float*)((char*)devPtr + r * pitch);
  11. for (int c = 0; c < width; ++c) {
  12. float element = row[c];
  13. }
  14. }
  15. }

下面的代码片段展示了一个尺寸为width x height x depth的三维数组的分配和使用方法：

[cpp] view plaincopy
  1. // Host code
  2. int width = 64, height = 64, depth = 64;
  3. cudaExtent extent = make_cudaExtent(width * sizeof(float), height, depth);
  4. cudaPitchedPtr devPitchedPtr;
  5. cudaMalloc3D(&devPitchedPtr, extent);
  6. MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);
  7. // Device code
  8. __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, int width, int height, int depth) {
  9. char* devPtr = devPitchedPtr.ptr;
  10. size_t pitch = devPitchedPtr.pitch;
  11. size_t slicePitch = pitch * height;
  12. for (int z = 0; z < depth; ++z) {
  13. char* slice = devPtr + z * slicePitch;
  14. for (int y = 0; y < height; ++y) {
  15. float* row = (float*)(slice + y * pitch);
  16. for (int x = 0; x < width; ++x)
  17. float element = row[x];
  18. }
  19. }
  20. }

更多详细的内容请查阅参考手册。

下面的代码示例展示了多种使用Runtime API访问全局变量的技术：

[cpp] view plaincopy
  1. __constant__ float constData[256];
  2. float data[256];
  3. cudaMemcpyToSymbol(constData, data, sizeof(data));
  4. cudaMemcpyFromSymbol(data, constData, sizeof(data));
  5. __device__ float devData;
  6. float value = 3.14f;
  7. cudaMemcpyToSymbol(devData, &value, sizeof(float));
  8. __device__ float* devPointer;
  9. float* ptr;
  10. cudaMalloc(&ptr, 256 * sizeof(float));
  11. cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

使用cudaGetSymbolAddress()函数可以获得被声明存储在全局内存中的变量地址。为了获得分配内存的大小，可以使用cudaGetSymbolSize()函数。

4.3 CUDA C Runtime

4.3.3 共享内存（Shared Memory）

下面是矩阵乘法的CUDA C主要实现代码：

[cpp] view plaincopy
  1. // Matrices are stored in row-major order:
  2. // M(row, col) = *(M.elements + row * M.width + col)
  3. typedef struct {
  4. int width;
  5. int height;
  6. float *elements;
  7. } Matrix;
  8. // Thread block size
  9. #define BLOCK_SIZE 16
  10. // Forward declaration of the matrix multiplication kernel
  11. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
  12. // Matrix multiplication - Host code
  13. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
  14. void MatMul(const Matrix A, const Matrix B, Matrix C) {
  15. // Load A and B to device memory
  16. Matrix d_A;
  17. d_A.width = A.width; d_A.height = A.height;
  18. size_t size = A.width * A.height * sizeof(float);
  19. cudaMalloc(&d_A.elements, size);
  20. cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
  21. Matrix d_B;
  22. d_B.width = B.width; d_B.height = B.height;
  23. size = B.width * B.height * sizeof(float);
  24. cudaMalloc(&d_B.elements, size);
  25. cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
  26. // Allocate C in device memory
  27. Matrix d_C;
  28. d_C.width = C.width; d_C.height = C.height;
  29. size = C.width * C.height * sizeof(float);
  30. cudaMalloc(&d_C.elements, size);
  31. // Invoke kernel
  32. dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  33. dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
  34. MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
  35. // Read C from device memory
  36. cudaMemcpy(C.elements, d_c.elements, size, cudaMemcpyDeviceToHost);
  37. // Free device memory
  38. cudaFree(d_A.elements);
  39. cudaFree(d_B.elements);
  40. cudaFree(d_C.elements);
  41. }
  42. // Matrix multiplication kernel called by MatMul()
  43. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
  44. // Each thread computes one element of C
  45. // by accumulating results into Cvalue
  46. float Cvalue = 0;
  47. int row  = blockIdx.y * blockDim.y + threadIdx.y;
  48. int col = blockIdx.x * blockDim.x + threadIdx.xl
  49. for (int e = 0; e < A.width; ++e)
  50. Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
  51. C.elements[row * C.width + col] = Cvalue;
  52. }
可以看出，为了计算矩阵C的任何一个元素，程序都需要从全局内存（global memory）中获得矩阵A的一行和矩阵B的一列。因此，完成这一计算矩阵A被读取了B.width次，矩阵B被读取了A.height次。
现在我们来使用共享内存（shared memory）实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub，我们使用一个线程块（thread block）来负责某一子方阵的计算，线程块中的每一个线程（thread）正好负责子方阵Csub中一个元素的计算。这样划分后，任何一个结果子方阵Csub'（尺寸为block_size * block_size）都是与该方阵具有相同行索引的尺寸为A.width * block_size的A的子矩阵Asub和与该方阵具有相同列索引的尺寸为block_size * B.height的B的子矩阵Bsub相乘所得到。
        为了匹配设备的计算资源，两个子矩阵Asub和Bsub被划分为尽可能多的分离的维度为block_size的子方阵，Csub的值便是这些子矩阵相乘后相加所得到的结果。子矩阵乘法的执行顺序都是首先将它们从全局内存（global memory）拷贝到共享内存（shared memory）（线程块中的每一个线程正好负责方阵一个元素的拷贝），然后由线程自己完成相应元素的计算任务，利用寄存器存储局部结果，最后将寄存器的内容与新得到的计算结果依此累加起来得到最终运算结果并将其传输到全局内存（global memory）中。
        通过使用这种分治的计算策略，共享内存得到了很好的利用，采用这种方案计算完成时全局内存中矩阵A被访问的次数为B.width / block_size，矩阵B被访问的次数为A.height / block_size，很明显，这为我们节省了非常多的全局内存带宽。优化后的矩阵计算示意图如下所示：

为了提升计算效率，我们为类型Matrix增加了一个成员变量stride。__device__函数用来获得和设置子矩阵的元素。下面是优化后的代码：

[cpp] view plaincopy
  1. // Matrices are stored in row-major order;
  2. // M(row, col) = *(M.elements + row * M.stride + col)
  3. typedef struct {
  4. int width;
  5. int height;
  6. int stride;
  7. float* elements;
  8. } Matrix;
  9. // Get a matrix element
  10. __device__ float GetElement(const Matrix A, int row, int col) {
  11. return A.elements[row * A.stride + col];
  12. }
  13. // Set a matrix element
  14. __device__ void SetElement(Matrix A, int row, int col, float value) {
  15. A.elements[row * A.stride + col] = value;
  16. }
  17. // Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
  18. // located col sub-matrices to the right and row sub-matrices down
  19. // from the upper-left corner of A
  20. __device__ Matrix GetSubMatrix(Matrix A, int row, int col) {
  21. Matrix Asub;
  22. Asub.width = BLOCK_SIZE;
  23. Asub.height = BLOCK_SIZE;
  24. Asub.stride = A.stride;
  25. Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];
  26. return Asub;
  27. }
  28. // Thread block size
  29. #define BLOCK_SIZE 16
  30. // Forward declaration of the matrix multiplication kernel
  31. __global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
  32. // Matrix multiplication - Host code
  33. // Matrix dimensions are assumed to be multiples of BLOCK_SIZE
  34. void MatMul(const Matrix A, const Matrix B, Matrix C) {
  35. // Load A and B to device memory
  36. Matrix d_A;
  37. d_A.width = d_A.stride = A.width;
  38. d_A.height = A.height;
  39. size_t size = A.width * A.height * sizeof(float);
  40. cudaMalloc(&d_A.elements, size);
  41. cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice);
  42. Matrix d_B;
  43. d_B.width = d_B.stride = B.width;
  44. d_B.height = B.height;
  45. size = B.width * B.height * sizeof(float);
  46. cudaMalloc(&d_B.elements, size);
  47. cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice);
  48. // Allocate C in device memory
  49. Matrix d_C;
  50. d_C.width = d_C.stride = C.width;
  51. d_C.height = C.height;
  52. size = C.width * C.height * sizeof(float);
  53. cudaMalloc(&d_C.elements, size);
  54. // Invoke kernel
  55. dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  56. dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
  57. MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
  58. // Read C from device memory
  59. cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost);
  60. // Free device memory
  61. cudaFree(d_A.elements);
  62. cudaFree(d_B.elements);
  63. cudaFree(d_C.elements);
  64. }
  65. // Matrix multiplication kernel called by MatMul()
  66. __global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) {
  67. // Block row and column
  68. int blockRow = blockIdx.y;
  69. int blockCol = blockIdx.x;
  70. // Each thread block computes one sub-matrix Csub of C
  71. Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
  72. // Each thread computes one element of Csub
  73. // by accumulating results into Cvalue
  74. float Cvalue = 0;
  75. // Thread row and column within Csub
  76. int row = threadIdx.y;
  77. int col = threadIdx.x;
  78. // Look over all the sub-matrices of A and B that are required to compute Csub
  79. // Multiply each pair of sub-matrices together and accumulate the results
  80. for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
  81. // Get sub-matrix Asub of A
  82. Matrix Asub = GetSubMatrix(A, blockRow, m);
  83. // Get sub-matrix Bsub of B
  84. Matrix Bsub = GetSubMatrix(B, m, blockCol);
  85. // Shared memory used to store Asub and Bsub respectively
  86. __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  87. __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
  88. // Load Asub and Bsub from device memory to shared memory
  89. // Each thread loads one element of each sub-matrix
  90. As[row][col] = GetElement(Asub, row, col);
  91. Bs[row][col] = GetElement(Bsub, row, col);
  92. // Synchronize to make sure the sub-matrices are loaded
  93. // before starting the computation
  94. __syncthreads();
  95. // Multiply Asub and Bsub together
  96. for (int e = 0; e < BLOCK_SIZE; ++e)
  97. Cvalue += As[row][e] * Bs[e][col];
  98. // Synchronize to make sure that the preceding computation is done before
  99. // loading two new sub-matrices of A and B in the next iteration
  100. __syncthreads();
  101. }
  102. // Write Csub to device memory
  103. // Each thread writes one element
  104. SetElement(Csub, row, col, Cvalue);
  105. }

异步并行执行

主机和设备间并行执行

将内核启动与数据传输重叠起来

对于一些计算能力等于或高于1.1的设备，它们可以将内核启动任务和锁页内存到设备内存的数据传输任务并行执行。应用程序可以检查设备属性中的asyncEngineCount项来确定设备是否支持这一功能。当该项值大于0时代表设备支持这一层次的并行。对于计算能力1.x的设备，该功能不支持通过cudaMallocPitch()函数分配的CUDA数组或2D数组。

并行内核执行

一些计算能力2.x或更高的设备可以同时并行执行多个内核函数。应用程序可以检查设备属性中的concurrentKernels项来确定设备是否支持这一功能，值为1代表支持。运算能力3.5的设备在同一时刻能够并行执行的最大内核函数数量为32，运算能力小于3.5的硬件则最多支持同时启动16个内核函数的执行。同时需要注意的是，在一个CUDA上下文中的内核函数不能与另一个CUDA上下文中的内核函数同时执行。使用很多纹理内存或者大量本地内存的内核函数也很可能无法与其它内核函数并行执行。

并行数据传输

一些计算能力为2.x或更高的设备可以将锁页内存到设备内存的数据传输和设备内存到锁页内存的数据传输并行执行。应用程序可检查设备属性中的asyncEngineCount项来确定这一功能的支持程度，等于2时表示支持。

流（Streams）

应用程序通过流来管理并行。一个流是一个顺次执行的命令序列。不同的流之间并行执行，没有固定的执行顺序。

1、流的创建与销毁

定义一个流的过程通常包括：创建一个流对象，然后指定它为内核启动或者主机设备间数据传输的流参数。下面的一段代码创建了两个流并且在锁页内存中分配了一块float类型的数组hostPtr：

[cpp] view plaincopy
  1. cudaStream_t stream[2];
  2. for (int i = 0; i < 2; ++i)
  3. cudaStreamCreate(&stream[i]);
  4. float *hostPtr;
  5. cudaMallocHost(&hostPtr, 2 * size);

下面的代码定义了每一个流的行为：从主机端拷贝数据到设备端，内核启动，从设备端拷贝数据到主机端：

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i) {
  2. cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
  3. MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
  4. cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
  5. }

这部分代码中有一点需要注意：为了并行化数据拷贝和内核执行，主机端内存必须分配为锁页（page-locked）内存。

要销毁一个流需要调用函数cudaStreamDestroy()

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i)
  2. cudaStreamDestroy(stream[i]);

cudaStreamDestroy() 函数等待之前流中的指令序列运行完成，然后销毁指定流，将控制权返还给主机端。

2、默认流（Default stream）

在内核启动或者数据拷贝过程中如果不指定流，或者设置流参数为0，则相应的指令将会运行在默认流上，它们也因此而顺次执行。

3、明同步（Explicit Synchronization）

在CUDA中有很多种方式可以用来同步流的执行：

cudaDeviceSynchronize()函数使得主机端线程阻塞直到所有流中的指令执行完成。

cudaStreamSynchronize()函数将一个流对象作为输入参数，用以等待指定流中的所有指令执行完成。

cudaStreamWaitEvent()函数将一个流对象和一个事件作为输入参数，它将延迟该函数调用后在指定流中所有新加入的命令的执行直到指定的事件完成为止。流参数可以为0，在该情形下所有流中的任何新加入的指令都必须等待指定事件的发生，然后才可以执行。

cudaStreamQuery()函数为应用程序提供了一个检测指定流中之前指令是否执行完成的方法。

为了避免同步带来的性能下降，所有上述同步函数最好用于计时目的或者分离错误的内核执行或数据拷贝。

4、暗同步（Implicit Synchronization）

如果任何一个流中正在执行以下操作，那么其它流是不能与其并行运行的：

a. 分配锁页内存空间

b. 设备内存分配

c. 设备内存置位

d. 同一设备两个不同地址间正在进行数据拷贝

e. 默认流中有指令正在执行

f. L1/shared内存配置的转换

对于支持并行内核执行并且计算能力3.0或以下的设备来说，任何一个需要检查依赖性以确定流内核启动是否完成的操作：

a. 只有当前CUDA上下文中所有流中所有之前的内核启动之后才能够启动执行。

b. 将会阻塞所有当前CUDA上下文中的任意流中新加入的内核调用直到内核检查完成。

需要进行依赖性检查的操作包括执行检查的内核启动所在流中的其它指令以及任何在该流上对cudaStreamQuery()函数的调用。因此，应用程序可以遵照以下指导原则来提升潜在并行性：

（1）所有非依赖操作应当比依赖性操作提前进行

（2）任何类型的同步越迟越好

5、重叠行为（Overlapping Behavior）

两个流间重叠行为的数量取决于以下几个因素：

（1）每个流中命令发出的次序

（2）设备是否支持内核启动与数据传输并行

（3）设备是否支持多内核并行启动

（4）设备是否支持多数据传输并行

例如，在不支持并行数据传输的设备上，“流的创建与销毁”章节中代码样例中的操作就不能并行，因为在stream[0]中发出设备端到主机端的数据拷贝后，stream[1]又发出主机端到设备端的数据拷贝命令，这两个命令式不能重叠执行的。假设设备支持数据传输与内核启动并行，那么如下代码：

[cpp] view plaincopy
  1. for (int i = 0; i < 2; ++i)
  2. cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
  3. for (int i = 0; i < 2; ++i)
  4. MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
  5. for (int i = 0; i < 2; ++i)
  6. cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);

可将stream[0]的内核启动和stream[1]从主机端到设备端的数据拷贝重叠起来并行执行。

6、回调函数

CUDA运行时提供了cudaStreamAddCallback()函数以在流中的任意位置插入一个回调函数点。回调函数运行于主机端，如果在默认流中插入回调函数，那么它将等待所有其它流中的命令执行完成之后才会开始执行。

下面的代码展示了回调函数技术的应用：

[cpp] view plaincopy
  1. void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void **data) {
  2. printf("Inside callback %d\n", (int)data);
  3. }
  4. ...
  5. for (int i = 0; i < 2; ++i) {
  6. cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
  7. MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
  8. cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
  9. cudaStreamAddCallback(stream[i], MyCallback, (void**)i, 0);
  10. }

上面的代码定义了两个流的操作，每个流都完成一次主机端到设备端的数据拷贝，一次内核启动，一次设备端到主机端的数据拷贝，最后增加了一个加入回调函数的操作。当设备端代码运行到回调函数点的时候，设备将控制权交还给主机端，主机端运行完成以后再将控制权返还给设备端，然后设备端继续运行。

值得注意的是，在一个回调函数中，一定不能进行任何CUDA API的调用，直接的或者间接的都是不可以的。

CUDA编程指南阅读笔记相关推荐

  1. CUDA编程指南阅读笔记（六）

    4. CUDA C语言编程接口 接上文. 4.3 CUDA C Runtime 4.3.3 共享内存(Shared Memory)         共享内存是CUDA设备中非常重要的一个存储区域,有效 ...

  2. Java并发编程艺术阅读笔记（一）

    Java并发编程艺术阅读笔记(一) 1.什么是上下文切换 CPU通过时间片分配算法循环执行任务,执行完一个任务的时间片后就会切换到下一个任务.但是在切换任务之前会保存上一个任务的状态,在切换回该任务, ...

  3. CUDA PTX ISA阅读笔记（一）

    不知道这是个啥的看这里:Parallel Thread Execution ISA Version 5.0. 简要来说,PTX就是.cu代码编译出来的一种东西,然后再由PTX编译生成执行代码.如果不想 ...

  4. iPhone编程指南学习笔记

    为什么80%的码农都做不了架构师?>>>    UIWindow 在创建应用程序窗口时,您应该总是将其初始的边框尺寸设置为整个屏幕的大小.如果您的窗口是从nib文件装载得到,Inte ...

  5. Google C++编程风格指南阅读笔记之命名、注释和格式

    文章目录 前言 命名约定 类型命名 变量命名 枚举命名 宏的命名 注释 注释风格 文件注释 类注释 函数注释 变量注释 类的数据成员 全局变量 实现注释 TODO注释 格式 行长度 空格还是制表符 函 ...

  6. 程序员健康指南阅读笔记

    编程需要程序员全神贯注,这就常常导致我们忽略生活的其他方面,其中最被忽略的就是,健康状况. 职业不应该让你犯职业病,只要用对了方法,它就不会造成伤害. 做出改变 努力健康起来吧! 永别了椅子 灵活的饮 ...

  7. python最佳实践指南试题_Python最佳实践指南 阅读笔记

    创建将0到19连接起来的字符串1 2 3 4 5 6 7 8nums = [] for n in range(20): nums.append(str(n)) print "".j ...

  8. [转载]高质量c/c++编程指南读书笔记

    一个strcpy函数的代码 能考查三个方面 (1) 编程风格 (2) 出错处理 (3) 算法复杂度分析(用于提供性能) 定义编程老手和编程高手 定义1:能长期稳定地编写出高质量程序的程序员称为编程老手 ...

  9. 编程修养 阅读笔记四

    转载:http://blog.csdn.net/haoel/article/details/2872 26.为常量声明宏 -------- 最好不要在程序中出现数字式的"硬编码"( ...

最新文章

  1. 漫画：据说很多搞软件的羡慕硬件工程师
  2. 国家微生物科学数据中心微生物组学数据汇交指南
  3. 微生物组数据库（http://egcloud.cib.cn）正式上线
  4. 行业观察 | 机器人Ameca挣脱「灵魂」枷锁觉醒？
  5. linux 报错 E: 无法定位软件包 python-lzma
  6. Redis - Redis command timed out nested exception is io.lettuce.core.RedisCommandTimeoutException
  7. 从 0 到 1，看我玩弄千万日志于股掌
  8. 谷歌紧急修复已遭在野利用的Chrome 0day
  9. 一个bootstrap.css的使用案例
  10. 以mysql为例的数据字典_建立数据字典
  11. 超火的快闪文字视频制作，用这个软件就能轻松学会
  12. DNA甲基化芯片探针的P值如何计算
  13. egret的WebView实现（基于egret2.5）
  14. c语言程序设计需要学多久,九江c语言编程学习,九江学c语言编程报班,九江学c语言编程一般要多久才能学会...
  15. 《大学物理》课程考试大纲
  16. [elixir! #0015][译] 学习 GenStage by Joseph Kain
  17. 现货、期货、期权、权证
  18. boost库之geometry
  19. python五子棋程序教程_Python 五子棋 编程
  20. SpingBoot—微服务初始化资源方法

热门文章

  1. 多线程中使用UNITY变量导致线程执行断掉却又不报错的问题
  2. Redis过期策略及实现原理
  3. ES6中的Promise使用方法与总结
  4. Zabbix SQL注入漏洞威胁预警通告
  5. Web认证及API的 使用TOKEN的一些思考
  6. ORACLE+RAC+ASM环境下添加redo日志组
  7. java_spring_依赖注入（构造器）
  8. N900超频和刷9G-home一气呵成
  9. Windows 7防火墙设置详解(三)
  10. numpy中newaxis的用法