• CUDA从入门到精通

    http://blog.csdn.net/augusdi/article/details/12833235

    CUDA从入门到精通（零）：写在前面

    在老板的要求下。本博主从2012年上高性能计算课程開始接触CUDA编程，随后将该技术应用到了实际项目中。使处理程序加速超过1K，可见基于图形显示器的并行计算对于追求速度的应用来说无疑是一个理想的选择。

    还有不到一年毕业。怕是毕业后这些技术也就随毕业而去，准备这个暑假开辟一个CUDA专栏，从入门到精通。步步为营。顺便分享设计的一些经验教训。希望能给学习CUDA的童鞋提供一定指导。

    个人能力所及，错误难免。欢迎讨论。

     

    PS：申请专栏好像须要先发原创帖超过15篇。

    。。算了，先写够再申请吧。到时候一并转过去。


    CUDA从入门到精通（一）：环境搭建

    NVIDIA于2006年推出CUDA（Compute Unified Devices Architecture），能够利用其推出的GPU进行通用计算。将并行计算从大型集群扩展到了普通显卡，使得用户仅仅须要一台带有Geforce显卡的笔记本就能跑较大规模的并行处理程序。

     

    使用显卡的优点是，和大型集群相比功耗非常低。成本也不高。但性能非常突出。以我的笔记本为例，Geforce 610M，用DeviceQuery程序測试。可得到例如以下硬件參数：

    计算能力达48X0.95 = 45.6 GFLOPS。而笔记本的CPU參数例如以下：

    CPU计算能力为（4核）：2.5G*4 = 10GFLOPS，可见，显卡计算性能是4核i5 CPU的4~5倍，因此我们能够充分利用这一资源来对一些耗时的应用进行加速。

     

    好了，工欲善其事必先利其器，为了使用CUDA对GPU进行编程，我们须要准备下面必备工具：

    1. 硬件平台，就是显卡，假设你用的不是NVIDIA的显卡。那么仅仅能说抱歉，其它都不支持CUDA。

    2. 操作系统，我用过windows XP，Windows 7都没问题。本博客用Windows7。

    3. C编译器，建议VS2008，和本博客一致。

    4. CUDA编译器NVCC。能够免费免注冊免license从官网下载CUDA ToolkitCUDA下载，最新版本号为5.0。本博客用的就是该版本号。

    5. 其它工具（如Visual Assist，辅助代码高亮）

     

    准备完成。開始安装软件。VS2008安装比較费时间，建议安装完整版（NVIDIA官网说Express版也能够）。过程不必详述。CUDA Toolkit 5.0里面包括了NVCC编译器、设计文档、设计例程、CUDA执行时库、CUDA头文件等必备的原材料。

    安装完成，我们在桌面上发现这个图标：

    不错，就是它。双击执行。能够看到一大堆例程。

    我们找到Simple OpenGL这个执行看看效果：

      点右边黄线标记处的Run就可以看到美妙的三维正弦曲面，鼠标左键拖动能够转换角度，右键拖动能够缩放。假设这个执行成功，说明你的环境基本搭建成功。

    出现故障的可能：

    1. 你使用远程桌面连接登录到还有一台server。该server上有显卡支持CUDA。但你远程终端不能执行CUDA程序。这是由于远程登录使用的是你本地显卡资源，在远程登录时看不到server端的显卡，所以会报错：没有支持CUDA的显卡！解决方法：1. 远程server装两块显卡。一块仅仅用于显示，还有一块用于计算。2.不要用图形界面登录，而是用命令行界面如telnet登录。

    2.有两个以上显卡都支持CUDA的情况。怎样区分是在哪个显卡上执行？这个须要你在程序里控制，选择符合一定条件的显卡。如较高的时钟频率、较大的显存、较高的计算版本号等。

    具体操作见后面的博客。

    好了。先说这么多。下一节我们介绍怎样在VS2008中给GPU编程。


    CUDA从入门到精通（二）：第一个CUDA程序

    书接上回，我们既然直接执行例程成功了。接下来就是了解怎样实现例程中的每一个环节。当然。我们先从简单的做起，一般编程语言都会找个helloworld样例，而我们的显卡是不会说话的，仅仅能做一些简单的加减乘除运算。所以，CUDA程序的helloworld。我想应该最合适只是的就是向量加了。

    打开VS2008，选择File->New->Project，弹出以下对话框，设置例如以下：

    之后点OK，直接进入project界面。

    project中。我们看到仅仅有一个.cu文件。内容例如以下：

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3.   
    4. #include <stdio.h>  
    5.   
    6. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    7.   
    8. __global__ void addKernel(int *c, const int *a, const int *b)  
    9. {  
    10.     int i = threadIdx.x;  
    11.     c[i] = a[i] + b[i];  
    12. }  
    13.   
    14. int main()  
    15. {  
    16.     const int arraySize = 5;  
    17.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    18.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    19.     int c[arraySize] = { 0 };  
    20.   
    21.     // Add vectors in parallel.  
    22.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);  
    23.     if (cudaStatus != cudaSuccess) {  
    24.         fprintf(stderr, "addWithCuda failed!");  
    25.         return 1;  
    26.     }  
    27.   
    28.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d} ",  
    29.         c[0], c[1], c[2], c[3], c[4]);  
    30.   
    31.     // cudaThreadExit must be called before exiting in order for profiling and  
    32.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    33.     cudaStatus = cudaThreadExit();  
    34.     if (cudaStatus != cudaSuccess) {  
    35.         fprintf(stderr, "cudaThreadExit failed!");  
    36.         return 1;  
    37.     }  
    38.   
    39.     return 0;  
    40. }  
    41.   
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess) {  
    53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    54.         goto Error;  
    55.     }  
    56.   
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess) {  
    60.         fprintf(stderr, "cudaMalloc failed!");  
    61.         goto Error;  
    62.     }  
    63.   
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess) {  
    66.         fprintf(stderr, "cudaMalloc failed!");  
    67.         goto Error;  
    68.     }  
    69.   
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess) {  
    72.         fprintf(stderr, "cudaMalloc failed!");  
    73.         goto Error;  
    74.     }  
    75.   
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess) {  
    79.         fprintf(stderr, "cudaMemcpy failed!");  
    80.         goto Error;  
    81.     }  
    82.   
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess) {  
    85.         fprintf(stderr, "cudaMemcpy failed!");  
    86.         goto Error;  
    87.     }  
    88.   
    89.     // Launch a kernel on the GPU with one thread for each element.  
    90.     addKernel<<<1, size>>>(dev_c, dev_a, dev_b);  
    91.   
    92.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    93.     // any errors encountered during the launch.  
    94.     cudaStatus = cudaThreadSynchronize();  
    95.     if (cudaStatus != cudaSuccess) {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.   
    100.     // Copy output vector from GPU buffer to host memory.  
    101.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    102.     if (cudaStatus != cudaSuccess) {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106.   
    107. Error:  
    108.     cudaFree(dev_c);  
    109.     cudaFree(dev_a);  
    110.     cudaFree(dev_b);  
    111.       
    112.     return cudaStatus;  
    113. }  
     能够看出，CUDA程序和C程序并无差别，仅仅是多了一些以"cuda"开头的一些库函数和一个特殊声明的函数：
    [cpp] view plaincopy
    1. __global__ void addKernel(int *c, const int *a, const int *b)  
    2. {  
    3.     int i = threadIdx.x;  
    4.     c[i] = a[i] + b[i];  
    5. }  

    这个函数就是在GPU上执行的函数，称之为核函数。英文名Kernel Function，注意要和操作系统内核函数区分开来。

    我们直接按F7编译。能够得到例如以下输出：

    [html] view plaincopy
    1. 1>------ Build started: Project: cuda_helloworld, Configuration: Debug Win32 ------    
    2. 1>Compiling with CUDA Build Rule...    
    3. 1>"C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0\bin vcc.exe"  -G   -gencode=arch=compute_10,code="sm_10,compute_10" -gencode=arch=compute_20,code="sm_20,compute_20"  --machine 32 -ccbin "C:Program Files (x86)Microsoft Visual Studio 9.0VCin"    -Xcompiler "/EHsc /W3 /nologo /O2 /Zi   /MT  "  -I"C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0\include" -maxrregcount=0   --compile -o "Debug/kernel.cu.obj" kernel.cu      
    4. 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu    
    5. 1>tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu    
    6. 1>tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu    
    7. 1>tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu    
    8. 1>kernel.cu    
    9. 1>kernel.cu    
    10. 1>tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp    
    11. 1>tmpxft_000000ec_00000000-24_kernel.compute_10.ii    
    12. 1>Linking...    
    13. 1>Embedding manifest...    
    14. 1>Performing Post-Build Event...    
    15. 1>copy "C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0\bincudart*.dll" "C:UsersDongXiaomanDocumentsVisual Studio 2008Projectscuda_helloworldDebug"    
    16. 1>C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0\bincudart32_50_35.dll    
    17. 1>C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0\bincudart64_50_35.dll    
    18. 1>已复制         2 个文件。    
    19. 1>Build log was saved at "file://c:UsersDongXiaomanDocumentsVisual Studio 2008Projectscuda_helloworldcuda_helloworldDebugBuildLog.htm"    
    20. 1>cuda_helloworld - 0 error(s), 105 warning(s)    
    21. ========== Build: 1 succeeded, 0 failed, 0 up-to-date, 0 skipped ==========    

    可见，编译.cu文件须要利用nvcc工具。该工具的具体使用见后面博客。

    直接执行，能够得到结果图例如以下：

    假设显示正确，那么我们的第一个程序宣告成功。

    刚入门CUDA，跑过几个官方提供的例程，看了看人家的代码。认为并不难。但自己动手写代码时，总是不知道要先干什么。后干什么，也不知道从哪个知识点学起。这时就须要有一本能提供指导的书籍或者教程，一步步跟着做下去。直到真正掌握。

    一般讲述CUDA的书，我觉得不错的有以下这几本：

    刚開始学习的人能够先看美国人写的这本《GPU高性能编程CUDA实战》。可操作性非常强，但不要期望能全看懂（Ps：里面有些概念事实上我如今还是不怎么懂），但不影响你进一步学习。假设想更全面地学习CUDA，《GPGPU编程技术》比較客观具体地介绍了通用GPU编程的策略。看过这本书，能够对显卡有更深入的了解，揭开GPU的神奇面纱。后面《OpenGL编程指南》全然是为了体验图形交互带来的乐趣，能够有选择地看；《GPU高性能运算之CUDA》这本是师兄给的，适合高速查询（感觉是将官方编程手冊翻译了一遍）一些关键技术和概念。

    有了这些指导材料还不够，我们在做项目的时候，遇到的问题在这些书上肯定找不到，所以还须要有以下这些利器：

    这里面有非常多工具的使用手冊，如CUDA_GDB，Nsight。CUDA_Profiler等，方便调试程序；另一些实用的库，如CUFFT是专门用来做高速傅里叶变换的。CUBLAS是专用于线性代数（矩阵、向量计算）的，CUSPASE是专用于稀疏矩阵表示和计算的库。这些库的使用能够减少我们设计算法的难度，提高开发效率。另外还有些新手教程也是值得一读的，你会对NVCC编译器有更近距离的接触。

    好了，前言就这么多，本博主计划按例如以下顺序来讲述CUDA：

    1.了解设备

    2.线程并行

    3.块并行

    4.流并行

    5.线程通信

    6.线程通信实例：规约

    7.存储模型

    8.常数内存

    9.纹理内存

    10.主机页锁定内存

    11.图形互操作

    12.优化准则

    13.CUDA与MATLAB接口

    14.CUDA与MFC接口

    前面三节已经对CUDA做了一个简单的介绍，这一节開始真正进入编程环节。

    首先，刚開始学习的人应该对自己使用的设备有较为扎实的理解和掌握。这样对后面学习并行程序优化非常有帮助。了解硬件具体參数能够通过上节介绍的几本书和官方资料获得。但假设仍然认为不够直观，那么我们能够自己动手获得这些内容。

    以第二节例程为模板，我们稍加修改的部分代码例如以下：

    [cpp] view plaincopy
    1. // Add vectors in parallel.  
    2. cudaError_t cudaStatus;  
    3. int num = 0;  
    4. cudaDeviceProp prop;  
    5. cudaStatus = cudaGetDeviceCount(&num);  
    6. for(int i = 0;i<num;i++)  
    7. {  
    8.     cudaGetDeviceProperties(&prop,i);  
    9. }  
    10. cudaStatus = addWithCuda(c, a, b, arraySize);  

    这个修改的目的是让我们的程序自己主动通过调用cuda API函数获得设备数目和属性，所谓“知己知彼，百战不殆”。

    cudaError_t 是cuda错误类型，取值为整数。

    cudaDeviceProp为设备属性结构体，其定义能够从cuda Toolkit安装文件夹中找到。我的路径为：C:Program FilesNVIDIA GPU Computing ToolkitCUDAv5.0includedriver_types.h。找到定义为：

    [cpp] view plaincopy
    1. /** 
    2.  * CUDA device properties 
    3.  */  
    4. struct __device_builtin__ cudaDeviceProp  
    5. {  
    6.     char   name[256];                  /**< ASCII string identifying device */  
    7.     size_t totalGlobalMem;             /**< Global memory available on device in bytes */  
    8.     size_t sharedMemPerBlock;          /**< Shared memory available per block in bytes */  
    9.     int    regsPerBlock;               /**< 32-bit registers available per block */  
    10.     int    warpSize;                   /**< Warp size in threads */  
    11.     size_t memPitch;                   /**< Maximum pitch in bytes allowed by memory copies */  
    12.     int    maxThreadsPerBlock;         /**< Maximum number of threads per block */  
    13.     int    maxThreadsDim[3];           /**< Maximum size of each dimension of a block */  
    14.     int    maxGridSize[3];             /**< Maximum size of each dimension of a grid */  
    15.     int    clockRate;                  /**< Clock frequency in kilohertz */  
    16.     size_t totalConstMem;              /**< Constant memory available on device in bytes */  
    17.     int    major;                      /**< Major compute capability */  
    18.     int    minor;                      /**< Minor compute capability */  
    19.     size_t textureAlignment;           /**< Alignment requirement for textures */  
    20.     size_t texturePitchAlignment;      /**< Pitch alignment requirement for texture references bound to pitched memory */  
    21.     int    deviceOverlap;              /**< Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. */  
    22.     int    multiProcessorCount;        /**< Number of multiprocessors on device */  
    23.     int    kernelExecTimeoutEnabled;   /**< Specified whether there is a run time limit on kernels */  
    24.     int    integrated;                 /**< Device is integrated as opposed to discrete */  
    25.     int    canMapHostMemory;           /**< Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer */  
    26.     int    computeMode;                /**< Compute mode (See ::cudaComputeMode) */  
    27.     int    maxTexture1D;               /**< Maximum 1D texture size */  
    28.     int    maxTexture1DMipmap;         /**< Maximum 1D mipmapped texture size */  
    29.     int    maxTexture1DLinear;         /**< Maximum size for 1D textures bound to linear memory */  
    30.     int    maxTexture2D[2];            /**< Maximum 2D texture dimensions */  
    31.     int    maxTexture2DMipmap[2];      /**< Maximum 2D mipmapped texture dimensions */  
    32.     int    maxTexture2DLinear[3];      /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */  
    33.     int    maxTexture2DGather[2];      /**< Maximum 2D texture dimensions if texture gather operations have to be performed */  
    34.     int    maxTexture3D[3];            /**< Maximum 3D texture dimensions */  
    35.     int    maxTextureCubemap;          /**< Maximum Cubemap texture dimensions */  
    36.     int    maxTexture1DLayered[2];     /**< Maximum 1D layered texture dimensions */  
    37.     int    maxTexture2DLayered[3];     /**< Maximum 2D layered texture dimensions */  
    38.     int    maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */  
    39.     int    maxSurface1D;               /**< Maximum 1D surface size */  
    40.     int    maxSurface2D[2];            /**< Maximum 2D surface dimensions */  
    41.     int    maxSurface3D[3];            /**< Maximum 3D surface dimensions */  
    42.     int    maxSurface1DLayered[2];     /**< Maximum 1D layered surface dimensions */  
    43.     int    maxSurface2DLayered[3];     /**< Maximum 2D layered surface dimensions */  
    44.     int    maxSurfaceCubemap;          /**< Maximum Cubemap surface dimensions */  
    45.     int    maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */  
    46.     size_t surfaceAlignment;           /**< Alignment requirements for surfaces */  
    47.     int    concurrentKernels;          /**< Device can possibly execute multiple kernels concurrently */  
    48.     int    ECCEnabled;                 /**< Device has ECC support enabled */  
    49.     int    pciBusID;                   /**< PCI bus ID of the device */  
    50.     int    pciDeviceID;                /**< PCI device ID of the device */  
    51.     int    pciDomainID;                /**< PCI domain ID of the device */  
    52.     int    tccDriver;                  /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */  
    53.     int    asyncEngineCount;           /**< Number of asynchronous engines */  
    54.     int    unifiedAddressing;          /**< Device shares a unified address space with the host */  
    55.     int    memoryClockRate;            /**< Peak memory clock frequency in kilohertz */  
    56.     int    memoryBusWidth;             /**< Global memory bus width in bits */  
    57.     int    l2CacheSize;                /**< Size of L2 cache in bytes */  
    58.     int    maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */  
    59. };  

    后面的凝视已经说明了其字段代表意义，可能有些术语对于刚開始学习的人理解起来还是有一定困难。没关系。我们如今仅仅须要关注下面几个指标：

    name：就是设备名称；

    totalGlobalMem：就是显存大小；

    major,minor：CUDA设备版本号号，有1.1, 1.2, 1.3, 2.0, 2.1等多个版本号；

    clockRate：GPU时钟频率。

    multiProcessorCount：GPU大核数。一个大核（专业点称为流多处理器，SM。Stream-Multiprocessor）包括多个小核（流处理器。SP，Stream-Processor）

    编译。执行。我们在VS2008project的cudaGetDeviceProperties()函数处放一个断点，单步执行这一函数，然后用Watch窗体。切换到Auto页，展开+。在我的笔记本上得到例如以下结果：

    能够看到，设备名为GeForce 610M。显存1GB。设备版本号2.1（比較高端了。哈哈），时钟频率为950MHz（注意950000单位为kHz），大核数为1。

    在一些高性能GPU上（如Tesla，Kepler系列）。大核数可能达到几十甚至上百。能够做更大规模的并行处理。

    PS：今天看SDK代码时发如今help_cuda.h中有个函数实现从CUDA设备版本号查询对应大核中小核的数目。认为非常实用，以后编程序能够借鉴，摘抄例如以下：

    [cpp] view plaincopy
    1. // Beginning of GPU Architecture definitions  
    2. inline int _ConvertSMVer2Cores(int major, int minor)  
    3. {  
    4.     // Defines for GPU Architecture types (using the SM version to determine the # of cores per SM  
    5.     typedef struct  
    6.     {  
    7.         int SM; // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version  
    8.         int Cores;  
    9.     } sSMtoCores;  
    10.   
    11.     sSMtoCores nGpuArchCoresPerSM[] =  
    12.     {  
    13.         { 0x10,  8 }, // Tesla Generation (SM 1.0) G80 class  
    14.         { 0x11,  8 }, // Tesla Generation (SM 1.1) G8x class  
    15.         { 0x12,  8 }, // Tesla Generation (SM 1.2) G9x class  
    16.         { 0x13,  8 }, // Tesla Generation (SM 1.3) GT200 class  
    17.         { 0x20, 32 }, // Fermi Generation (SM 2.0) GF100 class  
    18.         { 0x21, 48 }, // Fermi Generation (SM 2.1) GF10x class  
    19.         { 0x30, 192}, // Kepler Generation (SM 3.0) GK10x class  
    20.         { 0x35, 192}, // Kepler Generation (SM 3.5) GK11x class  
    21.         {   -1, -1 }  
    22.     };  
    23.   
    24.     int index = 0;  
    25.   
    26.     while (nGpuArchCoresPerSM[index].SM != -1)  
    27.     {  
    28.         if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor))  
    29.         {  
    30.             return nGpuArchCoresPerSM[index].Cores;  
    31.         }  
    32.   
    33.         index++;  
    34.     }  
    35.   
    36.     // If we don't find the values, we default use the previous one to run properly  
    37.     printf("MapSMtoCores for SM %d.%d is undefined.  Default to use %d Cores/SM ", major, minor, nGpuArchCoresPerSM[7].Cores);  
    38.     return nGpuArchCoresPerSM[7].Cores;  
    39. }  
    40. // end of GPU Architecture definitions  

    可见。设备版本号2.1的一个大核有48个小核，而版本号3.0以上的一个大核有192个小核！

    前文说到过。当我们用的电脑上有多个显卡支持CUDA时。怎么来区分在哪个上执行呢？这里我们看一下addWithCuda这个函数是怎么做的。

    [cpp] view plaincopy
    1. cudaError_t cudaStatus;  
    2.   
    3. // Choose which GPU to run on, change this on a multi-GPU system.  
    4. cudaStatus = cudaSetDevice(0);  
    5. if (cudaStatus != cudaSuccess) {  
    6.     fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?

      ");  

    7.     goto Error;  
    8. }  

    使用了cudaSetDevice(0)这个操作，0表示能搜索到的第一个设备号，假设有多个设备，则编号为0,1,2...。

    再看我们本节加入的代码。有个函数cudaGetDeviceCount(&num)，这个函数用来获取设备总数，这样我们选择执行CUDA程序的设备号取值就是0,1,...num-1，于是能够一个个枚举设备，利用cudaGetDeviceProperties(&prop)获得其属性,然后利用一定排序、筛选算法，找到最符合我们应用的那个设备号opt，然后调用cudaSetDevice(opt)就可以选择该设备。

    选择标准能够从处理能力、版本号控制、名称等各个角度出发。

    后面讲述流并发过程时，还要用到这些API。

    假设希望了解很多其它硬件内容能够结合http://www.geforce.cn/hardware获取。

    多线程我们应该都不陌生，在操作系统中，进程是资源分配的基本单元，而线程是CPU时间调度的基本单元（这里如果仅仅有1个CPU）。

    将线程的概念引申到CUDA程序设计中，我们能够觉得线程就是执行CUDA程序的最小单元，前面我们建立的project代码中，有个核函数概念不知各位童鞋还记得没有，在GPU上每一个线程都会执行一次该核函数。

    但GPU上的线程调度方式与CPU有非常大不同。CPU上会有优先级分配，从高到低，相同优先级的能够採用时间片轮转法实现线程调度。GPU上线程没有优先级概念。全部线程机会均等，线程状态仅仅有等待资源和运行两种状态。假设资源未就绪。那么就等待；一旦就绪。马上运行。

    当GPU资源非常充裕时。全部线程都是并发运行的。这样加速效果非常接近理论加速比；而GPU资源少于总线程个数时。有一部分线程就会等待前面运行的线程释放资源，从而变为串行化运行。

    代码还是用上一节的吧。修改非常少，再贴一遍：

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"           //CUDA执行时API  
    2. #include "device_launch_parameters.h"     
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = threadIdx.x;  
    8.     c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d} ",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // 重点理解这个函数  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0; //GPU设备端数据指针  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;     //状态指示  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  //选择执行平台  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?

      ");  

    55.         goto Error;  
    56.     }  
    57.     // 分配GPU设备端内存  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // 拷贝数据到GPU  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89.     // 执行核函数  
    90. <span style="BACKGROUND-COLOR: #ff6666"><strong>    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);</strong>  
    91. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    92.     // any errors encountered during the launch.  
    93.     cudaStatus = cudaThreadSynchronize();   //同步线程  
    94.     if (cudaStatus != cudaSuccess)   
    95.     {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.     // Copy output vector from GPU buffer to host memory.  
    100.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);      //拷贝结果回主机  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106. Error:  
    107.     cudaFree(dev_c);    //释放GPU设备端内存  
    108.     cudaFree(dev_a);  
    109.     cudaFree(dev_b);      
    110.     return cudaStatus;  
    111. }  

    红色部分即启动核函数的调用过程。这里看到调用方式和C不太一样。<<<>>>表示执行时配置符号。里面1表示仅仅分配一个线程组（又称线程块、Block）。size表示每一个线程组有size个线程（Thread）。本程序中size依据前面传递參数个数应该为5，所以执行的时候。核函数在5个GPU线程单元上分别执行了一次，总共执行了5次。

    这5个线程是怎样知道自己“身份”的？是靠threadIdx这个内置变量，它是个dim3类型变量。接受<<<>>>中第二个參数，它包括x,y,z 3维坐标，而我们传入的參数仅仅有一维。所以仅仅有x值是有效的。通过核函数中int i = threadIdx.x;这一句，每一个线程能够获得自身的id号，从而找到自己的任务去运行。

    CUDA从入门到精通（六）：块并行

     

    同一版本号的代码用了这么多次。有点过意不去，于是这次我要做较大的修改大笑。大家要擦亮眼睛，拭目以待。

    块并行相当于操作系统中多进程的情况。上节说到。CUDA有线程组（线程块）的概念，将一组线程组织到一起，共同分配一部分资源，然后内部调度运行。线程块与线程块之间。毫无瓜葛。这有利于做更粗粒度的并行。

    我们将上一节的代码改为块并行版本号例如以下：

    下节我们介绍块并行。

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7. <span style="BACKGROUND-COLOR: #ff0000">    int i = blockIdx.x;  
    8. </span>    c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d} ",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    55.         goto Error;  
    56.     }  
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89.     // Launch a kernel on the GPU with one thread for each element.  
    90.  <span style="BACKGROUND-COLOR: #ff0000">   addKernel<<<size,1 >>>(dev_c, dev_a, dev_b);  
    91. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    92.     // any errors encountered during the launch.  
    93.     cudaStatus = cudaThreadSynchronize();  
    94.     if (cudaStatus != cudaSuccess)   
    95.     {  
    96.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    97.         goto Error;  
    98.     }  
    99.     // Copy output vector from GPU buffer to host memory.  
    100.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106. Error:  
    107.     cudaFree(dev_c);  
    108.     cudaFree(dev_a);  
    109.     cudaFree(dev_b);      
    110.     return cudaStatus;  
    111. }  

    和上一节相比。仅仅有这两行有改变。<<<>>>里第一个參数改成了size，第二个改成了1，表示我们分配size个线程块，每一个线程块仅包括1个线程，总共还是有5个线程。这5个线程相互独立，运行核函数得到对应的结果。与上一节不同的是，每一个线程获取id的方式变为int i = blockIdx.x。这是线程块ID。

    于是有童鞋提问了。线程并行和块并行的差别在哪里？

    线程并行是细粒度并行，调度效率高；块并行是粗粒度并行，每次调度都要又一次分配资源。有时资源仅仅有一份。那么全部线程块都仅仅能排成一队，串行运行。

    那是不是我们全部时候都应该用线程并行，尽可能不用块并行？

    当然不是，我们的任务有时能够採用分治法，将一个大问题分解为几个小规模问题，将这些小规模问题分别用一个线程块实现，线程块内能够採用细粒度的线程并行。而块之间为粗粒度并行。这样能够充分利用硬件资源，减少线程并行的计算复杂度。

    适当分解。减少规模，在一些矩阵乘法、向量内积计算应用中能够得到充分的展示。

    实际应用中。经常是二者的结合。线程块、线程组织图例如以下所看到的。

    多个线程块组织成了一个Grid，称为线程格（经历了从一位线程，二维线程块到三维线程格的过程，立体感非常强啊）。

    好了，下一节我们介绍流并行。是更高层次的并行。

    前面我们没有讲程序的结构。我想有些童鞋可能迫不及待想知道CUDA程序究竟是怎么一个运行过程。好的，这一节在介绍流之前。先把CUDA程序结构简要说一下。

    CUDA程序文件后缀为.cu，有些编译器可能不认识这个后缀的文件，我们能够在VS2008的Tools->Options->Text Editor->File Extension里加入cu后缀到VC++中，例如以下图：

    一个.cu文件内既包括CPU程序（称为主机程序），也包括GPU程序（称为设备程序）。怎样区分主机程序和设备程序？依据声明。凡是挂有“__global__”或者“__device__”前缀的函数，都是在GPU上执行的设备程序。不同的是__global__设备程序可被主机程序调用，而__device__设备程序则仅仅能被设备程序调用。

    没有挂不论什么前缀的函数。都是主机程序。主机程序显示声明能够用__host__前缀。设备程序须要由NVCC进行编译。而主机程序仅仅须要由主机编译器（如VS2008中的cl.exe，Linux上的GCC）。

    主机程序主要完毕设备环境初始化。传输数据等必备过程，设备程序仅仅负责计算。

    主机程序中，有一些“cuda”打头的函数，这些都是CUDA Runtime API，即执行时函数，主要负责完毕设备的初始化、内存分配、内存拷贝等任务。

    我们前面第三节用到的函数cudaGetDeviceCount()。cudaGetDeviceProperties()。cudaSetDevice()都是执行时API。这些函数的详细參数声明我们不必一一记下来。拿出第三节的官方利器就能够轻松查询，让我们打开这个文件：

    打开后，在pdf搜索栏中输入一个执行时函数，比如cudaMemcpy，查到的结果例如以下：

    能够看到，该API函数的參数形式为，第一个表示目的地。第二个表示来源地。第三个參数表示字节数。第四个表示类型。假设对类型不了解，直接点击超链接，得到详解例如以下：

    可见，该API能够实现从主机到主机、主机到设备、设备到主机、设备到设备的内存拷贝过程。同一时候能够发现。利用该API手冊能够非常方便地查询我们须要用的这些API函数，所以以后编CUDA程序一定要把它打开，随时准备查询，这样能够大大提高编程效率。

    好了，进入今天的主题：流并行。

    前面已经介绍了线程并行和块并行。知道了线程并行为细粒度的并行，而块并行为粗粒度的并行，同一时候也知道了CUDA的线程组织情况，即Grid-Block-Thread结构。

    一组线程并行处理能够组织为一个block，而一组block并行处理能够组织为一个Grid，非常自然地想到。Grid仅仅是一个网格，我们能否够利用多个网格来完毕并行处理呢？答案就是利用流。

    流能够实如今一个设备上执行多个核函数。前面的块并行也好。线程并行也好。执行的核函数都是同样的（代码一样，传递參数也一样）。而流并行，能够执行不同的核函数，也能够实现对同一个核函数传递不同的參数，实现任务级别的并行。

    CUDA中的流用cudaStream_t类型实现。用到的API有下面几个：cudaStreamCreate(cudaStream_t * s)用于创建流，cudaStreamDestroy(cudaStream_t s)用于销毁流，cudaStreamSynchronize()用于单个流同步。cudaDeviceSynchronize()用于整个设备上的全部流同步，cudaStreamQuery()用于查询一个流的任务是否已经完毕。

    详细的含义能够查询API手冊。

    以下我们将前面的两个样例中的任务改用流实现。仍然是{1,2,3,4,5}+{10,20,30,40,50} = {11,22,33,44,55}这个样例。代码例如以下：

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = blockIdx.x;  
    8.     c[i] = a[i] + b[i];  
    9. }  
    10. int main()  
    11. {  
    12.     const int arraySize = 5;  
    13.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    14.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
    15.     int c[arraySize] = { 0 };  
    16.     // Add vectors in parallel.  
    17.     cudaError_t cudaStatus;  
    18.     int num = 0;  
    19.     cudaDeviceProp prop;  
    20.     cudaStatus = cudaGetDeviceCount(&num);  
    21.     for(int i = 0;i<num;i++)  
    22.     {  
    23.         cudaGetDeviceProperties(&prop,i);  
    24.     }  
    25.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    26.     if (cudaStatus != cudaSuccess)   
    27.     {  
    28.         fprintf(stderr, "addWithCuda failed!");  
    29.         return 1;  
    30.     }  
    31.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d} ",c[0],c[1],c[2],c[3],c[4]);  
    32.     // cudaThreadExit must be called before exiting in order for profiling and  
    33.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    34.     cudaStatus = cudaThreadExit();  
    35.     if (cudaStatus != cudaSuccess)   
    36.     {  
    37.         fprintf(stderr, "cudaThreadExit failed!");  
    38.         return 1;  
    39.     }  
    40.     return 0;  
    41. }  
    42. // Helper function for using CUDA to add vectors in parallel.  
    43. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    44. {  
    45.     int *dev_a = 0;  
    46.     int *dev_b = 0;  
    47.     int *dev_c = 0;  
    48.     cudaError_t cudaStatus;  
    49.   
    50.     // Choose which GPU to run on, change this on a multi-GPU system.  
    51.     cudaStatus = cudaSetDevice(0);  
    52.     if (cudaStatus != cudaSuccess)   
    53.     {  
    54.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    55.         goto Error;  
    56.     }  
    57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    59.     if (cudaStatus != cudaSuccess)   
    60.     {  
    61.         fprintf(stderr, "cudaMalloc failed!");  
    62.         goto Error;  
    63.     }  
    64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    65.     if (cudaStatus != cudaSuccess)   
    66.     {  
    67.         fprintf(stderr, "cudaMalloc failed!");  
    68.         goto Error;  
    69.     }  
    70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    71.     if (cudaStatus != cudaSuccess)   
    72.     {  
    73.         fprintf(stderr, "cudaMalloc failed!");  
    74.         goto Error;  
    75.     }  
    76.     // Copy input vectors from host memory to GPU buffers.  
    77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    78.     if (cudaStatus != cudaSuccess)   
    79.     {  
    80.         fprintf(stderr, "cudaMemcpy failed!");  
    81.         goto Error;  
    82.     }  
    83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    84.     if (cudaStatus != cudaSuccess)   
    85.     {  
    86.         fprintf(stderr, "cudaMemcpy failed!");  
    87.         goto Error;  
    88.     }  
    89. <span style="BACKGROUND-COLOR: #ff6666">  cudaStream_t stream[5];  
    90.     for(int i = 0;i<5;i++)  
    91.     {  
    92.         cudaStreamCreate(&stream[i]);   //创建流  
    93.     }  
    94. </span>    // Launch a kernel on the GPU with one thread for each element.  
    95. <span style="BACKGROUND-COLOR: #ff6666">  for(int i = 0;i<5;i++)  
    96.     {  
    97.         addKernel<<<1,1,0,stream[i]>>>(dev_c+i, dev_a+i, dev_b+i);    //运行流  
    98.     }  
    99.     cudaDeviceSynchronize();  
    100. </span>    // cudaThreadSynchronize waits for the kernel to finish, and returns  
    101.     // any errors encountered during the launch.  
    102.     cudaStatus = cudaThreadSynchronize();  
    103.     if (cudaStatus != cudaSuccess)   
    104.     {  
    105.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    106.         goto Error;  
    107.     }  
    108.     // Copy output vector from GPU buffer to host memory.  
    109.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    110.     if (cudaStatus != cudaSuccess)   
    111.     {  
    112.         fprintf(stderr, "cudaMemcpy failed!");  
    113.         goto Error;  
    114.     }  
    115. Error:  
    116. <span style="BACKGROUND-COLOR: #ff6666">  for(int i = 0;i<5;i++)  
    117.     {  
    118.         cudaStreamDestroy(stream[i]);   //销毁流  
    119.     }  
    120. </span>    cudaFree(dev_c);  
    121.     cudaFree(dev_a);  
    122.     cudaFree(dev_b);      
    123.     return cudaStatus;  
    124. }  

    注意到，我们的核函数代码仍然和块并行的版本号一样，仅仅是在调用时做了改变，<<<>>>中的參数多了两个，当中前两个和块并行、线程并行中的意义同样，仍然是线程块数（这里为1）、每一个线程块中线程数（这里也是1）。第三个为0表示每一个block用到的共享内存大小。这个我们后面再讲；第四个为流对象，表示当前核函数在哪个流上执行。我们创建了5个流，每一个流上都装载了一个核函数。同一时候传递參数有些不同，也就是每一个核函数作用的对象也不同。这样就实现了任务级别的并行，当我们有几个互不相关的任务时，能够写多个核函数。资源同意的情况下，我们将这些核函数装载到不同流上，然后执行，这样能够实现更粗粒度的并行。

    好了，流并行就这么简单，我们处理任务时。能够依据须要。选择最适合的并行方式。

    我们前面几节主要介绍了三种利用GPU实现并行处理的方式：线程并行，块并行和流并行。

    在这些方法中。我们一再强调，各个线程所进行的处理是互不相关的，即两个线程不回产生交集，每一个线程都仅仅关注自己的一亩三分地，对其它线程毫无兴趣，就当不存在。。。。

    当然。实际应用中，这种样例太少了，也就是遇到向量相加、向量相应点乘这类才会有如此高的并行度，而其它一些应用，如一组数求和，求最大（小）值，各个线程不再是相互独立的，而是产生一定关联，线程2可能会用到线程1的结果，这时就须要利用本节的线程通信技术了。

    线程通信在CUDA中有三种实现方式：

    1. 共享存储器；

    2. 线程 同步。

    3. 原子操作；

    最经常使用的是前两种方式，共享存储器，术语Shared Memory，是位于SM中的特殊存储器。还记得SM吗，就是流多处理器，大核是也。

    一个SM中不仅包括若干个SP（流处理器，小核），还包括一部分快速Cache，寄存器组。共享内存等，结构如图所看到的：

    从图中可看出，一个SM内有M个SP，Shared Memory由这M个SP共同占有。

    另外指令单元也被这M个SP共享，即SIMT架构（单指令多线程架构）。一个SM中全部SP在同一时间运行同一代码。

    为了实现线程通信，只靠共享内存还不够，须要有同步机制才干使线程之间实现有序处理。

    通常情况是这样：当线程A须要线程B计算的结果作为输入时。须要确保线程B已经将结果写入共享内存中，然后线程A再从共享内存中读出。同步不可缺少，否则。线程A可能读到的是无效的结果。造成计算错误。同步机制能够用CUDA内置函数：__syncthreads()；当某个线程运行到该函数时，进入等待状态。直到同一线程块（Block）中全部线程都运行到这个函数为止，即一个__syncthreads()相当于一个线程同步点，确保一个Block中全部线程都达到同步，然后线程进入运行状态。

    综上两点，我们能够写一段线程通信的伪代码例如以下：

    [cpp] view plaincopy
    1. //Begin  
    2. if this is thread B  
    3.      write something to Shared Memory;  
    4. end if  
    5. __syncthreads();  
    6. if this is thread A  
    7.     read something from Shared Memory;  
    8. end if  
    9. //End  

    上面代码在CUDA中实现时，因为SIMT特性。全部线程都运行相同的代码，所以在线程中须要推断自己的身份，以免误操作。

    注意的是，位于同一个Block中的线程才干实现通信，不同Block中的线程不能通过共享内存、同步进行通信，而应採用原子操作或主机介入。

    对于原子操作，假设感兴趣能够翻阅《GPU高性能编程CUDA实战》第九章“原子性”。

    本节完。下节我们给出一个实例来看线程通信的代码怎么设计。

    接着上一节，我们利用刚学到的共享内存和线程同步技术，来做一个简单的样例。

    先看下效果吧：

    非常easy。就是分别求出1~5这5个数字的和，平方和，连乘积。相信学过C语言的童鞋都能用for循环做出同上面一样的效果。但为了学习CUDA共享内存和同步技术，我们还是要把简单的东西复杂化(^_^)。

    简要分析一下，上面样例的输入都是一样的，1,2,3,4,5这5个数，但计算过程有些变化。并且每一个输出和全部输入都相关，不是前几节样例中那样，一个输出仅仅和一个输入有关。

    所以我们在利用CUDA编程时。须要针对特殊问题做些让步，把一些步骤串行化实现。

    输入数据原本位于主机内存，通过cudaMemcpy API已经复制到GPU显存（术语为全局存储器，Global Memory）。每一个线程执行时须要从Global Memory读取输入数据。然后完毕计算。最后将结果写回Global Memory。当我们计算须要多次同样输入数据时，大家可能想到，每次都分别去Global Memory读数据好像有点浪费，假设数据非常大。那么重复多次读数据会相当耗时间。

    索性我们把它从Global Memory一次性读到SM内部，然后在内部进行处理，这样能够节省重复读取的时间。

    有了这个思路，结合上节看到的SM结构图。看到有一片存储器叫做Shared Memory，它位于SM内部。处理时訪问速度相当快（差点儿相同每一个时钟周期读一次），而全局存储器读一次须要耗费几十甚至上百个时钟周期。于是，我们就制定A计划例如以下：

    线程块数：1，块号为0。（仅仅有一个线程块内的线程才干进行通信。所以我们仅仅分配一个线程块。详细工作交给每一个线程完毕）

    线程数：5，线程号分别为0~4。（线程并行。前面讲过）

    共享存储器大小：5个int型变量大小（5 * sizeof(int））。

    步骤一：读取输入数据。将Global Memory中的5个整数读入共享存储器，位置一一相应，和线程号也一一相应，所以能够同一时候完毕。

    步骤二：线程同步。确保全部线程都完毕了工作。

    步骤三：指定线程，对共享存储器中的输入数据完毕对应处理。

    代码例如以下：

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3.   
    4. #include <stdio.h>  
    5.   
    6. cudaError_t addWithCuda(int *c, const int *a, size_t size);  
    7.   
    8. __global__ void addKernel(int *c, const int *a)  
    9. {  
    10.     int i = threadIdx.x;  
    11. <span style="font-size:24px;"><strong>  extern __shared__ int smem[];</strong>  
    12. </span>   smem[i] = a[i];  
    13.     __syncthreads();  
    14.     if(i == 0)  // 0号线程做平方和  
    15.     {  
    16.         c[0] = 0;  
    17.         for(int d = 0; d < 5; d++)  
    18.         {  
    19.             c[0] += smem[d] * smem[d];  
    20.         }  
    21.     }  
    22.     if(i == 1)//1号线程做累加  
    23.     {  
    24.         c[1] = 0;  
    25.         for(int d = 0; d < 5; d++)  
    26.         {  
    27.             c[1] += smem[d];  
    28.         }  
    29.     }  
    30.     if(i == 2)  //2号线程做累乘  
    31.     {  
    32.         c[2] = 1;  
    33.         for(int d = 0; d < 5; d++)  
    34.         {  
    35.             c[2] *= smem[d];  
    36.         }  
    37.     }  
    38. }  
    39.   
    40. int main()  
    41. {  
    42.     const int arraySize = 5;  
    43.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
    44.     int c[arraySize] = { 0 };  
    45.     // Add vectors in parallel.  
    46.     cudaError_t cudaStatus = addWithCuda(c, a, arraySize);  
    47.     if (cudaStatus != cudaSuccess)   
    48.     {  
    49.         fprintf(stderr, "addWithCuda failed!");  
    50.         return 1;  
    51.     }  
    52.     printf(" 1+2+3+4+5 = %d 1^2+2^2+3^2+4^2+5^2 = %d 1*2*3*4*5 = %d ", c[1], c[0], c[2]);  
    53.     // cudaThreadExit must be called before exiting in order for profiling and  
    54.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    55.     cudaStatus = cudaThreadExit();  
    56.     if (cudaStatus != cudaSuccess)   
    57.     {  
    58.         fprintf(stderr, "cudaThreadExit failed!");  
    59.         return 1;  
    60.     }  
    61.     return 0;  
    62. }  
    63.   
    64. // Helper function for using CUDA to add vectors in parallel.  
    65. cudaError_t addWithCuda(int *c, const int *a,  size_t size)  
    66. {  
    67.     int *dev_a = 0;  
    68.     int *dev_c = 0;  
    69.     cudaError_t cudaStatus;  
    70.   
    71.     // Choose which GPU to run on, change this on a multi-GPU system.  
    72.     cudaStatus = cudaSetDevice(0);  
    73.     if (cudaStatus != cudaSuccess)   
    74.     {  
    75.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    76.         goto Error;  
    77.     }  
    78.   
    79.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    80.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    81.     if (cudaStatus != cudaSuccess)   
    82.     {  
    83.         fprintf(stderr, "cudaMalloc failed!");  
    84.         goto Error;  
    85.     }  
    86.   
    87.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    88.     if (cudaStatus != cudaSuccess)   
    89.     {  
    90.         fprintf(stderr, "cudaMalloc failed!");  
    91.         goto Error;  
    92.     }  
    93.     // Copy input vectors from host memory to GPU buffers.  
    94.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    95.     if (cudaStatus != cudaSuccess)   
    96.     {  
    97.         fprintf(stderr, "cudaMemcpy failed!");  
    98.         goto Error;  
    99.     }  
    100.     // Launch a kernel on the GPU with one thread for each element.  
    101. <span style="font-size:24px;"><strong>    addKernel<<<1, size, size * sizeof(int), 0>>>(dev_c, dev_a);</strong>  
    102. </span>  
    103.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    104.     // any errors encountered during the launch.  
    105.     cudaStatus = cudaThreadSynchronize();  
    106.     if (cudaStatus != cudaSuccess)   
    107.     {  
    108.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    109.         goto Error;  
    110.     }  
    111.   
    112.     // Copy output vector from GPU buffer to host memory.  
    113.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    114.     if (cudaStatus != cudaSuccess)   
    115.     {  
    116.         fprintf(stderr, "cudaMemcpy failed!");  
    117.         goto Error;  
    118.     }  
    119.   
    120. Error:  
    121.     cudaFree(dev_c);  
    122.     cudaFree(dev_a);      
    123.     return cudaStatus;  
    124. }  
    从代码中看到运行配置<<<>>>中第三个參数为共享内存大小（字节数），这样我们就知道了所有4个运行配置參数的意义。恭喜，你的CUDA最终入门了！


    入门后的进一步学习的内容。就是怎样优化自己的代码。我们前面的样例没有考虑不论什么性能方面优化，是为了更好地学习基本知识点，而不是其它细节问题。

    从本节開始，我们要从性能出发考虑问题，不断优化代码，使运行速度提高是并行处理的唯一目的。

    測试代码执行速度有非常多方法。C语言里提供了类似于SystemTime()这种API获得系统时间，然后计算两个事件之间的时长从而完毕计时功能。在CUDA中，我们有专门測量设备执行时间的API。以下一一介绍。

    翻开编程手冊《CUDA_Toolkit_Reference_Manual》，随时准备查询不懂得API。我们在执行核函数前后，做例如以下操作：

    [cpp] view plaincopy
    1. cudaEvent_t start, stop;<span style="white-space:pre">  </span>//事件对象  
    2. cudaEventCreate(&start);<span style="white-space:pre">  </span>//创建事件  
    3. cudaEventCreate(&stop);<span style="white-space:pre">       </span>//创建事件  
    4. cudaEventRecord(start, stream);<span style="white-space:pre">   </span>//记录開始  
    5. myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//运行核函数  
    6.   
    7. cudaEventRecord(stop,stream);<span style="white-space:pre"> </span>//记录结束事件  
    8. cudaEventSynchronize(stop);<span style="white-space:pre">   </span>//事件同步，等待结束事件之前的设备操作均已完毕  
    9. float elapsedTime;  
    10. cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长（单位为ms）  

    核函数运行时间将被保存在变量elapsedTime中。通过这个值我们能够评估算法的性能。

    以下给一个样例。来看怎么使用计时功能。

    前面的样例规模非常小，仅仅有5个元素，处理量太小不足以计时。以下将规模扩大为1024。此外将重复执行1000次计算总时间，这样预计不easy受随机扰动影响。我们通过这个样例对照线程并行和块并行的性能怎样。代码例如以下：

    [cpp] view plaincopy
    1. #include "cuda_runtime.h"  
    2. #include "device_launch_parameters.h"  
    3. #include <stdio.h>  
    4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
    5. __global__ void addKernel_blk(int *c, const int *a, const int *b)  
    6. {  
    7.     int i = blockIdx.x;  
    8.     c[i] = a[i]+ b[i];  
    9. }  
    10. __global__ void addKernel_thd(int *c, const int *a, const int *b)  
    11. {  
    12.     int i = threadIdx.x;  
    13.     c[i] = a[i]+ b[i];  
    14. }  
    15. int main()  
    16. {  
    17.     const int arraySize = 1024;  
    18.     int a[arraySize] = {0};  
    19.     int b[arraySize] = {0};  
    20.     for(int i = 0;i<arraySize;i++)  
    21.     {  
    22.         a[i] = i;  
    23.         b[i] = arraySize-i;  
    24.     }  
    25.     int c[arraySize] = {0};  
    26.     // Add vectors in parallel.  
    27.     cudaError_t cudaStatus;  
    28.     int num = 0;  
    29.     cudaDeviceProp prop;  
    30.     cudaStatus = cudaGetDeviceCount(&num);  
    31.     for(int i = 0;i<num;i++)  
    32.     {  
    33.         cudaGetDeviceProperties(&prop,i);  
    34.     }  
    35.     cudaStatus = addWithCuda(c, a, b, arraySize);  
    36.     if (cudaStatus != cudaSuccess)   
    37.     {  
    38.         fprintf(stderr, "addWithCuda failed!");  
    39.         return 1;  
    40.     }  
    41.   
    42.     // cudaThreadExit must be called before exiting in order for profiling and  
    43.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
    44.     cudaStatus = cudaThreadExit();  
    45.     if (cudaStatus != cudaSuccess)   
    46.     {  
    47.         fprintf(stderr, "cudaThreadExit failed!");  
    48.         return 1;  
    49.     }  
    50.     for(int i = 0;i<arraySize;i++)  
    51.     {  
    52.         if(c[i] != (a[i]+b[i]))  
    53.         {  
    54.             printf("Error in %d ",i);  
    55.         }  
    56.     }  
    57.     return 0;  
    58. }  
    59. // Helper function for using CUDA to add vectors in parallel.  
    60. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
    61. {  
    62.     int *dev_a = 0;  
    63.     int *dev_b = 0;  
    64.     int *dev_c = 0;  
    65.     cudaError_t cudaStatus;  
    66.   
    67.     // Choose which GPU to run on, change this on a multi-GPU system.  
    68.     cudaStatus = cudaSetDevice(0);  
    69.     if (cudaStatus != cudaSuccess)   
    70.     {  
    71.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
    72.         goto Error;  
    73.     }  
    74.     // Allocate GPU buffers for three vectors (two input, one output)    .  
    75.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
    76.     if (cudaStatus != cudaSuccess)   
    77.     {  
    78.         fprintf(stderr, "cudaMalloc failed!");  
    79.         goto Error;  
    80.     }  
    81.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
    82.     if (cudaStatus != cudaSuccess)   
    83.     {  
    84.         fprintf(stderr, "cudaMalloc failed!");  
    85.         goto Error;  
    86.     }  
    87.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
    88.     if (cudaStatus != cudaSuccess)   
    89.     {  
    90.         fprintf(stderr, "cudaMalloc failed!");  
    91.         goto Error;  
    92.     }  
    93.     // Copy input vectors from host memory to GPU buffers.  
    94.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
    95.     if (cudaStatus != cudaSuccess)   
    96.     {  
    97.         fprintf(stderr, "cudaMemcpy failed!");  
    98.         goto Error;  
    99.     }  
    100.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
    101.     if (cudaStatus != cudaSuccess)   
    102.     {  
    103.         fprintf(stderr, "cudaMemcpy failed!");  
    104.         goto Error;  
    105.     }  
    106.     cudaEvent_t start,stop;  
    107.     cudaEventCreate(&start);  
    108.     cudaEventCreate(&stop);  
    109.     cudaEventRecord(start,0);  
    110.     for(int i = 0;i<1000;i++)  
    111.     {  
    112. //      addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);  
    113.         addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);  
    114.     }  
    115.     cudaEventRecord(stop,0);  
    116.     cudaEventSynchronize(stop);  
    117.     float tm;  
    118.     cudaEventElapsedTime(&tm,start,stop);  
    119.     printf("GPU Elapsed time:%.6f ms. ",tm);  
    120.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
    121.     // any errors encountered during the launch.  
    122.     cudaStatus = cudaThreadSynchronize();  
    123.     if (cudaStatus != cudaSuccess)   
    124.     {  
    125.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel! ", cudaStatus);  
    126.         goto Error;  
    127.     }  
    128.     // Copy output vector from GPU buffer to host memory.  
    129.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
    130.     if (cudaStatus != cudaSuccess)   
    131.     {  
    132.         fprintf(stderr, "cudaMemcpy failed!");  
    133.         goto Error;  
    134.     }  
    135. Error:  
    136.     cudaFree(dev_c);  
    137.     cudaFree(dev_a);  
    138.     cudaFree(dev_b);      
    139.     return cudaStatus;  
    140. }  

    addKernel_blk是採用块并行实现的向量相加操作，而addKernel_thd是採用线程并行实现的向量相加操作。分别执行，得到的结果例如以下图所看到的：

    线程并行：

    块并行：

    可见性能居然相差近16倍！

    因此选择并行处理方法时，假设问题规模不是非常大。那么採用线程并行是比較合适的。而大问题分多个线程块处理时。每一个块内线程数不要太少，像本文中的仅仅有1个线程，这是对硬件资源的极大浪费。

    一个理想的方案是。分N个线程块。每一个线程块包括512个线程，将问题分解处理，效率往往比单一的线程并行处理或单一块并行处理高非常多。这也是CUDA编程的精髓。

    上面这样的分析程序性能的方式比較粗糙。仅仅知道大概执行时间长度。对于设备程序各部分代码执行时间没有一个深入的认识，这样我们就有个问题，假设对代码进行优化，那么优化哪一部分呢？是将线程数调节呢，还是改用共享内存？这个问题最好的解决方式就是利用Visual Profiler。

    以下内容摘自《CUDA_Profiler_Users_Guide》

    “Visual Profiler是一个图形化的剖析工具。能够显示你的应用程序中CPU和GPU的活动情况，利用分析引擎帮助你寻找优化的机会。”

    事实上除了可视化的界面，NVIDIA提供了命令行方式的剖析命令：nvprof。对于刚開始学习的人，使用图形化的方式比較easy上手。所以本节使用Visual Profiler。

    打开Visual Profiler，能够从CUDA Toolkit安装菜单处找到。主界面例如以下：

    我们点击File->New Session，弹出新建会话对话框，例如以下图所看到的：

    当中File一栏填入我们须要进行剖析的应用程序exe文件。后面能够都不填（假设须要命令行參数。能够在第三行填入）。直接Next。见下图：

    第一行为应用程序执行超时时间设定，可不填；后面三个单选框都勾上。这样我们分别使能了剖析，使能了并发核函数剖析。然后执行分析器。

    点Finish，開始执行我们的应用程序并进行剖析、分析性能。

    上图中，CPU和GPU部分显示了硬件和执行内容信息。点某一项则将时间条相应的部分高亮。便于观察。同一时候右边具体信息会显示执行时间信息。

    从时间条上看出，cudaMalloc占用了非常大一部分时间。以下分析器给出了一些性能提升的关键点。包含：低计算利用率（计算时间仅仅占总时间的1.8%。也难怪。加法计算复杂度本来就非常低呀！

    ）；低内存拷贝/计算交叠率（一点都没有交叠。全然是拷贝——计算——拷贝）；低存储拷贝尺寸（输入数据量太小了。相当于你淘宝买了个日记本。运费比实物价格还高！）；低存储拷贝吞吐率（仅仅有1.55GB/s）。这些对我们进一步优化程序是非常有帮助的。

    我们点一下Details，就在Analysis窗体旁边。得到结果例如以下所看到的：

    通过这个窗体能够看到每一个核函数运行时间，以及线程格、线程块尺寸。占用寄存器个数。静态共享内存、动态共享内存大小等參数，以及内存拷贝函数的运行情况。

    这个提供了比前面cudaEvent函数測时间更精确的方式，直接看到每一步的运行时间，精确到ns。

    在Details后面另一个Console。点一下看看。

    这个事实上就是命令行窗体，显示执行输出。看到增加了Profiler信息后，总执行时间变长了（原来线程并行版本号的程序执行时间仅仅需4ms左右）。

    这也是“測不准定理”决定的。假设我们希望測量更细微的时间，那么总时间肯定是不准的。假设我们希望測量总时间，那么细微的时间就被忽略掉了。

    后面Settings就是我们建立会话时的參数配置，不再详述。

    通过本节，我们应该能对CUDA性能提升有了一些想法。好。下一节我们将讨论怎样优化CUDA程序。

    http://blog.csdn.net/kkk584520/article/details/9413973

    http://blog.csdn.net/kkk584520/article/details/9414191

    http://blog.csdn.net/kkk584520/article/details/9415199

    http://blog.csdn.net/kkk584520/article/details/9417251

    http://blog.csdn.net/kkk584520/article/details/9420793

    http://blog.csdn.net/kkk584520/article/details/9428389

    http://blog.csdn.net/kkk584520/article/details/9428859

    http://blog.csdn.net/kkk584520/article/details/9449635

    http://blog.csdn.net/kkk584520/article/details/9472695

    http://blog.csdn.net/kkk584520/article/details/9473319

    http://blog.csdn.net/kkk584520/article/details/9490233
  • 相关阅读:
    紫书 例题8-18 UVa 1442 （扫描法）
    紫书 例题8-17 UVa 1609 （构造法）（详细注释）
    紫书 例题8-16 UVa 1608 （递归）
    紫书 例题8-15 UVa 12174 （滑动窗口）
    紫书 例题8-14 UVa 1607 （二分）
    紫书 例题8-13 UVa 11093 （反证法）
    紫书 例题8-12 UVa 12627 （找规律 + 递归）
    Codeforces Round #441 (Div. 2, by Moscow Team Olympiad)
    CodeForces
    CodeForces 444C 线段树
  • 原文地址：https://www.cnblogs.com/mengfanrong/p/5334857.html
Copyright © 2020-2023  润新知