cuda編程例子

① 如何進行CUDA C程序核函數的調試

由於CUDA調試工具的不完善、CUDA調試工具上手難度較高，並行思想本身就難調試等因素，CUDA調試一直都是一件很蛋疼的事情。寫CUDA也有三四年了，前段時間在群里見別人問CUDA調試的問題，突然有想法寫個CUDA調試的博客。自己經驗尚淺，希望各位大大看過後能夠在評論里指點一二，共同完善這篇博客。
本博客只針對邏輯bug。

1 定位bug

出現bug的第一想法自然是定位bug。cuda比較奇特的地方在於，有時報錯bug在500行，但500行出的代碼沒有錯誤，而是在1000行的地方邏輯錯了，十分頭疼。
下面介紹三種我總結的定位bug方法：

1.1 二分法

一半一半的注釋代碼，定位bug。比較笨拙和麻煩，但是十分好用。

1.2 輸出定位法

將整體代碼分為幾個模塊，正常的CUDA代碼大概可以分為數據返純初始化，內存申請，內存拷貝，核函數執行，結果拷貝等模塊。在每個模塊結束後輸出標志，示例如圖1。這樣在調試時就可以根據輸出快速定位bug大約在什麼位置。如下圖：

② CUDA編程基礎——Grid、Block、Thread

本消尺余文主要介紹三者之間的關系。

三者之間關系如圖所示，從中可以看出，三者存在包含關系。每個grid分為多個block，每困棗個block分為多個Thread，grid和block最多可拿滾以是三維的。

③ cuda程序無法傳數據到顯存

方法如下：
1、CPU和GPU之間
1）CPU->GPU
從CPU向GPU傳輸數據，最為人熟知的就是cudaMemcpy了。
默認情況下，數據是從系統的分頁內存先到鎖頁內存，然後再到GPU顯存。因此如果顯式指定使用鎖頁內存，是可以加快數據傳輸速度的。
（鎖頁內存，在cuda編程里使用CudaHostMalloc分配。實質上和linux的mlock系統調用一樣，就是給內存頁打上標記，不讓操作系統將其從物理內存交換到硬碟）
至於為什麼cuda要這樣設計，個人理解是為了實現的方便。因為操作系統已經處理了硬碟和物理內存間的頁交換等情況，顯卡驅動只需要實現物理內存到GPU顯存這一種數據傳輸即可，不需要把操作系統內存管理的事情再做一遍。
2) GPU->CPU
GPU向CPU拷貝數據時，鎖頁內存同樣比分頁內存快
值得一提的是，適當使用pinned memory顯然可以加快IO速度。但是並不是越多越好，因為鎖頁內存是完全獨佔住了物理內存，操作系統無法調度，可能會影響系統整體性能。
3)同一張GPU卡內部
同一張卡內兩塊顯存對拷，實測P40上高達~285GB/s。也比較接近於GPU卡本身的訪存速度
4）數據拷貝的overhead
在上面的測試數據中，可以看到傳輸數據量從1M->32M增長的過程中，測得的傳輸帶寬是有逐漸增加的。
這是因為每次調用cuda api進行數據傳輸都有overhead，在數據量小的時候這個overhead在數據傳輸時間中的佔比就顯得很高。這也提示我們盡量合並小數據的傳輸
2、同機的GPU之間
一般可以通過cudaMemcpyPeer/cudaMemcpyPeerAsync函數進行顯存拷貝
1）cudaMemcpyPeer withoutP2P
/********代碼示例*******/

cudaSetDevice(1);

cudaMalloc((int**)&dest, bytes);

cudaSetDevice(2);

cudaMalloc((int**)&dsrc, bytes);

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes);
通過nvprof+nvpp可以看到：禁用GPU P2P時，數據是先從GPU2拷貝到系統內存（DtoH），然後再從系統內存拷貝到GPU1（HtoD）

當然，這里是在一個進程內做GPU之間的數據拷貝。如果是2個進程分別運行在GPU1和GPU2上，那在CPU上這2個進程間可以通過共享內存或者socket通信來完成數據的拷貝。

2）cudaMemcpyPeer withP2P
/********代碼示例*******/

cudaSetDevice(1);

cudaMalloc((int**)&dest, bytes);

cudaSetDevice(2);

cudaMalloc((int**)&dsrc, bytes);

cudaDeviceEnablePeerAccess(1,0);

cudaDeviceEnablePeerAccess(2,0);

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes);
啟用GPU P2P時，數據直接從GPU2拷貝到了GPU1，不再經過系統內存。

3)通過變數賦值方式傳輸數據
深度學習中，卡之間傳遞的數據其實很多都是參數數值，因此也可以直接用一個GPU內的變數給另一個GPU上的變數賦值來進行數據傳輸

