CUDA從入門到精通

時間  2019-11-16
標籤 cuda 入門 精通 简体版
原文   原文鏈接

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

CUDA從入門到精通（零）：寫在前面
算法

在老闆的要求下。本博主從2012年上高性能計算課程開始接觸CUDA編程，隨後將該技術應用到了實際項目中。使處理程序加速超過1K，可見基於圖形顯示器的並行計算對於追求速度的應用來講無疑是一個理想的選擇。編程

還有不到一年畢業。怕是畢業後這些技術也就隨畢業而去，準備這個暑假開闢一個CUDA專欄，從入門到精通。步步爲營。順便分享設計的一些經驗教訓。但願能給學習CUDA的童鞋提供必定指導。windows

我的能力所及，錯誤不免。歡迎討論。多線程

 

PS：申請專欄好像需要先發原創帖超過15篇。架構

。。算了，先寫夠再申請吧。到時候一併轉過去。併發


CUDA從入門到精通（一）：環境搭建app

NVIDIA於2006年推出CUDA（Compute Unified Devices Architecture），可以利用其推出的GPU進行通用計算。將並行計算從大型集羣擴展到了普通顯卡，使得用戶僅僅需要一臺帶有Geforce顯卡的筆記本就能跑較大規模的並行處理程序。dom

 

使用顯卡的優勢是，和大型集羣相比功耗很是低。成本也不高。但性能很是突出。以個人筆記本爲例，Geforce 610M，用DeviceQuery程序測試。可獲得例如如下硬件參數：async

計算能力達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 plain copy
  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}\n",  
  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!\n", 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 plain copy
  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 plain copy
  1. 1>------ Build started: Project: cuda_helloworld, Configuration: Debug Win32 ------    
  2. 1>Compiling with CUDA Build Rule...    
  3. 1>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\nvcc.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.0\VC\bin"    -Xcompiler "/EHsc /W3 /nologo /O2 /Zi   /MT  "  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.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 Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart*.dll" "C:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\Debug"    
  16. 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll    
  17. 1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll    
  18. 1>已複製         2 個文件。    
  19. 1>Build log was saved at "file://c:\Users\DongXiaoman\Documents\Visual Studio 2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.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 plain copy
  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 Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\driver_types.h。找到定義爲：

[cpp]  view plain copy
  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 plain copy
  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\n", 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 plain copy
  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 plain copy
  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}\n",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!\n", 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 plain copy
  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}\n",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!\n", 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 plain copy
  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}\n",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!\n", 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 plain copy
  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 plain copy
  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("\t1+2+3+4+5 = %d\n\t1^2+2^2+3^2+4^2+5^2 = %d\n\t1*2*3*4*5 = %d\n\n\n\n\n\n", 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!\n", 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 plain copy
  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 plain copy
  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\n",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.\n",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!\n", 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
相關文章
相關標籤/搜索
每日一句
    每一个你不满意的现在，都有一个你没有努力的曾经。
本站公眾號
   歡迎關注本站公眾號,獲取更多信息
聯繫我們 最近搜索 最新文章 沪ICP备13005482号-10 MyBatis教程 SQL 教程 MySQL教程 Java 教程 Thymeleaf 教程 Hibernate教程 Spring教程 Redis教程