CUDA編程2——共享內(nèi)存的優(yōu)勢

這里解決一個問題。通過兩個程序掌呜,討論共享內(nèi)存的優(yōu)勢滓玖。

共享內(nèi)存預(yù)計比全局內(nèi)存快得多。它可以用作暫存器內(nèi)存(或軟件托管的高速緩存)质蕉,以最大程度地減少來自CUDA塊的全局內(nèi)存訪問.

一 全局內(nèi)存


// Matrices are stored in row-major order:

// M(row, col) = *(M.elements + row * M.width + col)

typedef struct {

int width;

int height;

float* elements;

} Matrix;

// Thread block size

#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel

__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code

// Matrix dimensions are assumed to be multiples of BLOCK_SIZE

void MatMul(const Matrix A, const Matrix B, Matrix C)

{

// Load A and B to device memory

Matrix d_A;

d_A.width = A.width; d_A.height = A.height;

size_t size = A.width * A.height * sizeof(float);

cudaMalloc(&d_A.elements, size);

cudaMemcpy(d_A.elements, A.elements, size,

cudaMemcpyHostToDevice);

Matrix d_B;

d_B.width = B.width; d_B.height = B.height;

size = B.width * B.height * sizeof(float);

cudaMalloc(&d_B.elements, size);

cudaMemcpy(d_B.elements, B.elements, size,

cudaMemcpyHostToDevice);

// Allocate C in device memory

Matrix d_C;

d_C.width = C.width; d_C.height = C.height;

size = C.width * C.height * sizeof(float);

cudaMalloc(&d_C.elements, size);

// Invoke kernel

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);

dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);

MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);

// Read C from device memory

cudaMemcpy(C.elements, d_C.elements, size,

cudaMemcpyDeviceToHost);

// Free device memory

cudaFree(d_A.elements);

cudaFree(d_B.elements);

cudaFree(d_C.elements);

}

// Matrix multiplication kernel called by MatMul()

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)

{

// Each thread computes one element of C

// by accumulating results into Cvalue

float Cvalue = 0;

int row = blockIdx.y * blockDim.y + threadIdx.y;

int col = blockIdx.x * blockDim.x + threadIdx.x;

for (int e = 0; e < A.width; ++e)

Cvalue += A.elements[row * A.width + e]

* B.elements[e * B.width + col];

C.elements[row * C.width + col] = Cvalue;

?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者
  • 序言:七十年代末势篡,一起剝皮案震驚了整個濱河市翩肌,隨后出現(xiàn)的幾起案子,更是在濱河造成了極大的恐慌禁悠,老刑警劉巖念祭,帶你破解...
    沈念sama閱讀 206,839評論 6 482
  • 序言:濱河連續(xù)發(fā)生了三起死亡事件,死亡現(xiàn)場離奇詭異碍侦,居然都是意外死亡粱坤,警方通過查閱死者的電腦和手機,發(fā)現(xiàn)死者居然都...
    沈念sama閱讀 88,543評論 2 382
  • 文/潘曉璐 我一進店門瓷产,熙熙樓的掌柜王于貴愁眉苦臉地迎上來站玄,“玉大人,你說我怎么就攤上這事拦英⊙咽玻” “怎么了?”我有些...
    開封第一講書人閱讀 153,116評論 0 344
  • 文/不壞的土叔 我叫張陵疤估,是天一觀的道長灾常。 經(jīng)常有香客問我,道長铃拇,這世上最難降的妖魔是什么钞瀑? 我笑而不...
    開封第一講書人閱讀 55,371評論 1 279
  • 正文 為了忘掉前任,我火速辦了婚禮慷荔,結(jié)果婚禮上雕什,老公的妹妹穿的比我還像新娘。我一直安慰自己显晶,他們只是感情好贷岸,可當我...
    茶點故事閱讀 64,384評論 5 374
  • 文/花漫 我一把揭開白布。 她就那樣靜靜地躺著磷雇,像睡著了一般偿警。 火紅的嫁衣襯著肌膚如雪。 梳的紋絲不亂的頭發(fā)上唯笙,一...
    開封第一講書人閱讀 49,111評論 1 285
  • 那天螟蒸,我揣著相機與錄音,去河邊找鬼崩掘。 笑死七嫌,一個胖子當著我的面吹牛,可吹牛的內(nèi)容都是我干的苞慢。 我是一名探鬼主播诵原,決...
    沈念sama閱讀 38,416評論 3 400
  • 文/蒼蘭香墨 我猛地睜開眼,長吁一口氣:“原來是場噩夢啊……” “哼!你這毒婦竟也來了皮假?” 一聲冷哼從身側(cè)響起鞋拟,我...
    開封第一講書人閱讀 37,053評論 0 259
  • 序言:老撾萬榮一對情侶失蹤骂维,失蹤者是張志新(化名)和其女友劉穎惹资,沒想到半個月后,有當?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體航闺,經(jīng)...
    沈念sama閱讀 43,558評論 1 300
  • 正文 獨居荒郊野嶺守林人離奇死亡褪测,尸身上長有42處帶血的膿包…… 初始之章·張勛 以下內(nèi)容為張勛視角 年9月15日...
    茶點故事閱讀 36,007評論 2 325
  • 正文 我和宋清朗相戀三年,在試婚紗的時候發(fā)現(xiàn)自己被綠了潦刃。 大學時的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片侮措。...
    茶點故事閱讀 38,117評論 1 334
  • 序言:一個原本活蹦亂跳的男人離奇死亡,死狀恐怖乖杠,靈堂內(nèi)的尸體忽然破棺而出分扎,到底是詐尸還是另有隱情,我是刑警寧澤胧洒,帶...
    沈念sama閱讀 33,756評論 4 324
  • 正文 年R本政府宣布畏吓,位于F島的核電站,受9級特大地震影響卫漫,放射性物質(zhì)發(fā)生泄漏菲饼。R本人自食惡果不足惜,卻給世界環(huán)境...
    茶點故事閱讀 39,324評論 3 307
  • 文/蒙蒙 一列赎、第九天 我趴在偏房一處隱蔽的房頂上張望宏悦。 院中可真熱鬧,春花似錦包吝、人聲如沸饼煞。這莊子的主人今日做“春日...
    開封第一講書人閱讀 30,315評論 0 19
  • 文/蒼蘭香墨 我抬頭看了看天上的太陽砖瞧。三九已至,卻和暖如春掺喻,著一層夾襖步出監(jiān)牢的瞬間芭届,已是汗流浹背。 一陣腳步聲響...
    開封第一講書人閱讀 31,539評論 1 262
  • 我被黑心中介騙來泰國打工感耙, 沒想到剛下飛機就差點兒被人妖公主榨干…… 1. 我叫王不留褂乍,地道東北人。 一個月前我還...
    沈念sama閱讀 45,578評論 2 355
  • 正文 我出身青樓即硼,卻偏偏與公主長得像逃片,于是被迫代替她去往敵國和親。 傳聞我的和親對象是個殘疾皇子,可洞房花燭夜當晚...
    茶點故事閱讀 42,877評論 2 345