cuda學(xué)習(xí)筆記

2019-11-08 02:27:11

字體：大中小

來源：轉(zhuǎn)載

供稿：網(wǎng)友

cuda學(xué)習(xí)筆記

標(biāo)簽（空格分隔）：學(xué)習(xí)筆記

一、學(xué)習(xí)平臺(tái)

1.1學(xué)習(xí)平臺(tái)搭建與學(xué)習(xí)平臺(tái)基本信息

1.1.1學(xué)習(xí)平臺(tái)搭建

此篇文檔的整理基于nvidia公司出品的GeForce GTX 950 GPU，在電腦主機(jī)當(dāng)中安裝好獨(dú)立顯卡之后，安裝cuda7.0至軟件盤（不用再單獨(dú)安裝顯卡驅(qū)動(dòng)程序）。在vs下新建cuda工程，就可以編寫cuda程序了。

1.1.2學(xué)習(xí)平臺(tái)基本信息

在編寫cuda程序時(shí)，程序的頭文件應(yīng)該包括 “cuda_runtime.h”和”device_launch_parameters.h”；以下一段代碼用來查看顯卡gpu的計(jì)算性能和架構(gòu)

int main(){ cudaDevicePRop prop; int count; cudaGetDeviceCount(&count); for (int i = 0; i < count; ++i){ cudaGetDeviceProperties(&prop, i); printf(" --- Genaral Information for Device %d ---/n", i); printf("Name : %s/n",prop.name); printf("Compute capability : %d.%d/n", prop.major, prop.minor); printf("Clock rate: %d/n",prop.clockRate); printf("Device copy overlap: "); if (prop.deviceOverlap){ printf("Enabled/n"); } else{ printf("Disabled/n"); } printf("Kernel execition timeout : "); if (prop.kernelExecTimeoutEnabled) printf("Enabled/n"); else printf("Disabled/n"); printf(" ---Memory Information for Device %d ---/n",i); printf("Total global mem: %ld/n",prop.totalGlobalMem); printf("Total const mem : %ld/n", prop.totalConstMem); printf("Max mem pitch : %ld/n", prop.memPitch); printf("Texture Alignment : %ld/n",prop.textureAlignment); printf(" ---MP Information for device %d ---/n", i); printf("Multiprocessor count : %d/n", prop.multiProcessorCount); printf("shared mem per mp: %d/n", prop.sharedMemPerBlock); printf("Register per mp: %d/n",prop.regsPerBlock); printf("Threads in warp: %d/n", prop.warpSize); printf(" Max threads per block :%d/n", prop.maxThreadsPerBlock); printf("Max thread dimentions : (%d, %d, %d)/n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]); printf("Max grid dimensions:(%d, %d, %d)/n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2] ); printf("/n/n/n"); }}

950顯卡運(yùn)行結(jié)果如下圖所示：

此處輸入圖片的描述

二、cuda線程模型

2.1線程索引相關(guān)問題

我們把在GPU上啟動(dòng)的線程塊集合稱為一個(gè)線程格。從名字的含義可以看出，線程格既可以使一維的線程塊集合，也可以是二維的線程塊集合。核函數(shù)的每個(gè)副本都可以通過內(nèi)置變量blockIdx來判斷哪個(gè)線程塊正在執(zhí)行它。同樣，它還可以通過內(nèi)置變量gridDim來獲得線程塊的大小。通過這兩個(gè)變量來計(jì)算每個(gè)線程塊需要的數(shù)據(jù)索引。則當(dāng)前線程塊的索引如下：線程塊的索引=行索引*線程格的數(shù)目+列索引 blockIdx.y * gridDim.x+blockIdx.x; 同樣的：線程索引 = 行索引*線程塊的數(shù)目+列索引 threadIdx.y * blockIdx.x + threadIdx.x; int offset = x + y * Dim;在這里Dim表示線程塊的大小（也就是線程的數(shù)量），y為線程塊索引，并且x為線程塊中的線程索引，所以計(jì)算得到如下索引:

int tid = blockDim.x*blockIdx.x + threadIdx.x;tid += blockDim.x*gridDim.x;//每個(gè)線程塊中的數(shù)量乘以線程格中線程塊的總數(shù)量，即為當(dāng)前線程格中運(yùn)行的線程總數(shù)量。

對(duì)于二維線程的索引，有如下代碼：

int x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;

2.2線程開辟

