本文共 797 字，大约阅读时间需要 2 分钟。

上篇博客的快排中用到了基于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个线程实现对向量相乘的并行

转载地址：http://cyeti.baihongyu.com/