现在的位置: 首页 > 综合 > 正文

cuda版本的word2vec

2018年11月09日 ⁄ 综合 ⁄ 共 755字 ⁄ 字号 评论关闭

上篇博客的快排中用到了基于warp的cuda操作用于分隔数组, 为什么要将控制线程的级别定义为warp呢?

在一个warp内,线程的可以通过__ballot函数,并发的获取这32个数中于pivot的比较结果,然后通过ptx类似汇编的语句asm("mov.u32 %0, %%lanemask_lt;" : "=r"(lane_mask)) 获得线程在warp内的位置的掩码,再按位与之后调用 __popc函数就可以获得在这个warp内这个线程之前有多少个线程对应的数大于或者小于Pivot,就可以获得这个线程对应的数的偏移,进而就实现了分割数组。

这里的所获得的启发就是,一个看似只能串行的扫描操作,也可以通过控制warp实现并行

到这里联想到之前研究风辰大神对word2vec的cuda改写,也实现了对warp的精细控制,进而获得的极大的加速 https://github.com/fengChenHPC/word2vec_cbow

int blockSize = 256;

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

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

一个block有8个warp,一个warp处理一个字,一个block可以处理8个字,一共有sentence_length个字,所以需要gridSize个block

一个字对应这一个特征向量的相乘操作,例如

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

比如一个字对应了几百维的特征向量,这个字又对应一个warp内的32个线程,可以用32个线程实现对向量相乘的并行

抱歉!评论已关闭.