3

这是我试图并行化的伪代码(取自 word2vec C 代码)。首先,我将列出具有相应大小的数据结构,然后是伪代码:

1.  long long sen[MAX_SENTENCE_LENGTH]  
// In the C code, MAX_SENTENCE_LENGTH = 1000. Increasing this should be  
//fine.

2.  float neu1[N] (hidden layer values)
//N is the length of each vector. For now, max N = 400

3.  float neu1e[N] (hidden layer error values)

4.  float syn0[V * N] (input to hidden layer weight matrix)
// For now, we can assume that V * N is small enough to be stored on the GPU
   // In the test data, V = 72k words

5.  float syn1neg[V * N] (back propagation weights used during negative  
sampling)

6. float exptable[1000] 

程序的输入是一个文本文件。然后程序一次处理一个单词以建立一个词汇表。例如,如果我的文本文件有一个句子

“并行编程非常有趣”</p>

那么词汇表看起来像这样(因为代码根据单词的频率对词汇表进行排序):

            {“Very:2”, “Parallel:1”, “programming:1”, “is:1”,    “interesting:1”}
                   0      1               2              3                4

构建词汇表后,代码再次开始处理文本,一次处理 1000 个单词。前 1000 个单词存储在 中sen[MAX_SENTENCE_LENGTH],然后针对 中的所有单词训练一个神经网络sen,这个过程一直持续到文件末尾。对于上面的句子,sen将如下所示[1,2,3,0,0,4]

假设训练只在一次迭代中完成,伪代码如下:

for sen in text
{ 
    for word in sen
    {

        for (c = 0; c < N; c++) 
            neu1[c] = 0;

        for (c = 0; c < N; c++) 
            neu1e[c] = 0;   

       /*The variable window is a user supplied parameter. 
        It is used to consider the context  around a word in a sentence. 
        For example, if I am looking at the first word in the sentence
        (target word is word1), and window = 5, then the words in the 
        window = {word2, word3, word4, word5}. 
        If I am looking at the third word in the sentence 
        (target word is word3), then window = {word1, word2, word4, word5}*/    

        for word in window
        {
            for (c = 0; c < N; c++) 
            neu1[c] += syn0[c + word * N];
        }

        for (c = 0; c < N; c++) 
            neu1[c] /= window;

        //negative: number of negative samples to provide (assume it to be 
             //between 5 to 25)
        for (d = 0; d < negative + 1; d++) 
        {

            target = sen[random_index]  
            l2 = target * N;
            f = 0;
            for (c = 0; c < N; c++) 
            f += neu1[c] * syn1neg[c + l2];

           gradient = exptable[function of f] //f is calculated in the loop above

           for (c = 0; c < N; c++) 
              neu1e[c] += gradient * syn1neg[c + l2];

           for (c = 0; c < N; c++) 
              syn1neg[c + l2] += gradient * neu1[c];

          } //Negative Sampling ends    

        for word in window
        {
             for (c = 0; c < N; c++) 
                syn0[c + word * N] += neu1e[c];
        }

   } // word in sen loop ends

 } // sen in text loop ends

我认为并行化的最佳方法是并行处理句子中的单词。考虑到所有循环,我认为我应该N每个单词使用线程,以便单个线程在syn0, syn1neg每个循环中只访问一次全局内存()。此外,由于所有neu1neu1e更新都是独立的,它们可以驻留在线程的私有内存中并独立更新。

我现在主要担心以下几点:

  1. 全局内存访问以随机方式发生,因为syn0syn1neg是根据word变量的值(词汇表中的索引)访问的。而且,正如我们所见,句子中的单词不会以任何顺序出现。

这是一个大问题吗?或者,我们可以通过为 GPU 提供足够数量的线程来隐藏内存延迟吗?另外,我不明白这种访问模式是否真的是随机的,因为 N 个线程/字将访问 syn0 和 syn1neg 中的顺序数据,但下一组 N 个线程可能会访问位于内存中很远的顺序数据。

  1. 在负采样循环中,需要进行归约操作。该变量f是点积的总和。问题是我打算存储neu1在每个线程的私有内存中,而syn1neg在全局内存中。

负采样是否需要单独的内核?看起来它需要一种与仅启动 N 个线程/单词不同的方法,但我不确定哪种方法最有效。

除了这些问题,请提出我处理此代码的方式是否存在问题。