/********代碼示例*******/

(&numBlocks, &blockSize, p2p_float);

p2p_float<<>>(

(float *)dest, (float *)src, num_elems);

__global__ void p2p_float(float *__restrict__ dest, float const *__restrict__ src,

size_t num_elems) {undefined

size_t globalId = blockIdx.x * blockDim.x + threadIdx.x;

size_t gridSize = blockDim.x * gridDim.x;

#pragma unroll(5)

for (size_t i = globalId; i < num_elems; i += gridSize) {undefined

dest[i] = src[i];

}

}

④ 【CUDA 編程】bank 與bank沖突

以下內容摘笑塌抄於Nvida 官方教程碰則圓 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-hierarchy

Shared memory has 32 banks that are organized such that successive 32-bit words map to successive banks. Each bank has a bandwidth of 32 bits per clock cycle.
共享內存由連續的32bit單元映射到32個bank， 每個時鍾周期內， 每個bank都有32bit的帶寬

A shared memory request for a warp does not generate a bank conflict between two threads that access any address within the same 32-bit word (even though the two addresses fall in the same bank). In that case, for read accesses, the word is broadcast to the requesting threads and for write accesses, each address is written by only one of the threads (which thread performs the write is undefined).
一個wrap內的兩個線程訪問同一個在共享內存中的32bit數據， 這樣並不會引起bank沖突 （怎麼感覺和下面的圖middle的描述不一樣？？）盯橋 。對於讀操作， 32bit數據會被廣播給請求的線程； 對於寫線程， 數據僅僅會被一個線程寫(這樣會產生未定義的行為)。

Figure 17 shows some examples of strided access.
Figure 18 shows some examples of memory read accesses that involve the broadcast mechanism.
Figure 17. Strided Shared Memory Accesses. Examples for devices of compute capability 3.x (in 32-bit mode) or compute capability 5.x and 6.x

Left Linear addressing with a stride of one 32-bit word (no bank conflict).
Middle Linear addressing with a stride of two 32-bit words (two-way bank conflict).
Right Linear addressing with a stride of three 32-bit words (no bank conflict).
Figure 18. Irregular Shared Memory Accesses. Examples for devices of compute capability 3.x, 5.x, or 6.x.

Left Conflict-free access via random permutation.
Middle Conflict-free access since threads 3, 4, 6, 7, and 9 access the same word within bank 5.
Right Conflict-free broadcast access (threads access the same word within a bank).

⑤ c語言 gpu

U越來越強大，GPU為顯示圖像做了優化之外，在計算上已經超越了通用的CPU。如此強大的晶元如果只是作為顯卡就太浪費了，因此NVidia推出CUDA，讓顯卡可以用於圖像計算以外的目的，也就是超於游戲，使得GPU能夠發揮其強大的運算能力。

一年前NVIDIA發布CUDA，這是一種專門針對GPU的C語言開發工具。與以往採用圖形API介面指揮GPU完成各種運算處理功能不同，CUDA的出現使研究人員和工程師可以在熟悉的C語言環境下，自由地輸入代碼調用GPU的並行處理架構。這使得原先需要花費數天數周才能出結果的運算大大縮短到數幾小時，甚至幾分鍾之內。

CUDA是用於GPU計算的開發環境，它是一個全新的軟硬體架構，可以將GPU視為一個並行數據計算的設備，對所進行的計算進行分配和管理。在CUDA的架構中，這些計算不再像過去所謂的GPGPU架構那樣必須將計算映射到圖形API（OpenGL和Direct 3D）中，因此對於開發者來說，CUDA的開發門檻大大降低了。CUDA的GPU編程語言基於標準的C語言，因此任何有C語言基礎的用戶都很容易地開發CUDA的應用程序。

那麼，如何使得CPU與GPU之間很好的進行程序之間的銜接呢？以GPGPU的概念來看，顯卡仍然需要以傳統的DirectX和OpenGL這樣的API來實現，對於編程人員來說，這樣的方法非常繁瑣，而CUDA正是以GPGPU這個概念衍生而來的新的應用程序介面，不過CUDA則提供了一個更加簡便的方案——C語言。我們回顧一下CUDA的發展歷史。

⑥ 如何使用CUDA 顯卡編程

cuda是利用gpu編程。你需要先去下一個visual
studio,然後去cuda官網下一個現在版本的cuda7.0。全部安裝好後就可以編程了。cuda是並行編程，利用線程組織架構。有很多教程，你可以去看看。最好的一本書是cuda
by
example。

⑦ 如何學習cuda c

1、CUDAC編寫WindowsConsoleApplication

下面我們從一個簡單的例子開始學習CUDAC。

