cuda - 并行处理伪代码以在GPU上工作:克服未对齐的内存访问

Question

Closed. This question needs to be more focused。它当前不接受答案。












想改善这个问题吗？更新问题，使其仅通过 editing this post专注于一个问题。

3年前关闭。



Improve this question





这是我试图并行化的伪代码（取自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] 

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


“并行编程非常有趣”


则词汇表看起来像这样（因为代码根据单词的出现频率对词汇表进行了排序）：

            {“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）。另外，由于所有 neu1和 neu1e更新都是独立的，因此它们可以驻留在线程的专用内存中并独立进行更新。

我现在主要关心的是以下几点：


全局内存访问是以随机方式发生的，因为根据 syn0变量的值（词汇表中的索引）访问 syn1neg和 word。而且，正如我们所看到的，句子中的单词不会以任何顺序出现。


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


在负采样循环中，需要执行归约运算。变量 f是点积的总和。问题是我打算将<​​cc>存储在每个线程的专用内存中，而 neu1则存储在全局内存中。


否定采样是否需要单独的内核？看起来它需要的方法不同于仅启动N个线程/字的方法，但是我不确定哪种方法最有效。

除了这些问题，请提出建议，以了解我处理此代码的方式是否存在问题。

最佳答案

序言：您已打开一罐蠕虫（即使没有SLOC礼物），因此希望您可以接受按部分显示的注释，因为象切片似乎是解决整个复杂主题的唯一方法，而不是“从主要问题逃脱”到实现域的各个代码行的舒适区域，在该区域中，即使尚未丢失先验知识，通常都会错过The Big Picture。



A1:是的，这是一个主要事实（又名“问题”）。

GPU设备是在硅中作为SIMD单指令多数据硬件架构进行设计和优化的，因此它们在代码+数据布局方面都表现最佳，而在整个生命周期中都无需超过-span）确实很小的内存区域（千字节），适合SIMD SM内核的片上内存区域（GPU-SM-REGISTER -s无溢出，由LRU维护的L1缓存），因此不会引入任何“理想化”的性能每次about 350-700 ns访问造成的破坏性潜伏期罚款gloMEM。


[B]-Fig.2状态：TESLA具有SM，每个SMX具有8个[SM]内核，每个SFU一对[SM]，多线程访存和发行[MTI] TPC（纹理/处理器群集）
每个 per 1个16KB shaMEM银行
每个 SM一个只读的 conL1cache（更快/更低的合并冲突）
SM TECH.GPU：NVIDIA CUDA C编程指南，[PG-02829-001_v7.0]； 2015/03


这对于光栅图像处理（以小块有组织的矩阵卷积演算样式处理2D数据数组）的工作效果很好，其中所有线程在相同的时间（肯定是 [B]分段的）执行相同的 （最好是）非冲突数据单元上的指令-这最适合 warpSize。

这也意味着，任何不允许这种独特的逐步渐进式完全对齐的 SIMD -ops的实际操作都将自然地会阻塞性能（线程只能等待线程间的差异，因为（重新） -syncing-barriers，用于远程内存访问，直到最终完成数据传递为止，因此延迟屏蔽越来越少地掩盖了这些自然障碍，使人们看不到真正的 GPGPU代码执行错觉。



SIMD不，实际上不多。可用于定量评估其影响并证明执行时间范围限制的现场基准测试方法和证明。

虽然内核部署指令 PARALLEL有一些帮助：

__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”），对代码通用假设的影响的论文，设计不会被过高估计，因为结果始终是特定于内核的（在实践中，只需探索在 A2:-访问层次结构内的 __launch_bounds__()硅上有限资源部署期间哪种3D几何形状受最小影响）。

尽管可以修改代码执行的3D线程“几何”，但在延迟方面最关键的资源（ SMX s）是有限的，在上下文交换期间，其他线程无法“理想地重用”排程。您计划的线程越多，用于线程的 GPU-off-chip-MEM就越少（对于各个Compute Compatibility XY来说，整体的静态限制不是问题），并且在执行过程中将发生更多的片外内存访问真正的代码执行（无需给出 GPU-SM-REGISTER即可提供此事实，只需遵循已发布的体系结构文档）。



GPU-SM-REGISTER否。内核分离的想法可能会带来从另一种线程3D几何安排中获得潜在好处的幻觉，但是您的代码在为加载/共享协同处理/发布支付额外的性能成本时会遇到更多问题它是数据结构。在完全由 SLOC设计的代码执行中，内核分隔是有意义的。

关于cuda - 并行处理伪代码以在GPU上工作:克服未对齐的内存访问，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/31750210/

上一篇：ant - 用Ant就地过滤文件？

下一篇：flash - 如何使嵌入式视频不自动播放

相关文章：

cuda - CUDA 线程如何工作

Matlab 和 GPU/CUDA 编程

c++ - 无法在动态链接库中找到过程入口点 clCreateCommandQueueWithProperties

c++ - 具有推力内部状态的仿函数

c++ - OpenGL 着色器与 CUDA

python - 并行化一系列生成器

c - 如何使用缩减并行化这个 for 循环？

PHP - 使用 multi_curl 进行并行处理是个好主意吗？

c++ - Intel i5 处理器优于 ATI HD Radeon GPU？

OpenCL 传输速率超过 PCI-e 带宽

Answer 1

序言：您已打开一罐蠕虫（即使没有SLOC礼物），因此希望您可以接受按部分显示的注释，因为象切片似乎是解决整个复杂主题的唯一方法，而不是“从主要问题逃脱”到实现域的各个代码行的舒适区域，在该区域中，即使尚未丢失先验知识，通常都会错过The Big Picture。



A1:是的，这是一个主要事实（又名“问题”）。

GPU设备是在硅中作为SIMD单指令多数据硬件架构进行设计和优化的，因此它们在代码+数据布局方面都表现最佳，而在整个生命周期中都无需超过-span）确实很小的内存区域（千字节），适合SIMD SM内核的片上内存区域（GPU-SM-REGISTER -s无溢出，由LRU维护的L1缓存），因此不会引入任何“理想化”的性能每次about 350-700 ns访问造成的破坏性潜伏期罚款gloMEM。


[B]-Fig.2状态：TESLA具有SM，每个SMX具有8个[SM]内核，每个SFU一对[SM]，多线程访存和发行[MTI] TPC（纹理/处理器群集）
每个 per 1个16KB shaMEM银行
每个 SM一个只读的 conL1cache（更快/更低的合并冲突）
SM TECH.GPU：NVIDIA CUDA C编程指南，[PG-02829-001_v7.0]； 2015/03


这对于光栅图像处理（以小块有组织的矩阵卷积演算样式处理2D数据数组）的工作效果很好，其中所有线程在相同的时间（肯定是 [B]分段的）执行相同的 （最好是）非冲突数据单元上的指令-这最适合 warpSize。

这也意味着，任何不允许这种独特的逐步渐进式完全对齐的 SIMD -ops的实际操作都将自然地会阻塞性能（线程只能等待线程间的差异，因为（重新） -syncing-barriers，用于远程内存访问，直到最终完成数据传递为止，因此延迟屏蔽越来越少地掩盖了这些自然障碍，使人们看不到真正的 GPGPU代码执行错觉。



SIMD不，实际上不多。可用于定量评估其影响并证明执行时间范围限制的现场基准测试方法和证明。

虽然内核部署指令 PARALLEL有一些帮助：

__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”），对代码通用假设的影响的论文，设计不会被过高估计，因为结果始终是特定于内核的（在实践中，只需探索在 A2:-访问层次结构内的 __launch_bounds__()硅上有限资源部署期间哪种3D几何形状受最小影响）。

尽管可以修改代码执行的3D线程“几何”，但在延迟方面最关键的资源（ SMX s）是有限的，在上下文交换期间，其他线程无法“理想地重用”排程。您计划的线程越多，用于线程的 GPU-off-chip-MEM就越少（对于各个Compute Compatibility XY来说，整体的静态限制不是问题），并且在执行过程中将发生更多的片外内存访问真正的代码执行（无需给出 GPU-SM-REGISTER即可提供此事实，只需遵循已发布的体系结构文档）。



GPU-SM-REGISTER否。内核分离的想法可能会带来从另一种线程3D几何安排中获得潜在好处的幻觉，但是您的代码在为加载/共享协同处理/发布支付额外的性能成本时会遇到更多问题它是数据结构。在完全由 SLOC设计的代码执行中，内核分隔是有意义的。