4

1 回答 1

-2

序言:你已经打开了一罐蠕虫(即使没有SLOC礼物),所以希望你能接受每部分提出的评论,因为大象切片似乎是解决整个复杂主题的唯一方法,而不是“逃避”主要问题进入实现域的各个代码行的舒适区,如果没有先验丢失,通常会错过大图。


A1:的,这是一个主要事实(又名“问题”)。

GPUSIMD-设备硅中设计和优化为单指令多数据硬件架构因此它们代码+数据布局方面表现最佳,无需超过(在其整个生命周期内) )确实很小的内存区域(千字节),适合SIMD SM内核的片上内存区域(GPU-SM-REGISTER没有溢出的-s,例如LRU维护的L1缓存),因此不会引入任何“理想化”的性能破坏性延迟惩罚about 350-700 ns每次gloMEM访问。

[B]-Fig.2状态:每个

TESLA具有8 个 [ ] 核心SM每个[ ]对,多线程提取和问题 [ ] TPC(纹理/处理器集群)每个 16KB库每个只读 (更快/更低的合并冲突)TECH.GPU:NVIDIA CUDA C 编程指南,[PG-02829-001_v7.0];2015/03
SMXSM
SFUSM
MTIper
shaMEMSM
conL1cache SM
[B]

这适用于光栅图像处理(以小图块组织的矩阵卷积演算风格处理二维数据数组),其中所有线程warpSize在相同(当然,-分段)时间执行相同的指令理想情况下) 非碰撞数据单元——这最适合。SIMDGPGPU

这也意味着,任何不允许这种独特的逐步渐进完全对齐的操作的现实生活中的操作都SIMD将自然而然地阻塞性能(线程可以但等待线程间的分歧,用于(重新)同步-障碍,用于远程内存访问,直到最终发生数据交付,因此延迟屏蔽越来越少地隐藏PARALLEL这些自然障碍,使他们无法看到真正的代码执行的错觉)。


A2:不,几乎没有。可用于定量评估其影响并证明结果执行时间范围的限制的现场基准测试方法和证明。

虽然内核部署指令有一些帮助__launch_bounds__()

__global__ void
__launch_bounds__( 1,     // dim3tbGridSIZE <maxThreadsPerBLOCK>         COMPILE-TIME_ADVICE_FOR_OPTIMISING_COMPILER_____________REGISTERs, CACHE_, FETCH_, PROXIMITY_PATTERNs ANALYSES
                   1  /*, // dim3tBlockSIZE <minBlocksPerMULTIPROCESSOR> COMPILE-TIME_ADVICE_FOR_OPTIMISING_COMPILER_____________OPTIMUM_SCHEDULE_TO_FILL_FETCH_LATENCIES
                   ?,     // iAsyncSeqOfCmdsQUEUE_Stream_ID <<- TO LET BE FREELY ASSIGNABLE ... NON-BLOCKING EXEC'd KERNEL
                   0  */  // iSharedMemSIZE
                   )
                 Device_printf_GPU_CLK( int const iTag ){
                        ...
                        return;
}

有很多关于广泛(蛮力扫描等)“优化”的论文发表,相当机械调整,各种启动参数化的内核编译(threading-3D-“geometry”),对代码的一般假设的影响 -设计不应被高估,因为结果始终是内核特定的(并且在实践中,只需探索在SMX硅内部GPU-off-chip-MEM访问层次结构上的有限资源部署期间哪个 3D 几何结构将受到的影响最小)。

虽然可以修改代码执行的线程 3D“几何”,但延迟方面最关键的资源 ( GPU-SM-REGISTERs) 是有限的,并且不能在调度期间被上下文交换上的其他线程“理想地重用”。您计划的线程越多,GPU-SM-REGISTER专用于线程的 s 就越少(各个计算兼容性 XY 的总体静态限制不是问题),并且在实际代码期间将发生更多的片外内存访问 -执行(无需SLOC为给出这一事实而编写,只需遵循已发布的架构文档即可)。


A3:。内核分离的想法可能会引入另一种线程 3D 几何排列的潜在好处的错觉,但是您的代码将在为加载/共享协同处理/发布它的数据而支付额外的性能成本方面遇到更多问题-结构。RDMA内核分离在完全设计的代码执行中是有意义的。

于 2016-03-29T10:30:49.057 回答