一種解決方案是將線程塊的大小設(shè)置為某個(gè)固定數(shù)值BLOCKSIZE，然后啟動(dòng)N/BLOCKSIZE個(gè)線程塊，這樣就相當(dāng)啟動(dòng)了N個(gè)線程同時(shí)運(yùn)行。通常我們?cè)O(shè)置的線程塊的個(gè)數(shù)為（N+BLOCKSIZE-1）/BLOCKSIZE來防止0線程的開辟問題。

2.3多維數(shù)據(jù)線程調(diào)用

線性存儲(chǔ)器也可以通過cudaMallocPitch（）和cudaMalloc3D分配。在分配二維和三維數(shù)組的時(shí)候，推薦使用，因?yàn)樯鲜稣{(diào)用保證了GPU的最佳性能。返回的（pitch,stride）必須用于訪問數(shù)組元素。下面的代碼分配了一個(gè)尺寸為weight*height的二維浮點(diǎn)數(shù)組，同時(shí)演示了怎么在設(shè)備代碼中遍歷數(shù)組元素

//host codeint width =64,height = 64;float *dexPtr;int pitch;cudaMallocPitch((void **)&devPtr,&pitch,width*sizeof(float),height);kernel<<<100,512>>>(devPtr,pitch,widtf,height);//device code__global__void kernel(float* devPtr,int pitch,int width,int height){for(int i =0;i<helght;++i){float* row =(float*) ((char*)devPtr+i*pitch);for(int j = 0;j<width;++j){float element = row[i]; } }}

下面的代碼演示分配一個(gè)尺寸為width*height*depth的三維浮點(diǎn)數(shù)組，同時(shí)演示了怎么在設(shè)備代碼中遍歷數(shù)組元素。

//host codecudaPitchedPtr devPitchedPtr;cudaExtent extent = make_cudaExtent(64,64,64);cudaMalloc3D(&devPitchedPtr,extent);kernel<<<100,512>>>(devPitchedPtr,extent);//device code__global__ void kernel(cudaPitchedPtr devPitchedPtr,cudaExtent extent){char *devPtr=devPitchedPtr.ptr;size_t pitch =devPitchedPtr.pitch;size_t slicePitch = pitch*extent.height;for(int i=0;i<extent.depth;++i){char *slice = devPtr + i*slicePitch;for(int j=0;j<extent.height;++j){float *row=(float*)(slice + y*pitch);for(int x = 0;x<extent.width;++x){float element = row[x]; } }}

三、矩陣乘法

3.1使用全局內(nèi)存

#include<iostream>#include<fstream>using namespace std;#define AROWS 16#define ACOLS 4#define BROWS 4#define BCOLS 16void randomInit(float* _data, int _size){ for (int i = 0; i < _size; ++i) { _data[i] = rand() %5; //_data[i] = (float)(rand() / (float)RAND_MAX); }}__global__ void MatMul(float *dev_a, float *dev_b, float *dev_c){ int col = threadIdx.x; int row = threadIdx.y; float cvalue = 0; for (int i = 0; i<ACOLS; i++) { cvalue += dev_a[row*ACOLS+i] * dev_b[BCOLS*i + col]; } dev_c[row*BCOLS + col] = cvalue;}int main(){ float *a, *b, *c, *dev_a, *dev_b, *dev_c; a = (float*)malloc(sizeof(float)*AROWS*ACOLS); b = (float*)malloc(sizeof(float)*BROWS*BCOLS); c = (float*)malloc(sizeof(float)*AROWS*BCOLS); memset(c, 0, sizeof(float)*AROWS*BCOLS); cudaMalloc((void **)&dev_a, sizeof(float)*AROWS*ACOLS); cudaMalloc((void **)&dev_b, sizeof(float)*BROWS*BCOLS); cudaMalloc((void **)&dev_c, sizeof(float)*AROWS*BCOLS); cudaMemset(dev_c, 0, sizeof(float)*AROWS*BCOLS); //給矩陣AB賦初值 randomInit(a, AROWS*ACOLS); randomInit(b, BROWS*BCOLS); for (int i = 0; i < AROWS*ACOLS; i++) { cout << "a[" << i << "]=" << a[i] << endl; } for (int i = 0; i < BROWS*BCOLS; i++) { cout << "b[" << i << "]=" << b[i] << endl; } cudaMemcpy(dev_a, a, sizeof(float)*AROWS*ACOLS, cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, sizeof(float)*BROWS*BCOLS, cudaMemcpyHostToDevice); dim3 blocks(16, 16); dim3 threads(AROWS+15, BCOLS+15); MatMul << <blocks, threads >> >(dev_a, dev_b, dev_c); cudaMemcpy(c, dev_c, sizeof(float)*AROWS*BCOLS, cudaMemcpyDeviceToHost); //打印C矩陣到文件中 ofstream outFile; outFile.open("result.txt"); outFile << fixed; outFile.precision(2); outFile.setf(ios_base::showpoint); for (int i = 0; i < AROWS*BCOLS; i++) { outFile << "C[" << i << "]=" << c[i] << endl; } system("pause"); cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); free(a); free(b); free(c); return 0;}

使用全局內(nèi)存實(shí)現(xiàn)矩陣的一維向量乘法如上程序所示，但這種實(shí)現(xiàn)方式并沒有充分利用gpu的優(yōu)勢(shì)。下面的代碼是使用共享內(nèi)存實(shí)現(xiàn)矩陣乘法。

3.2使用共享內(nèi)存

四、關(guān)于cuda的基本架構(gòu)

在CUDA架構(gòu)下，線程的最小單元是thread，多個(gè)thread組成一個(gè)block，多個(gè)block再組成一個(gè)grid。每一個(gè)block中開辟的所有線程共享同一個(gè)shared memory。block里面的thread之間的通信和同步所帶來的開銷是比較大的。SM以 32 個(gè) Thread 為一組的 Warp 來執(zhí)行 Thread。Warp內(nèi)的線程是靜態(tài)的，即在屬于同一個(gè)warp內(nèi)的thread之間進(jìn)行通信，不需要進(jìn)行柵欄同步(barrier)。

4.1共享內(nèi)存

每個(gè)block中開辟的所有線程共享一個(gè)shared memory。使用共享內(nèi)存變量的時(shí)候，需要在聲明的時(shí)候加上shared關(guān)鍵詞修飾，使用共享內(nèi)存的時(shí)候應(yīng)注意同一個(gè)線程塊中的線程都執(zhí)行結(jié)束才能進(jìn)行下一步操作，所以需要使用__syncthread關(guān)鍵詞使得block中的線程同步。

4.2常量?jī)?nèi)存