打開VS，新建一個CUDAWinApp項目，項目名稱為Vector，解決方案名稱為CUDADemo。依次點擊「確定」，「下一步」，選擇Emptyproject。點擊「Finished」。這樣一個CUDA的項目就建成了。

右鍵點擊Vector項目，依次選擇「添加」、「新建項」、「代碼」、「CUDA」。在名稱中輸入要添加的拿唯則文件名。如Vector.cu。然後點擊添加。

下面在Vector.cu文件里實現兩個向量相加的程序。

//添加系統庫

#include

#include

//添加CUDA支持

#include

__global__voidVecAdd(float*A,float*B,float*C);

__host__voidrunVecAdd(intargc,char**argv);

intmain(intargc,char**argv)

{

runVecAdd(argc,argv);

CUT_EXIT(argc,argv);

}

__host__voidrunVecAdd(intargc,char**argv)

{//初始化host端內存數據

constunsignedintN=8;//向量維數

constunsignedintmemSize=sizeof(float)*N;//需要空間的位元組數

float*h_A=(float*)malloc(memSize);

float*h_B=(float*)malloc(memSize);

float*h_C=(float*)malloc(memSize);

for(unsignedinti=0;i<N;i++)

{h_A[i]=i;h_B[i]山唯=i;}

//設備端顯存空間

float*d_A,*d_B,*d_C;

//初始化Device

CUT_DEVICE_INIT(argc,argv);

CUDA_SAFE_CALL(cudaMalloc((void**)&d_A,memSize));

CUDA_SAFE_CALL(cudaMalloc((void**)&d_B,memSize));

CUDA_SAFE_CALL(cudaMalloc((void**)&d_C,memSize));

CUDA_SAFE_CALL(cudaMemcpy(d_A,h_A,memSize,cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpy(d_B,h_B,memSize,cudaMemcpyHostToDevice));

VecAdd<<<1,N,memSize>>>(d_A,d_B,d_C);

CUT_CHECK_ERROR("Kernelexecutionfailed");

CUDA_SAFE_CALL(cudaMemcpy(h_C,d_C,memSize,cudaMemcpyDeviceToHost));

for(unsignedinti=0;i<N;i++)

{printf("%.0f",h_C[i]);}

free(h_A);free(h_B);free(h_C);

CUDA_SAFE_CALL(cudaFree(d_A));

CUDA_SAFE_CALL(cudaFree(d_B));

CUDA_SAFE_CALL(cudaFree(d_C));

}

__global__voidVecAdd(float*A,float*B,float*C)

{

//分配sharedmemory

extern__shared__floats_A[];

extern__shared__floats_B[];

extern__shared__floats_C[];

//從globalmemory拷貝到sharedmemory

constunsignedinti=threadIdx.x;

s_A[i]=A[i];

s_B[i]=B[i];

//計算

s_C[i]=s_A[i]+s_B[i];

//拷貝到globalmemory

C[i]=s_C[i];

}

由於這里不是講CUDA編程的，關於它的編程模型已經超出了我要介紹的范圍，您可以閱讀消棚《GPU高性能運算之CUDA》來獲得CUDA編程模型的知識。

編譯Vector項目，執行此項目後會得到圖1如下輸出：

圖1Vector項目執行結果

2、CUDAC編寫DLL模塊

更多情況下的您的軟體可能只是使用CUDA來實現一段程序的加速，這種情況下我們可以使用CUDAC編寫DLL來提供介面。下面我們就將例1編譯成DLL。

在剛才的CUDADemo解決方案目錄下添加一個新的CUDA項目（當然您也可以重新建立一個解決方案）。項目名為VecAdd_dynamic。ApplicationType選為DLL，AdditionalOptions選擇EmptyProject。

第一步，添加頭文件，文件名最好與工程名同名，這樣便於您的維護工作。這里我向項目中添加了VecAdd_dynamic.h，在此頭文件中添加如下代碼

#ifndef_VECADD_DYNAMIC_H_

#define_VECADD_DYNAMIC_H_

//並行計算N維向量的加法

__declspec(dllexport)voidVecAdd(float*h_A,float*h_B,float*h_C,intN);

#endif

第二步，添加cpp文件，文件名為VecAdd_dynamic.cpp，在此文件中添加如下代碼

#include

#include"VecAdd_dynamic.h"

#ifdef_MANAGED

#pragmamanaged(push,off)

#endif

BOOLAPIENTRYDllMain(HMODULEhMole,DWORDul_reason_for_call,LPVOIDlpReserved)

{

returnTRUE;

}

#ifdef_MANAGED

#pragmamanaged(pop)

#endif

第三步，添加def文件，此文件的功能就是確保其它廠商的編譯器能夠調用此DLL里的函數。這一點非常關鍵，因為您的程序可能用到多個廠家的編譯器。文件名為VecAdd_dynamic.def。向該文件中添加：

EXPORTS

VecAdd

第四步，添加cu文件，文件名為VecAdd_dynamic.cu。注意此文件最好直接添加到項目目錄下，不要添加到源文件選項卡或其它已有的選項卡下。

在cu文件里添加如下代碼，實現要導出的函數。

#include

#include

#include

#if__DEVICE_EMULATION__

boolInitCUDA(void)

{returntrue;}

#else

boolInitCUDA(void)

{

intcount=0;

inti=0;

cudaGetDeviceCount(&count);

if(count==0)

{

fprintf(stderr,"Thereisnodevice./n");

returnfalse;

}

for(i=0;i<count;i++)

{

cudaDevicePropprop;

if(cudaGetDeviceProperties(&prop,i)==cudaSuccess)

{

if(prop.major>=1)

{break;}

}

}

if(i==count)

{

fprintf(stderr,"ThereisnodevicesupportingCUDA./n");

returnfalse;

}

cudaSetDevice(i);

printf("CUDAinitialized./n");

returntrue;

}

#endif

__global__voidD_VecAdd(float*g_A,float*g_B,float*g_C,intN)

{

unsignedinti=threadIdx.x;

if(i<N)

{g_C[i]=g_A[i]+g_B[i];}

}

voidVecAdd(float*h_A,float*h_B,float*h_C,intN)

{

if(!InitCUDA())

{return;}

float*g_A,*g_B,*g_C;

unsignedintsize=N*sizeof(float);

CUDA_SAFE_CALL(cudaMalloc((void**)&g_A,size));

CUDA_SAFE_CALL(cudaMalloc((void**)&g_B,size));

CUDA_SAFE_CALL(cudaMalloc((void**)&g_C,size));

CUDA_SAFE_CALL(cudaMemcpy(g_A,h_A,size,cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaMemcpy(g_B,h_B,size,cudaMemcpyHostToDevice));

D_VecAdd<<<1,N>>>(g_A,g_B,g_C,N);

CUDA_SAFE_CALL(cudaMemcpy(h_C,g_C,size,cudaMemcpyDeviceToHost));

cudaFree(g_A);cudaFree(g_B);cudaFree(g_C);

}

第五步，如果您已經正確完成了以上四步，那麼剩下的就只有編譯，只要您用過VS，這一步就不需要我介紹了吧。成功之後，在您的解決方案文件目錄下的Debug文件夾下會有一個VecAdd_dynamic.dll文件。

3、在.NET中使用CUDAC編寫的DLL

下面介紹在託管程序中如何使用VecAdd_dynamic.dll。

第一步，在上面的解決方案CUDADemo下添加一個C++/CLR的Windows窗體應用程序，工程名為NETDemo（當然您也可以重新建一個解決方案，工程名也是隨意的）。

第二步，在窗體上添加一個按鈕，名字隨意，我將它的現實文本改為「調用CUDA_DLL」，給這個按鈕添加click事件。我們的代碼將在這個事件里添加調用VecAdd()的程序。在窗體上添加一個文本框用來顯示調用輸出的結果。

第三步，代碼實現。為工程NETDemo添加一個頭文件，我將它命名為Win32.h，這個文件中主要是實現VecAdd()函數的導入。在此文件中添加如下代碼

#pragmaonce

namespaceWin32

{

usingnamespaceSystem::Runtime::InteropServices;

[DllImport("VecAdd_dynamic.dll",EntryPoint="VecAdd",CharSet=CharSet::Auto)]

extern"C"voidVecAdd(float*h_A,float*h_B,float*h_C,intN);

}

在Form1.h中，#pragmaonce之後namespaceNETDemo之前添加以下代碼。

#include"Win32.h"

#include

在button1_Click()中添加如下代碼

intN=8;

float*h_A=(float*)malloc(N*sizeof(float));

float*h_B=(float*)malloc(N*sizeof(float));

float*h_C=(float*)malloc(N*sizeof(float));

for(inti=0;i<N;i++)

{h_A[i]=i;h_B[i]=i;}

Win32::VecAdd(h_A,h_B,h_C,N);

String^reslut;

for(inti=0;i<N;i++)

{reslut+=Convert::ToString(h_C[i])+",";}

this->textBox1->Text=Convert::ToString(reslut);

free(h_A);free(h_B);free(h_C);

第四步、執行NETDemo項目。點擊「調用CUDA_DLL」，您會看到圖3所示的結果

圖3NETDemo運行結果

到現在為止您已經完全可以正確使用CUDA了。