這一年都在編寫CUDA的程序，用了很多優(yōu)化的手段，發(fā)現(xiàn)大部分其實還是官方的指南里面的手段

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
https://developer.download.nvidia.cn/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

至于代碼，多看CUDA自帶的example就好了，挺好的代碼

學(xué)會使用trust庫

后面才發(fā)現(xiàn)有這個庫的，具體效率沒看過，但是發(fā)現(xiàn)很多我手寫的代碼trust里面都有，下次可以用用。

內(nèi)存方面

CUDA里面，用nsight一看，有大部分時間都花在cudaMalloc,cudaMemcpy這一類的函數(shù)里面，所以如何減少這些函數(shù)調(diào)用是很重要的

減少cudaMalloc,cudaMemcpy的調(diào)用

cudaMalloc是很耗時的一個操作，我在優(yōu)化雷達數(shù)據(jù)計算的時候，一份數(shù)據(jù)用CPU算要80ms，但是用GPU算只要10，如果用GTX1080ti的話也才6ms左右，但是單單cudaMalloc就用了10幾ms.特別在多線程下，這個問題特別嚴重
所以不要輕易回收申請的GPU內(nèi)存，能復(fù)用就復(fù)用，能實現(xiàn)一個簡單的內(nèi)存池最好

多線程同時操作內(nèi)存

void func() {char *p = nullptr;cudaMalloc((void**)&p, 1024 * 1024 * 10);cudaFree(p);}int main() {std::thread t(func);std::thread t2(func);std::thread t3(func);t.join();t2.join();t3.join(); }

上面的代碼同時開啟3個線程，申請了三塊顯存, 可以看到，每個malloc的申請都在100ms之間.我估計應(yīng)該是產(chǎn)生競爭了

后來寫代碼發(fā)現(xiàn)，是每個線程在第一次執(zhí)行CUDA的代碼的時候，會比較慢。修改了代碼，在代碼前面加入create stream，發(fā)現(xiàn)createstream比較慢，而cudaMalloc快了點，但是依然有4-10ms的消耗，所以我估計還是有一定的競爭
而每個線程第一次調(diào)用CUDA函數(shù)會很慢是因為CUDA的每個線程需要初始化一些東西吧，找了版本也沒找到相關(guān)的說明，有人說是context，但是從nsight里面看到，大部分的context都是一樣的

---

將代碼修改成串行申請內(nèi)存，可以看到內(nèi)存申請只有第一個比較耗時，第二個第三個都是很快的

void func() {char *p = nullptr;cudaMalloc((void**)&p, 1024 * 1024 * 10);cudaFree(p); }int main() {func();func();func(); }

拷貝內(nèi)存

static const int malloc_size = 1024 * 1024 * 10; void func() {cudaStream_t stream;cudaStreamCreate(&stream);char *gpu_mem = nullptr;char *host_mem = (char*)malloc(malloc_size);cudaMalloc((void**)&gpu_mem, malloc_size);cudaMemcpy(gpu_mem, host_mem, malloc_size, cudaMemcpyHostToDevice);cudaMemcpy(host_mem, gpu_mem, malloc_size, cudaMemcpyDeviceToHost);cudaFree(gpu_mem);free(host_mem); }

開多線程拷貝

單線程拷貝內(nèi)存

可以看到單線程的內(nèi)存拷貝比較穩(wěn)定，而多線程在2-10ms之間，不知道是不是因為多線程操作內(nèi)存引起的部分等待

網(wǎng)上說 可以使用stream進行異步傳輸，于是修改代碼，使用async的memcpy

cudaStream_t stream;cudaStreamCreate(&stream);char *gpu_mem = nullptr;char *host_mem = (char*)malloc(malloc_size);cudaMalloc((void**)&gpu_mem, malloc_size);cudaMemcpyAsync(gpu_mem, host_mem, malloc_size, cudaMemcpyHostToDevice,stream);cudaMemcpyAsync(host_mem, gpu_mem, malloc_size, cudaMemcpyDeviceToHost,stream);cudaFree(gpu_mem);free(host_mem);

最終結(jié)果，表示沒啥鳥用，速度慢還是慢，從幾張圖的開始時間和結(jié)束時間來看，也沒有看出什么異步的優(yōu)勢

pinned 內(nèi)存

申請pinned內(nèi)存耗時

拷貝pinned內(nèi)存與拷貝非pinned內(nèi)存對比

從上面圖中可以看出，拷貝pinn內(nèi)存基本是拷貝非pinn內(nèi)存速度的7-9倍左右，效率提升還是很可觀的

zero copy, shared memory

zero copy對我來講幾乎沒什么用的東西，太過頻繁在CPU和GPU寫和讀內(nèi)存反而導(dǎo)致速度下降得很嚴重，最終被放棄了

shared memory對我來講用處并不大，而且要修改太多的代碼。我們代碼里面主要時間還是花在GPU和CPU的交互，即使修改成使用shared memory最終估計也就優(yōu)化幾ms，沒必要那么折騰

IO優(yōu)化

IO優(yōu)化主要是在三方面，知乎上的答案講得很清楚了

gpu的io問題一直是比較嚴重的瓶頸，我自己總結(jié)有三種解決方案。
1.優(yōu)化算法，選擇更有效的存儲方式，減少顯存的使用。只拷貝必要數(shù)據(jù)，優(yōu)化數(shù)據(jù)結(jié)構(gòu)等等。
2.使用cuda內(nèi)存優(yōu)化技巧。比如顯存整塊申請和拷貝，使用pinned memory，shared memory,constant memory，使用流來掩蓋內(nèi)存延遲等等。
3.硬件解決方案，放棄pci-e，使用ibm的power架構(gòu)和p100，直接用nv-link進行內(nèi)存?zhèn)鬏敗＿@個方案是一勞永逸的，但是要準備好money，一整套下來很貴的。第三個方案一般整不起，前兩個方案是算法和語言技巧方面，與算法本身有很大的關(guān)系。個人經(jīng)驗上來講方案1提升的潛力比較大。
作者：楊偉光
鏈接：https://www.zhihu.com/question/52472621/answer/130949946
來源：知乎
著作權(quán)歸作者所有。商業(yè)轉(zhuǎn)載請聯(lián)系作者獲得授權(quán)，非商業(yè)轉(zhuǎn)載請注明出處。

https://www.cnblogs.com/1024incn/p/5891051.html

最重要的

• 學(xué)會使用nvprof工具，優(yōu)化這個才是關(guān)鍵
• 了解自己的程序是否適合做并行化，不了解自己的程序是否適合并行化的話，多讀幾本多線程相關(guān)的書籍，或者CPU內(nèi)存操作優(yōu)化的書籍,其實兩者的優(yōu)化思路都差不多.
• 注意一點，如果不熟悉CUDA的話，特別是那種執(zhí)行本來只需要幾百毫秒的算法，一部小心寫出來的代碼可能比CPU版本還慢，我們的雷達處理器就是一個很好的例子。最早要優(yōu)化運算，CPU版本是200ms，我同事寫出來的CUDA代碼居然要300ms,后來經(jīng)過我的優(yōu)化，最后是60ms，主要是內(nèi)存操作的優(yōu)化

轉(zhuǎn)載于:https://www.cnblogs.com/linyilong3/p/11061680.html

總結(jié)

以上是生活随笔為你收集整理的CUDA杂谈的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

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