常量?jī)?nèi)存用于保存在核函數(shù)執(zhí)行期間不會(huì)發(fā)生變化的數(shù)據(jù)，定義常量?jī)?nèi)存的時(shí)候應(yīng)該使用關(guān)鍵詞constant進(jìn)行修飾。當(dāng)從主機(jī)內(nèi)存復(fù)制到GPU上的常量?jī)?nèi)存時(shí)，需要使用cudaMemcpyToSymbol()復(fù)制數(shù)據(jù)。常量?jī)?nèi)存讀取數(shù)據(jù)可以節(jié)約內(nèi)存帶寬。注：只要當(dāng)一個(gè)warp的半線程束中的所有16個(gè)線程有相同的讀取請(qǐng)求時(shí)，才值得使用常量?jī)?nèi)存。

4.3紋理內(nèi)存

紋理內(nèi)存是專門為那些在內(nèi)存訪問模式中存在大量空間局部性的圖形應(yīng)用程序而設(shè)計(jì)的。以下內(nèi)容參考博文http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015755.html

4.3.1、概述

紋理存儲(chǔ)器中的數(shù)據(jù)以一維、二維或者三維數(shù)組的形式存儲(chǔ)在顯存中，可以通過緩存加速訪問，并且可以聲明大小比常數(shù)存儲(chǔ)器要大的多。在kernel中訪問紋理存儲(chǔ)器的操作稱為紋理拾取(texture fetching)。將顯存中的數(shù)據(jù)與紋理參照系關(guān)聯(lián)的操作，稱為將數(shù)據(jù)與紋理綁定(texture binding). 顯存中可以綁定到紋理的數(shù)據(jù)有兩種，分別是普通的線性存儲(chǔ)器和cuda數(shù)組。注：線性存儲(chǔ)器只能與一維或二維紋理綁定,采用整型紋理拾取坐標(biāo)，坐標(biāo)值與數(shù)據(jù)在存儲(chǔ)器中的位置相同； CUDA數(shù)組可以與一維、二維、三維紋理綁定，紋理拾取坐標(biāo)為歸一化或者非歸一化的浮點(diǎn)型，并且支持許多特殊功能。

4.3.2、紋理緩存：

（1）、紋理緩存中的數(shù)據(jù)可以被重復(fù)利用（2）、紋理緩存一次預(yù)取拾取坐標(biāo)對(duì)應(yīng)位置附近的幾個(gè)象元，可以實(shí)現(xiàn)濾波模式。

4.3.3紋理存儲(chǔ)器的使用：

