博主由于工作當(dāng)中的需要，開始學(xué)習(xí) GPU 上面的編程，主要涉及到的是基于 GPU 的深度學(xué)習(xí)方面的知識，鑒于之前沒有接觸過 GPU 編程，因此在這里特地學(xué)習(xí)一下 GPU 上面的編程。有志同道合的小伙伴，歡迎一起交流和學(xué)習(xí)，我的郵箱:caijinping220@gmail.com 。使用的是自己的老古董筆記本上面的 Geforce 103m 顯卡，雖然顯卡相對于現(xiàn)在主流的系列已經(jīng)非常的弱，但是對于學(xué)習(xí)來說，還是可以用的。本系列博文也遵從由簡單到復(fù)雜，記錄自己學(xué)習(xí)的過程。

---

0. 目錄

• GPU 編程入門到精通（一）之 CUDA 環(huán)境安裝
• GPU 編程入門到精通（二）之 運行第一個程序
• GPU 編程入門到精通（三）之 第一個 GPU 程序
• GPU 編程入門到精通（四）之 GPU 程序優(yōu)化
• GPU 編程入門到精通（五）之 GPU 程序優(yōu)化進階

1. CUDA 初始化函數(shù)

由于是使用 Runtime API, 所以在文件開頭要加入 cuda_runtime.h 頭文件。
初始化函數(shù)包括一下幾個步驟：

1.1. 獲取 CUDA 設(shè)備數(shù)

可以通過 cudaGetDeviceCount 函數(shù)獲取 CUDA 的設(shè)備數(shù)，具體用法，如下所示：

// get the cuda device count cudaGetDeviceCount(&count); if (count == 0) {fprintf(stderr, "There is no device.\n");return false; }

函數(shù)通過引用傳遞 count 值，獲取當(dāng)前支持的 CUDA 設(shè)備數(shù)。

1.2. 獲取 CUDA 設(shè)備屬性

可以通過 cudaGetDeviceProperties 函數(shù)獲取 CUDA 設(shè)備的屬性，具體用法，如下所示：

// find the device >= 1.X int i; for (i = 0; i < count; ++i) {cudaDeviceProp prop;if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {if (prop.major >= 1) {printDeviceProp(prop);break;}} }// if can't find the device if (i == count) {fprintf(stderr, "There is no device supporting CUDA 1.x.\n");return false; }

函數(shù)通過引用傳遞 prop 關(guān)于屬性的結(jié)構(gòu)體，并且列出主設(shè)備號大于 1 的設(shè)備屬性，其中設(shè)備屬性通過函數(shù) printDeviceProp 打印。打印函數(shù)如下所示：

// function printDeviceProp void printDeviceProp(const cudaDeviceProp &prop) {printf("Device Name : %s.\n", prop.name);printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);printf("regsPerBlock : %d.\n", prop.regsPerBlock);printf("warpSize : %d.\n", prop.warpSize);printf("memPitch : %d.\n", prop.memPitch);printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);printf("totalConstMem : %d.\n", prop.totalConstMem);printf("major.minor : %d.%d.\n", prop.major, prop.minor);printf("clockRate : %d.\n", prop.clockRate);printf("textureAlignment : %d.\n", prop.textureAlignment);printf("deviceOverlap : %d.\n", prop.deviceOverlap);printf("multiProcessorCount : %d.\n", prop.multiProcessorCount); }

1.3. 設(shè)置 CUDA 設(shè)備

通過函數(shù) cudaSetDevice 就可以設(shè)置 CUDA 設(shè)備了，具體用法，如下所示：

// set cuda device cudaSetDevice(i);

1.4. CUDA 初始化完整代碼

/* ******************************************************************* ##### File Name: first_cuda.cu ##### File Func: initial CUDA device and print device prop ##### Author: Caijinping ##### E-mail: caijinping220@gmail.com ##### Create Time: 2014-4-21 * ********************************************************************/#include <stdio.h> #include <cuda_runtime.h>void printDeviceProp(const cudaDeviceProp &prop) {printf("Device Name : %s.\n", prop.name);printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);printf("regsPerBlock : %d.\n", prop.regsPerBlock);printf("warpSize : %d.\n", prop.warpSize);printf("memPitch : %d.\n", prop.memPitch);printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);printf("totalConstMem : %d.\n", prop.totalConstMem);printf("major.minor : %d.%d.\n", prop.major, prop.minor);printf("clockRate : %d.\n", prop.clockRate);printf("textureAlignment : %d.\n", prop.textureAlignment);printf("deviceOverlap : %d.\n", prop.deviceOverlap);printf("multiProcessorCount : %d.\n", prop.multiProcessorCount); }bool InitCUDA() {//used to count the device numbersint count;// get the cuda device countcudaGetDeviceCount(&count);if (count == 0) {fprintf(stderr, "There is no device.\n");return false;}// find the device >= 1.Xint i;for (i = 0; i < count; ++i) {cudaDeviceProp prop;if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {if (prop.major >= 1) {printDeviceProp(prop);break;}}}// if can't find the deviceif (i == count) {fprintf(stderr, "There is no device supporting CUDA 1.x.\n");return false;}// set cuda devicecudaSetDevice(i);return true; }int main(int argc, char const *argv[]) {if (InitCUDA()) {printf("CUDA initialized.\n");}return 0; }

