原文鏈接：http://blog.csdn.net/augusdi/article/details/12205435
作者專欄：http://blog.csdn.net/augusdi/article/category/1663451
（一）： VS10設置
? ? 在開始學習之前，首先要做的就是找到一本好的教材，要知道一本好的教材可以讓我們更加輕松地入門。在看了一些個CUDA編程相關的教材之后，我向大家推薦的一本教材叫做《GPU高性能編程CUDA實戰》。本教材相比其他的教材而言，它講得比較細，對于一些我們可能不太明白的知識點做了詳細的說明 。而且這本教材以層層深入的方式向我們展示了GPU的世界，從而引領我們進入CUDA編程的大門。
?????? 其他的教材的話我建議的是《深入淺出談CUDA》以及《CUDA編程指南》，前一本書可以和《GPU高性能編程CUDA實戰》這本書一起看，我建議先看《深入淺出談CUDA》，然后再看《GPU高性能編程CUDA實戰》，這樣理解起來能夠更加的透徹。
????? ? CUDA5.5是集成的CUDA ToolKit和SDK于一身的軟件，可以到NVIDIA Developer Zone去下載，下載地址https://developer.nvidia.com/cuda-downloads，有了工具我們就可以進行CUDA編程了。
????? ? CUDA安裝好之后，就可以配置VS2010的cuda環境了，此時cuda已經和你的vs綁定好了，默認安裝的VS10支持包。于是我們就可以開始我們的第一個cuda程序：
1.創建一個cuda項目
??????? 由于安裝cuda5.5的時候已經和vs2010綁定，所以直接可以創建一個cuda項目。選擇cuda 5.5 runtime就可以。
??????? 此時vs2010會自動為這個項目添加一個cuda程序：kernel.cu。直接編譯這個.cu文件，可能會出現"轉換到 COFF 期間失敗: 文件無效或損壞"這樣的錯誤信息，修改如下所示：
??????? 右鍵->工程屬性->配置屬性-> 清單工具->輸入和輸出->嵌入清單，選擇[否]
2.CUDA C/C++關鍵字和函數高亮顯示
??????? 在上面.cu文件中發現CUDA C/C++的關鍵字__global__等沒有高亮顯示，而且還有下劃曲線。下面進行CUDA C/C++關鍵字和函數的語法高亮顯示，配置Visual AssistX 函數高亮，代碼提示等功能。
下面是關于代碼高亮的設置。共三個設置
2.1.cu文件中C/C++關鍵字高亮
這個設置是讓VS2010編輯.cu文件時，把.cu文件里的C/C++語法高亮。
設置方法: 在VS2010的菜單 依次點擊：“工具|選項|文本編輯器|文件擴展名”，在這個界面里： “編輯器”下拉框選“Microsoft Visual C++”，在“Extension”文本框輸入.cu 點擊“添加”按鈕 ，重復工作把.cuh 添加為vc++類型，點擊確定按鈕。 把全部.cu文件關閉，再打開，.cu文件C++關鍵字就高亮了。如果不行就重啟VS2010。但是CUDA的關鍵字還是黑色的，下一步把CUDA關鍵自高亮顯示。
2.2.CUDA關鍵字高亮
為了讓CUDA的關鍵字，如__device__、dim3之類的文字高亮，需要如下步驟：
把SDK_PATH\C\doc\syntax_highlighting\usertype.dat復制到X:\Program Files (X86)\Microsoft Visual Studio 10.0\Common7\IDE\ 目錄下?
這里 X:是安裝VS2010盤符，這是win7 64位下的路徑。
注意：win7 64系統不要復制到這個目錄里，復制到這里是不會CUDA關鍵字高亮的：
X:\Program Files \Microsoft Visual Studio 10.0\Common7\IDE\
再次強調：SDK_PATH要換成你安裝SDK的實際路徑，不要直接使用這個字符串。
VS2010需要重啟，重啟后打開.cu文件，CUDA的關鍵字應該變成藍色了。
2.3. CUDA 函數高亮，及CUDA函數輸入代碼提示
a) 讓Visual AssistX支持CUDA函數高亮和代碼完成。在VS2010菜單里依次點擊: “VAssistX->Visual assist X Options->Projects->C/C++Directories ”在這個界面的 “Platform”下拉框選 Custom，在“Show Directories for..“下拉框選 Other include files， 然后在下面的輸入框里，新建、添加三個路徑，點擊確定，三個路徑分別如下：
SDK_PATH的\c\common\inc
這里SDK_PATH要用你安裝的SDK路徑代替，X:是你的CUDA toolkit安裝盤符。
這樣分別在這三個目錄里面的.h文件定義的函數和類型VAssistX會幫我們在編輯文件時候實現高亮和完成代碼。如果少包含一個路徑，這個目錄里的.h定義的函數和類型就不會高亮和提示了。
b)讓Visual AssistX支持.cu文件，也就讓Visual AssistX幫我們實現.cu文件高亮和代碼完成功能，需要編輯注冊表。首先關閉VS2010，點擊開始菜單輸入regedit打開注冊表，找到如下位置：HKEY_CURRENT_USER\Software\Whole Tomato\Visual Assist X\VANet10。在右邊找到ExtSource項目，鼠標右鍵選修改，在原有文字后添加如下文字：.cu;.cuh; 確定后
經過上面的安裝，開發程序應該沒問題了，但會發現有如下的問題：
1. VS2010顯示.cu文件時，有很多CUDA函數下有波浪線
2. 在.cu文件里使用VS2010的“Go to definition”功能查無法找到CUDA函數的定義。
這兩個問題的原因是VS2010不認識這些函數，沒有指定它們在那個頭文件里包含的。解決方法其實是一樣的，把CUDA tookit 和CUDA SDK的頭文件路徑加到VS2010的VC++目錄里面。
?
??????? CUDA編程主要做的就是和GPU打交道，在和這樣的一個陌生的家伙交流之前，我們需要做的就是先得認識和熟悉這個家伙。
在深入研究如何編寫設備代碼之前，我們需要通過某種機制來判斷計算機中當前有哪些設備，以及每個設備都支持哪些功能。幸運的是，可以通過一個非常簡單的接口來獲得這種信息。首先，我們希望知道在系統中有多少個設備是支持CUDA架構的，并且這些設備能夠運行基于CUDA C編寫的核函數。要獲得CUDA設備的數量．可以調用cudaGetDeviceCount（）。這個函數的作用從色的名字就可以看出來。在調用cudaGetDeviceCount（）后，可以對每個設備進行迭代、井查詢各個設備的相關信息。CUDA運行時將返回一個cudaDevice Prop類型的結構，其中包含了設備的相關屬性。我們可以獲得哪些屬性？從CUDA 3.0開始，在cudaDeviceProp結構中包含了以下信息：
[cpp] view plaincopy
struct ?cudaDeviceProp???{?? ????char ?name[256];??????????? ????size_t ?totalGlobalMem;???????????? ????size_t ?sharedMemPerBlock;????? ????int ?regsPerBlock;????????? ????int ?warpSize;????????????? ????size_t ?memPitch;????????/在內存復制中最大的修正量（Pitch），單位為字節?? ????int ?maxThreadsPerBlock;??????? ????int ?maxThreadsDim[3];????????? ????int ?maxGridSize?[3];?????????? ????size_t ?totalConstMem;????????? ????int ?major;???????????? ????int ?minor;???????????? ????int ?clockRate;???????????? ????size_t ?textureAlignment;?????? ????int ?deviceoverlap;???????? ????int ?multiProcessorCount;?????? ????int ?kernelExecTimeoutEnabled;????? ????int ?integrated;??????????? ????int ?canMapHostMemory;????????? ????int ?computeMode;?????????? ????int ?maxTexture1D;????????? ????int ?maxTexture2D[2];?????????? ????int ?maxTexture3D[3];?????????? ????int ?maxTexture2DArray[3];????? ????int ?concurrentKernels?;??????? };??
設備屬性的使用
通過上面的結構體，我們大致了解了設備的屬性，然后我們就可以通過這個結構體來查詢設備屬性了。可能會有人問，到底我們需要這些設備屬性來干嘛，別著急，以后在編寫相關性能優化的代碼的時候，就知道了解這些屬性的好處了。現在我們只需要知道方法就可以了。
首先我們可以通過兩個函數，第一個就是上面的cudaGetDeviceCount（）來選擇設備，然后循環地通過getDeviceProperties（）來獲得設備的屬性，之后我們就可以通過這樣的一個結構體變量將設備的屬性值獲取出來。
[cpp] view plaincopy
#include?<cuda_runtime.h> ??#include?<iostream> ??using ?namespace ?std;???? int ?main()??{?? ????cudaDeviceProp?prop;?? ?? ????int ?count;?? ????cudaGetDeviceCount(&count);?? ?? ????for (int ?i?=?0?;?i?<?count?;?i++)?? ????{?? ????????cudaGetDeviceProperties(&prop,i);?? ????????cout<<"the?information?for?the?device?:?" <<i<<endl;?? ????????cout<<"name:" <<prop.name<<endl;?? ????????cout<<"the?memory?information?for?the?device?:?" <<i<<endl;?? ????????cout<<"total?global?memory:" <<prop.totalGlobalMem<<endl;?? ????????cout<<"total?constant?memory:" <<prop.totalConstMem<<endl;?? ????????cout<<"threads?in?warps:" <<prop.warpSize<<endl;?? ????????cout<<"max?threads?per?block:" <<prop.maxThreadsPerBlock<<endl;?? ????????cout<<"max?threads?dims:" <<prop.maxThreadsDim[0]<<"??" <<prop.maxThreadsDim[1]<<?? ????????????"??" <<prop.maxThreadsDim[2]<<endl;?? ????????cout<<"max?grid?dims:" <<prop.maxGridSize[0]<<"??" <<?? ????????????prop.maxGridSize[1]<<"??" <<prop.maxGridSize[2]<<endl;?? ?? ????}?? ????return ?0;?? }??
我這邊只是獲取一部分的屬性值，只是和大家介紹一下，具體的屬性值可以按照這樣的方法來獲取。
（二）：從入門到入門
CUDA從入門到精通（一）：環境搭建
?
?
1. 硬件平臺，就是顯卡，如果你用的不是NVIDIA的顯卡，那么只能說抱歉，其他都不支持CUDA。
2. 操作系統，我用過windows XP，Windows 7都沒問題，本博客用Windows7。
3. C編譯器，建議VS2008，和本博客一致。
4. CUDA編譯器NVCC，可以免費免注冊免license從官網下載CUDA ToolkitCUDA下載，最新版本為5.0，本博客用的就是該版本。
5. 其他工具（如Visual Assist，輔助代碼高亮）
?
不錯，就是它，雙擊運行，可以看到一大堆例程。我們找到Simple OpenGL這個運行看看效果：
? 點右邊黃線標記處的Run即可看到美妙的三維正弦曲面，鼠標左鍵拖動可以轉換角度，右鍵拖動可以縮放。如果這個運行成功，說明你的環境基本搭建成功。
1. 你使用遠程桌面連接登錄到另一臺服務器，該服務器上有顯卡支持CUDA，但你遠程終端不能運行CUDA程序。這是因為遠程登錄使用的是你本地顯卡資源，在遠程登錄時看不到服務器端的顯卡，所以會報錯：沒有支持CUDA的顯卡！解決方法：1. 遠程服務器裝兩塊顯卡，一塊只用于顯示，另一塊用于計算；2.不要用圖形界面登錄，而是用命令行界面如telnet登錄。
2.有兩個以上顯卡都支持CUDA的情況，如何區分是在哪個顯卡上運行？這個需要你在程序里控制，選擇符合一定條件的顯卡，如較高的時鐘頻率、較大的顯存、較高的計算版本等。詳細操作見后面的博客。
好了，先說這么多，下一節我們介紹如何在VS2008中給GPU編程。
CUDA從入門到精通（二）：第一個CUDA程序
之后點OK，直接進入工程界面。
工程中，我們看到只有一個.cu文件，內容如下：
[cpp] view plaincopy
#include?"cuda_runtime.h" ??#include?"device_launch_parameters.h" ???? #include?<stdio.h> ???? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size);?? ?? __global__?void ?addKernel(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?threadIdx.x;?? ????c[i]?=?a[i]?+?b[i];?? }?? ?? int ?main()??{?? ????const ?int ?arraySize?=?5;?? ????const ?int ?a[arraySize]?=?{?1,?2,?3,?4,?5?};?? ????const ?int ?b[arraySize]?=?{?10,?20,?30,?40,?50?};?? ????int ?c[arraySize]?=?{?0?};?? ?? ?????? ????cudaError_t?cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ?? ????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]);?? ?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ?? ????return ?0;?? }?? ?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size)?? {?? ????int ?*dev_a?=?0;?? ????int ?*dev_b?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;?? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?? ????cudaStatus?=?cudaMalloc((void **)&dev_b,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?? ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?? ?????? ????addKernel<<<1,?size>>>(dev_c,?dev_a,?dev_b);?? ?? ?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);?? ????if ?(cudaStatus?!=?cudaSuccess)?{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?? Error:?? ????cudaFree(dev_c);?? ????cudaFree(dev_a);?? ????cudaFree(dev_b);?? ?????? ????return ?cudaStatus;?? }??
?可以看出，CUDA程序和C程序并無區別，只是多了一些以"cuda"開頭的一些庫函數和一個特殊聲明的函數：
[cpp] view plaincopy
__global__?void ?addKernel(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?threadIdx.x;?? ????c[i]?=?a[i]?+?b[i];?? }??
這個函數就是在GPU上運行的函數，稱之為核函數，英文名Kernel Function，注意要和操作系統內核函數區分開來。
我們直接按F7編譯，可以得到如下輸出：
[html] view plaincopy
1> ------?Build?started:?Project:?cuda_helloworld,?Configuration:?Debug?Win32?------???? 1> Compiling?with?CUDA?Build?Rule...???? 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?????? 1> tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.gpu???? 1> tmpxft_000000ec_00000000-14_kernel.compute_10.cudafe2.gpu???? 1> tmpxft_000000ec_00000000-5_kernel.compute_20.cudafe1.gpu???? 1> tmpxft_000000ec_00000000-17_kernel.compute_20.cudafe2.gpu???? 1> kernel.cu???? 1> kernel.cu???? 1> tmpxft_000000ec_00000000-8_kernel.compute_10.cudafe1.cpp???? 1> tmpxft_000000ec_00000000-24_kernel.compute_10.ii???? 1> Linking...???? 1> Embedding?manifest...???? 1> Performing?Post-Build?Event...???? 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"???? 1> C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\cudart32_50_35.dll???? 1> C:\Program?Files\NVIDIA?GPU?Computing?Toolkit\CUDA\v5.0\\bin\cudart64_50_35.dll???? 1> 已復制?????????2?個文件。???? 1> Build?log?was?saved?at?"file://c:\Users\DongXiaoman\Documents\Visual?Studio?2008\Projects\cuda_helloworld\cuda_helloworld\Debug\BuildLog.htm"???? 1> cuda_helloworld?-?0?error(s),?105?warning(s)???? ==========?Build:?1?succeeded,?0?failed,?0?up-to-date,?0?skipped ?==========????
可見，編譯.cu文件需要利用nvcc工具。該工具的詳細使用見后面博客。
直接運行，可以得到結果圖如下：
如果顯示正確，那么我們的第一個程序宣告成功！
CUDA從入門到精通（三）：必備資料
1.了解設備
2.線程并行
3.塊并行
4.流并行
5.線程通信
6.線程通信實例：規約
7.存儲模型
8.常數內存
9.紋理內存
10.主機頁鎖定內存
11.圖形互操作
12.優化準則
13.CUDA與MATLAB接口
14.CUDA與MFC接口
CUDA從入門到精通（四）：加深對設備的認識
[cpp] view plaincopy
?? cudaError_t?cudaStatus;?? int ?num?=?0;??cudaDeviceProp?prop;?? cudaStatus?=?cudaGetDeviceCount(&num);?? for (int ?i?=?0;i<num;i++)??{?? ????cudaGetDeviceProperties(&prop,i);?? }?? cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);??
cudaError_t 是cuda錯誤類型，取值為整數。
cudaDeviceProp為設備屬性結構體，其定義可以從cuda Toolkit安裝目錄中找到，我的路徑為：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\include\driver_types.h，找到定義為：
[cpp] view plaincopy
? ? ?? struct ?__device_builtin__?cudaDeviceProp??{?? ????char ???name[256];???????????????????? ????size_t ?totalGlobalMem;??????????????? ????size_t ?sharedMemPerBlock;???????????? ????int ????regsPerBlock;????????????????? ????int ????warpSize;????????????????????? ????size_t ?memPitch;????????????????????? ????int ????maxThreadsPerBlock;??????????? ????int ????maxThreadsDim[3];????????????? ????int ????maxGridSize[3];??????????????? ????int ????clockRate;???????????????????? ????size_t ?totalConstMem;???????????????? ????int ????major;???????????????????????? ????int ????minor;???????????????????????? ????size_t ?textureAlignment;????????????? ????size_t ?texturePitchAlignment;???????? ????int ????deviceOverlap;???????????????? ????int ????multiProcessorCount;?????????? ????int ????kernelExecTimeoutEnabled;????? ????int ????integrated;??????????????????? ????int ????canMapHostMemory;????????????? ????int ????computeMode;?????????????????? ????int ????maxTexture1D;????????????????? ????int ????maxTexture1DMipmap;??????????? ????int ????maxTexture1DLinear;??????????? ????int ????maxTexture2D[2];?????????????? ????int ????maxTexture2DMipmap[2];???????? ????int ????maxTexture2DLinear[3];???????? ????int ????maxTexture2DGather[2];???????? ????int ????maxTexture3D[3];?????????????? ????int ????maxTextureCubemap;???????????? ????int ????maxTexture1DLayered[2];??????? ????int ????maxTexture2DLayered[3];??????? ????int ????maxTextureCubemapLayered[2];?? ????int ????maxSurface1D;????????????????? ????int ????maxSurface2D[2];?????????????? ????int ????maxSurface3D[3];?????????????? ????int ????maxSurface1DLayered[2];??????? ????int ????maxSurface2DLayered[3];??????? ????int ????maxSurfaceCubemap;???????????? ????int ????maxSurfaceCubemapLayered[2];?? ????size_t ?surfaceAlignment;????????????? ????int ????concurrentKernels;???????????? ????int ????ECCEnabled;??????????????????? ????int ????pciBusID;????????????????????? ????int ????pciDeviceID;?????????????????? ????int ????pciDomainID;?????????????????? ????int ????tccDriver;???????????????????? ????int ????asyncEngineCount;????????????? ????int ????unifiedAddressing;???????????? ????int ????memoryClockRate;?????????????? ????int ????memoryBusWidth;??????????????? ????int ????l2CacheSize;?????????????????? ????int ????maxThreadsPerMultiProcessor;?? };??
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）
[cpp] view plaincopy
?? inline ?int ?_ConvertSMVer2Cores(int ?major,?int ?minor)??{?? ?????? ????typedef ?struct ?? ????{?? ????????int ?SM;??? ????????int ?Cores;?? ????}?sSMtoCores;?? ?? ????sSMtoCores?nGpuArchCoresPerSM[]?=?? ????{?? ????????{?0x10,??8?},??? ????????{?0x11,??8?},??? ????????{?0x12,??8?},??? ????????{?0x13,??8?},??? ????????{?0x20,?32?},??? ????????{?0x21,?48?},??? ????????{?0x30,?192},??? ????????{?0x35,?192},??? ????????{???-1,?-1?}?? ????};?? ?? ????int ?index?=?0;?? ?? ????while ?(nGpuArchCoresPerSM[index].SM?!=?-1)?? ????{?? ????????if ?(nGpuArchCoresPerSM[index].SM?==?((major?<<?4)?+?minor))?? ????????{?? ????????????return ?nGpuArchCoresPerSM[index].Cores;?? ????????}?? ?? ????????index++;?? ????}?? ?? ?????? ????printf("MapSMtoCores?for?SM?%d.%d?is?undefined.??Default?to?use?%d?Cores/SM\n" ,?major,?minor,?nGpuArchCoresPerSM[7].Cores);?? ????return ?nGpuArchCoresPerSM[7].Cores;?? }?? ??
[cpp] view plaincopy
cudaError_t?cudaStatus;?? ?? ?? cudaStatus?=?cudaSetDevice(0);?? if ?(cudaStatus?!=?cudaSuccess)?{??????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????goto ?Error;?? }??
CUDA從入門到精通（五）：線程并行
[cpp] view plaincopy
#include?"cuda_runtime.h"???????????//CUDA運行時API ??#include?"device_launch_parameters.h"??? ??#include?<stdio.h> ??cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size);?? __global__?void ?addKernel(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?threadIdx.x;?? ????c[i]?=?a[i]?+?b[i];?? }?? int ?main()??{?? ????const ?int ?arraySize?=?5;?? ????const ?int ?a[arraySize]?=?{?1,?2,?3,?4,?5?};?? ????const ?int ?b[arraySize]?=?{?10,?20,?30,?40,?50?};?? ????int ?c[arraySize]?=?{?0?};?? ?????? ????cudaError_t?cudaStatus;?? ????int ?num?=?0;?? ????cudaDeviceProp?prop;?? ????cudaStatus?=?cudaGetDeviceCount(&num);?? ????for (int ?i?=?0;i<num;i++)?? ????{?? ????????cudaGetDeviceProperties(&prop,i);?? ????}?? ????cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ????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]);?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ????return ?0;?? }?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size)?? {?? ????int ?*dev_a?=?0;??? ????int ?*dev_b?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;??????? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);???? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_b,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?????? <span?style="BACKGROUND-COLOR:?#ff6666" ><strong>????addKernel<<<1,?size>>>(dev_c,?dev_a,?dev_b);</strong>?? </span>?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();????? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);???????? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? Error:?? ????cudaFree(dev_c);?????? ????cudaFree(dev_a);?? ????cudaFree(dev_b);?????? ????return ?cudaStatus;?? }??
CUDA從入門到精通（六）：塊并行
?
[cpp] view plaincopy
#include?"cuda_runtime.h" ??#include?"device_launch_parameters.h" ??#include?<stdio.h> ??cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size);?? __global__?void ?addKernel(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? <span?style="BACKGROUND-COLOR:?#ff0000" >????int ?i?=?blockIdx.x;?? </span>????c[i]?=?a[i]?+?b[i];?? }?? int ?main()??{?? ????const ?int ?arraySize?=?5;?? ????const ?int ?a[arraySize]?=?{?1,?2,?3,?4,?5?};?? ????const ?int ?b[arraySize]?=?{?10,?20,?30,?40,?50?};?? ????int ?c[arraySize]?=?{?0?};?? ?????? ????cudaError_t?cudaStatus;?? ????int ?num?=?0;?? ????cudaDeviceProp?prop;?? ????cudaStatus?=?cudaGetDeviceCount(&num);?? ????for (int ?i?=?0;i<num;i++)?? ????{?? ????????cudaGetDeviceProperties(&prop,i);?? ????}?? ????cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ????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]);?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ????return ?0;?? }?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size)?? {?? ????int ?*dev_a?=?0;?? ????int ?*dev_b?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;?? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_b,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ?<span?style="BACKGROUND-COLOR:?#ff0000" >???addKernel<<<size,1?>>>(dev_c,?dev_a,?dev_b);?? </span>?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? Error:?? ????cudaFree(dev_c);?? ????cudaFree(dev_a);?? ????cudaFree(dev_b);?????? ????return ?cudaStatus;?? }??
CUDA從入門到精通（七）：流并行
[cpp] view plaincopy
#include?"cuda_runtime.h" ??#include?"device_launch_parameters.h" ??#include?<stdio.h> ??cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size);?? __global__?void ?addKernel(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?blockIdx.x;?? ????c[i]?=?a[i]?+?b[i];?? }?? int ?main()??{?? ????const ?int ?arraySize?=?5;?? ????const ?int ?a[arraySize]?=?{?1,?2,?3,?4,?5?};?? ????const ?int ?b[arraySize]?=?{?10,?20,?30,?40,?50?};?? ????int ?c[arraySize]?=?{?0?};?? ?????? ????cudaError_t?cudaStatus;?? ????int ?num?=?0;?? ????cudaDeviceProp?prop;?? ????cudaStatus?=?cudaGetDeviceCount(&num);?? ????for (int ?i?=?0;i<num;i++)?? ????{?? ????????cudaGetDeviceProperties(&prop,i);?? ????}?? ????cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ????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]);?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ????return ?0;?? }?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size)?? {?? ????int ?*dev_a?=?0;?? ????int ?*dev_b?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;?? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_b,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? <span?style="BACKGROUND-COLOR:?#ff6666" >??cudaStream_t?stream[5];?? ????for (int ?i?=?0;i<5;i++)?? ????{?? ????????cudaStreamCreate(&stream[i]);????? ????}?? </span>?????? <span?style="BACKGROUND-COLOR:?#ff6666" >??for (int ?i?=?0;i<5;i++)?? ????{?? ????????addKernel<<<1,1,0,stream[i]>>>(dev_c+i,?dev_a+i,?dev_b+i);?????? ????}?? ????cudaDeviceSynchronize();?? </span>?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? Error:?? <span?style="BACKGROUND-COLOR:?#ff6666" >??for (int ?i?=?0;i<5;i++)?? ????{?? ????????cudaStreamDestroy(stream[i]);????? ????}?? </span>????cudaFree(dev_c);?? ????cudaFree(dev_a);?? ????cudaFree(dev_b);?????? ????return ?cudaStatus;?? }??
UDA從入門到精通（八）：線程通信
1. 共享存儲器；
2. 線程?同步；
3. 原子操作；
[cpp] view plaincopy
?? if ?this ?is?thread ?B???????write?something?to?Shared?Memory;?? end?if ?? __syncthreads();?? if ?this ?is?thread ?A??????read?something?from?Shared?Memory;?? end?if ?? ??
對于原子操作，如果感興趣可以翻閱《GPU高性能編程CUDA實戰》第九章“原子性”。
CUDA從入門到精通（九）：線程通信實例
線程塊數：1，塊號為0；（只有一個線程塊內的線程才能進行通信，所以我們只分配一個線程塊，具體工作交給每個線程完成）
線程數：5，線程號分別為0~4；（線程并行，前面講過）
共享存儲器大小：5個int型變量大小（5 * sizeof(int））。
步驟一：讀取輸入數據。將Global Memory中的5個整數讀入共享存儲器，位置一一對應，和線程號也一一對應，所以可以同時完成。
步驟二：線程同步，確保所有線程都完成了工作。
步驟三：指定線程，對共享存儲器中的輸入數據完成相應處理。
[cpp] view plaincopy
#include?"cuda_runtime.h" ??#include?"device_launch_parameters.h" ???? #include?<stdio.h> ???? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?size_t ?size);?? ?? __global__?void ?addKernel(int ?*c,?const ?int ?*a)?? {?? ????int ?i?=?threadIdx.x;?? <span?style="font-size:24px;" ><strong>??extern ?__shared__?int ?smem[];</strong>?? </span>???smem[i]?=?a[i];?? ????__syncthreads();?? ????if (i?==?0)???? ????{?? ????????c[0]?=?0;?? ????????for (int ?d?=?0;?d?<?5;?d++)?? ????????{?? ????????????c[0]?+=?smem[d]?*?smem[d];?? ????????}?? ????}?? ????if (i?==?1)?? ????{?? ????????c[1]?=?0;?? ????????for (int ?d?=?0;?d?<?5;?d++)?? ????????{?? ????????????c[1]?+=?smem[d];?? ????????}?? ????}?? ????if (i?==?2)???? ????{?? ????????c[2]?=?1;?? ????????for (int ?d?=?0;?d?<?5;?d++)?? ????????{?? ????????????c[2]?*=?smem[d];?? ????????}?? ????}?? }?? ?? int ?main()??{?? ????const ?int ?arraySize?=?5;?? ????const ?int ?a[arraySize]?=?{?1,?2,?3,?4,?5?};?? ????int ?c[arraySize]?=?{?0?};?? ?????? ????cudaError_t?cudaStatus?=?addWithCuda(c,?a,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ????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]);?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ????return ?0;?? }?? ?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,??size_t ?size)?? {?? ????int ?*dev_a?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;?? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?????? <span?style="font-size:24px;" ><strong>????addKernel<<<1,?size,?size?*?sizeof (int ),?0>>>(dev_c,?dev_a);</strong>?? </span>?? ?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ?? Error:?? ????cudaFree(dev_c);?? ????cudaFree(dev_a);?????? ????return ?cudaStatus;?? }??
CUDA從入門到精通（十）：性能剖析和Visual Profiler
[cpp] view plaincopy
cudaEvent_t?start,?stop;<span?style="white-space:pre" >??</span>?? cudaEventCreate(&start);<span?style="white-space:pre" >??</span>?? cudaEventCreate(&stop);<span?style="white-space:pre" >???????</span>?? cudaEventRecord(start,?stream);<span?style="white-space:pre" >???</span>?? myKernel<<<dimg,dimb,size_smem,stream>>>(parameter?list);?? ?? cudaEventRecord(stop,stream);<span?style="white-space:pre" >?</span>?? cudaEventSynchronize(stop);<span?style="white-space:pre" >???</span>?? float ?elapsedTime;??cudaEventElapsedTime(&elapsedTime,start,stop);??
[cpp] view plaincopy
#include?"cuda_runtime.h" ??#include?"device_launch_parameters.h" ??#include?<stdio.h> ??cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size);?? __global__?void ?addKernel_blk(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?blockIdx.x;?? ????c[i]?=?a[i]+?b[i];?? }?? __global__?void ?addKernel_thd(int ?*c,?const ?int ?*a,?const ?int ?*b)?? {?? ????int ?i?=?threadIdx.x;?? ????c[i]?=?a[i]+?b[i];?? }?? int ?main()??{?? ????const ?int ?arraySize?=?1024;?? ????int ?a[arraySize]?=?{0};?? ????int ?b[arraySize]?=?{0};?? ????for (int ?i?=?0;i<arraySize;i++)?? ????{?? ????????a[i]?=?i;?? ????????b[i]?=?arraySize-i;?? ????}?? ????int ?c[arraySize]?=?{0};?? ?????? ????cudaError_t?cudaStatus;?? ????int ?num?=?0;?? ????cudaDeviceProp?prop;?? ????cudaStatus?=?cudaGetDeviceCount(&num);?? ????for (int ?i?=?0;i<num;i++)?? ????{?? ????????cudaGetDeviceProperties(&prop,i);?? ????}?? ????cudaStatus?=?addWithCuda(c,?a,?b,?arraySize);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"addWithCuda?failed!" );?? ????????return ?1;?? ????}?? ?? ?????? ?????? ????cudaStatus?=?cudaThreadExit();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadExit?failed!" );?? ????????return ?1;?? ????}?? ????for (int ?i?=?0;i<arraySize;i++)?? ????{?? ????????if (c[i]?!=?(a[i]+b[i]))?? ????????{?? ????????????printf("Error?in?%d\n" ,i);?? ????????}?? ????}?? ????return ?0;?? }?? ?? cudaError_t?addWithCuda(int ?*c,?const ?int ?*a,?const ?int ?*b,?size_t ?size)?? {?? ????int ?*dev_a?=?0;?? ????int ?*dev_b?=?0;?? ????int ?*dev_c?=?0;?? ????cudaError_t?cudaStatus;?? ?? ?????? ????cudaStatus?=?cudaSetDevice(0);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMalloc((void **)&dev_c,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_a,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMalloc((void **)&dev_b,?size?*?sizeof (int ));?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMalloc?failed!" );?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof (int ),?cudaMemcpyHostToDevice);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? ????cudaEvent_t?start,stop;?? ????cudaEventCreate(&start);?? ????cudaEventCreate(&stop);?? ????cudaEventRecord(start,0);?? ????for (int ?i?=?0;i<1000;i++)?? ????{?? ?? ????????addKernel_thd<<<1,size>>>(dev_c,?dev_a,?dev_b);?? ????}?? ????cudaEventRecord(stop,0);?? ????cudaEventSynchronize(stop);?? ????float ?tm ;?? ????cudaEventElapsedTime(&tm ,start,stop);?? ????printf("GPU?Elapsed?time:%.6f?ms.\n" ,tm );?? ?????? ?????? ????cudaStatus?=?cudaThreadSynchronize();?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n" ,?cudaStatus);?? ????????goto ?Error;?? ????}?? ?????? ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof (int ),?cudaMemcpyDeviceToHost);?? ????if ?(cudaStatus?!=?cudaSuccess)??? ????{?? ????????fprintf(stderr,?"cudaMemcpy?failed!" );?? ????????goto ?Error;?? ????}?? Error:?? ????cudaFree(dev_c);?? ????cudaFree(dev_a);?? ????cudaFree(dev_b);?????? ????return ?cudaStatus;?? }??
線程并行：
塊并行：
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
總結
以上是生活随笔 為你收集整理的CUDA5.5入门文章：VS10设置 的全部內容，希望文章能夠幫你解決所遇到的問題。
如果覺得生活随笔 網站內容還不錯，歡迎將生活随笔 推薦給好友。