使用紋理存儲(chǔ)器時(shí)，首先要在主機(jī)端聲明要綁定到紋理的線性存儲(chǔ)器或CUDA數(shù)組 (1)聲明紋理參考系

texture<Type, Dim, ReadMode> texRef;//Type指定數(shù)據(jù)類型，特別注意：不支持3元組//Dim指定紋理參考系的維度，默認(rèn)為1//ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默認(rèn))

注：紋理參照系必須定義在所有函數(shù)體外 (2) 聲明CUDA數(shù)組，分配空間 CUDA數(shù)組可以通過cudaMalloc3DArray()或者cudaMallocArray()函數(shù)分配。前者可以分配1D、2D、3D的數(shù)組，后者一般用于分配2D的CUDA數(shù)組。使用完畢，要用cudaFreeArray()函數(shù)釋放顯存。

//1數(shù)組 cudaMalloc((void**)&dev_A, data_size); cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice); cudaFree(dev_A); //2維數(shù)組 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>() cudaArray *cuArray; cudaMallocArray(&cuArray, &channelDesc, 64, 32); //64x32 cudaMemcpyToArray(cuArray, 0, 0, h_data, sizeof(float)*width*height, cudaMemcpyHostToDevice); cudaFreeArray(cuArray); //3維數(shù)組 64x32x16 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>(); cudaArray *d_volumeArray; cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumSize); cudaMemcpy3DParms copyParams = {0}; copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height); copyParams.dstArray = d_volumeArray; copyParams.extent = volumeSize; copyParams.kind = cudaMemcpyHostToDevice; cudaMemcpy3D(&copyParams); tex.normalized = true; tex.filterMode = cudaFilterModeLinear; tex.addressMode[0] = cudaAddressModeWrap; tex.addressMode[1] = cudaAddressModeWrap; tex.addressMode[2] = cudaAddressModeWrap;

（3）設(shè)置運(yùn)行時(shí)紋理參照系屬性

struct textureReference{ int normalized; enum cudaTextureFilterMode filterMode; enum cudaTextureAddressMode addressMode[3]; struct cudaChannelFormatDesc channelDesc;}

normalized設(shè)置是否對(duì)紋理坐標(biāo)歸一化 filterMode用于設(shè)置紋理的濾波模式 addressMode說明了尋址方式

(4)紋理綁定通過cudaBindTexture() 或 cudaBindTextureToArray()將數(shù)據(jù)與紋理綁定。通過cudaUnbindTexture()用于解除紋理參照系的綁定注：與紋理綁定的數(shù)據(jù)的類型必須與聲明紋理參照系時(shí)的參數(shù)匹配 (I).cudaBindTexture() //將1維線性內(nèi)存綁定到1維紋理

cudaError_t cudaBindTexture( size_t * offset, const struct textureReference * texref, const void * devPtr, const struct cudaChannelFormatDesc * desc, size_t size = UINT_MAX )

(II).cudaBindTexture2D //將1維線性內(nèi)存綁定到2維紋理

cudaError_t cudaBindTexture2D( size_t * offset, const struct textureReference * texref, const void * devPtr, const struct cudaChannelFormatDesc * desc, size_t width, size_t height, size_t pitch )

(III). cudaBindTextureToArray() //將cuda數(shù)組綁定到紋理

cudaError_t cudaBindTextureToArray ( const struct textureReference * texref, const struct cudaArray * array, const struct cudaChannelFormatDesc * desc )

(5)紋理拾取　　對(duì)于線性存儲(chǔ)器綁定的紋理，使用tex1Dfetch()訪問，采用的紋理坐標(biāo)是整型。由cudaMallocPitch() 或者 cudaMalloc3D()分配的線性空間實(shí)際上仍然是經(jīng)過填充、對(duì)齊的一維線性空　　　間，因此也用tex1Dfetch() 　　對(duì)與一維、二維、三維cuda數(shù)組綁定的紋理，分別使用tex1D(), tex2D() 和 tex3D()函數(shù)訪問，并且使用浮點(diǎn)型紋理坐標(biāo)。

4.4cuda流

五、使用事件對(duì)cuda程序進(jìn)行計(jì)時(shí)

創(chuàng)建事件

cudaEvent_t start,stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start,0);

……..此處省略cuda程序

事件結(jié)束

cudaEventRecord(stop,0);cudaEventSynchronize(stop);float elapsedTime;cudaEventElapsedTime(&elapsedTime,start,stop);printf("Time to generate : %3.1f ms/n",elapsedTime);cudaEventDestroy(start);cudaEventDestroy(stop);