如果之前沒有用過gdb, 可以速學一下, 就幾個指令. 想要用cuda-gdb對程式進行除錯, 首先你要確保你的gpu沒有在執行作業系統介面, 比方說, 我用的是ubuntu, 我就需要用sudo service lightdm stop關閉圖形介面, 進入tty1這種字元介面. 當然用ssh遠端訪問也是可以的. 接下來, 使用第二篇中矩陣加法的例子. 但是注意, 編譯的使用需要改變一下, 加入**-g -g**引數, 其實和gdb是相似的.
nvcc -g -g cudaadd.cu -o cudaadd.o
複製**
然後使用cuda-gdb cudaadd.o即可對程式進行除錯.
在除錯之前, 我把**貼出來:
#include
__global__ void add(float * x, float *y, float * z, int n)
}int main
() dim3 blocksize(256);
// 4096
dim3 gridsize((n + blocksize.x - 1) / blocksize.x);
add << < gridsize, blocksize >> >(x, y, z, n);
cudadevicesynchronize();
float maxerror = 0.0;
for (int i = 0; i < n; i++)
printf ("max default: %.4f\n", maxerror);
cudafree(x);
cudafree(y);
cudafree(z);
return 0;
}複製**
之後就是常規操作了, 新增斷點, 執行, 下一步, 檢視想看的資料. 不同點是cuda的指令, 例如cuda block(1,0,0)可以從一開始block(0,0,0)切換到block(1,0,0).
如果按照常規的思路, 兩兩進行進行加法運算. 每次步長翻倍即可, 從演算法的角度來說, 這是沒啥問題的. 但是沒有依照gpu架構進行設計.
#include
const int threadsperblock = 512;
const int n = 2048;
const int blockspergrid = (n + threadsperblock - 1) / threadsperblock; /* 4 */
__global__ void reductionsum( float * d_a, float * d_partial_sum )
/* 將當前block的計算結果寫回輸出陣列 */
if ( tid == 0 )
d_partial_sum[blockidx.x] = partialsum[0];
}int main
()複製**
其實需要改動的地方非常小, 改變步長即可.
__global__ void reductionsum( float * d_a, float * d_partial_sum )
// 相同, 略去
}複製**
之前的文章裡面也說過warp. warp: gpu執行程式時的排程單位, 目前cuda的warp的大小為32, 同在乙個warp的執行緒, 以不同資料資源執行相同的指令, 這就是所謂simt. 說人話就是, 這32個執行緒必須要幹相同的事情, 如果有執行緒動作不一致, 就需要等待一波執行緒完成自己的工作, 然後再去做另外一件事情. 所以, 用圖說話就是, 第二種方案可以更快將warp閒置, 交給gpu排程, 所以, 肯定是第二種更快.
圖一在運算依次之後, 沒有warp可以空閒, 而圖二直接空閒2個warp. 圖一到了第二次可以空閒2個warp, 而圖二已經空閒3個warp. 我這副圖只是示意圖, 如果是實際的, 差距會更大.
所以來看下執行耗時, 會發現差距還是很大的, 幾乎是差了一倍. 不過gpu確實算力太猛, 這樣看還不太明顯, 有意放大資料量會更加明顯.
GPU 程式設計入門到精通（四）之 GPU 程式優化
gpu 程式設計入門到精通 三 之 第乙個 gpu 程式 中講到了怎樣利用 cuda5.5 在 gpu 中執行乙個程式。通過程式的執行。我們看到了 gpu 確實能夠作為乙個運算器。可是，我們在前面的樣例中並沒有正真的發揮 gpu 並行處理程式的能力。也就是說之前的樣例僅僅利用了 gpu 的乙個執行緒...
GPU程式設計優化筆記
參考 硬體上，gpu由多個sm steaming multiprocessor 構成，sm有多個warp，warp有多個sp streaming processor 乙個sp對應乙個執行緒。乙個warp中的sp執行相同的指令。block內部可以使用sm提供的shared memory和 syncth...
MIC程式設計優化（1） 並行度優化
在計算機體系結構中，並行度是指指令並行執行的最大條數。在設計並行程式時，我們可以簡單地把並行度認為是在多核 眾核處理器上能同時執行的執行緒數 程序數。對於同乙個程式，並行度設計方法的不同將會嚴重影響到程式的效能。mic上的並行度優化主要涉及並行執行緒 程序的數目 並行層級 並行粒度等方面。1 並行度...