上篇博客的快排中用到了基于warp的cuda操作用于分隔數(shù)組, 為什么要將控制線程的級(jí)別定義為warp呢？

在一個(gè)warp內(nèi)，線程的可以通過(guò)__ballot函數(shù)，并發(fā)的獲取這32個(gè)數(shù)中于pivot的比較結(jié)果，然后通過(guò)ptx類(lèi)似匯編的語(yǔ)句asm("mov.u32 %0, %%lanemask_lt;" : "=r"(lane_mask)) 獲得線程在warp內(nèi)的位置的掩碼，再按位與之后調(diào)用 __popc函數(shù)就可以獲得在這個(gè)warp內(nèi)這個(gè)線程之前有多少個(gè)線程對(duì)應(yīng)的數(shù)大于或者小于Pivot，就可以獲得這個(gè)線程對(duì)應(yīng)的數(shù)的偏移，進(jìn)而就實(shí)現(xiàn)了分割數(shù)組。

這里的所獲得的啟發(fā)就是，一個(gè)看似只能串行的掃描操作，也可以通過(guò)控制warp實(shí)現(xiàn)并行

到這里聯(lián)想到之前研究風(fēng)辰大神對(duì)word2vec的cuda改寫(xiě)，也實(shí)現(xiàn)了對(duì)warp的精細(xì)控制，進(jìn)而獲得的極大的加速?https://github.com/fengChenHPC/word2vec_cbow

int blockSize = 256;

int gridSize = (sentence_length)/(blockSize/32);

cbow_kernel<1><<<gridSize, blockSize, smsize>>>()

一個(gè)block有8個(gè)warp，一個(gè)warp處理一個(gè)字，一個(gè)block可以處理8個(gè)字，一共有sentence_length個(gè)字，所以需要gridSize個(gè)block

一個(gè)字對(duì)應(yīng)這一個(gè)特征向量的相乘操作，例如

for (int c = idInWarp; c < layer1_size; c += warpSize) neu1[c] += syn0[c + last_word * layer1_size];

比如一個(gè)字對(duì)應(yīng)了幾百維的特征向量，這個(gè)字又對(duì)應(yīng)一個(gè)warp內(nèi)的32個(gè)線程，可以用32個(gè)線程實(shí)現(xiàn)對(duì)向量相乘的并行

總結(jié)

以上是生活随笔為你收集整理的cuda版本的word2vec的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

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