2. Runtime API 函數(shù)解析

2.1. cudaGetDeviceCount

cudaGetDeviceCount

返回具有計算能力的設(shè)備的數(shù)量

函數(shù)原型：

cudaError_t cudaGetDeviceCount( int* count )

函數(shù)說明：

以 *count 形式返回可用于執(zhí)行的計算能力大于等于 1.0 的設(shè)備數(shù)量。如果不存在此類設(shè)備，將返回 1

返回值：

cudaSuccess，注意，如果之前是異步啟動，該函數(shù)可能返回錯誤碼。

2.2. cudaGetDeviceProperties

cudaGetDeviceProperties

返回關(guān)于計算設(shè)備的信息

函數(shù)原型：

cudaError_t cudaGetDeviceProperties( struct cudaDeviceProp* prop,int dev )

函數(shù)說明：

以*prop形式返回設(shè)備dev的屬性。

返回值：

cudaSuccess、cudaErrorInvalidDevice，注，如果之前是異步啟動，該函數(shù)可能返回錯誤碼。

另外 cudaDeviceProp 結(jié)構(gòu)定義如下:

struct 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]; size_t totalConstMem; int major; int minor; int clockRate; size_t textureAlignment; int deviceOverlap; int multiProcessorCount; }

cudaDeviceProp 結(jié)構(gòu)中的各個變量意義如下：

name
用于標(biāo)識設(shè)備的ASCII字符串;
totalGlobalMem
設(shè)備上可用的全局存儲器的總量,以字節(jié)為單位;
sharedMemPerBlock
線程塊可以使用的共享存儲器的最大值,以字節(jié)為單位;多處理器上的所有線程塊可以同時共享這些存儲器;
regsPerBlock
線程塊可以使用的32位寄存器的最大值;多處理器上的所有線程塊可以同時共享這些寄存器;
warpSize
按線程計算的warp塊大小;
memPitch
允許通過cudaMallocPitch()為包含存儲器區(qū)域的存儲器復(fù)制函數(shù)分配的最大間距(pitch),以字節(jié)為單位;
maxThreadsPerBlock
每個塊中的最大線程數(shù)
maxThreadsDim[3]
塊各個維度的最大值:
maxGridSize[3]
網(wǎng)格各個維度的最大值;
totalConstMem
設(shè)備上可用的不變存儲器總量,以字節(jié)為單位;
major,minor
定義設(shè)備計算能力的主要修訂號和次要修訂號;
clockRate
以千赫為單位的時鐘頻率;
textureAlignment
對齊要求;與textureAlignment字節(jié)對齊的紋理基址無需對紋理取樣應(yīng)用偏移;
deviceOverlap
如果設(shè)備可在主機和設(shè)備之間并發(fā)復(fù)制存儲器,同時又能執(zhí)行內(nèi)核,則此值為 1;否則此值為 0;
multiProcessorCount
設(shè)備上多處理器的數(shù)量。

2.3. cudaGetDeviceCount

cudaSetDevice

設(shè)置設(shè)備以供GPU執(zhí)行使用

函數(shù)原型：

cudaError_t cudaSetDevice(int dev)

函數(shù)說明：

將dev記錄為活動主線程將執(zhí)行設(shè)備碼的設(shè)備。

返回值：

cudaSuccess、cudaErrorInvalidDevice，注，如果之前是異步啟動，該函數(shù)可能返回錯誤碼。

3. nvcc 編譯代碼

nvcc 是 CUDA 的編譯工具，它可以將 .cu 文件解析出在 GPU 和 host 上執(zhí)行的部分，也就是說，它會幫忙把 GPU 上執(zhí)行和主機上執(zhí)行的代碼區(qū)分開來，不許要我們手動去做了。在 GPU 執(zhí)行的部分會通過 NVIDIA 提供的 編譯器編譯成中介碼，主機執(zhí)行的部分則調(diào)用 gcc 編譯。

通過如下命令，可以編譯之前寫的 first_cuda.cu 程序：

nvcc -o first_cuda first_cuda.cu

通過上述編譯，生成可執(zhí)行文件 first_cuda

運行結(jié)果如下所示：

總結(jié)

以上是生活随笔為你收集整理的GPU 编程入门到精通（二）之 运行第一个程序的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。