gcc-via-nvcc 是不是矢量化这些总和和最大减少？

Posted 2023-02-22

【中文标题】gcc-via-nvcc 是不是矢量化这些总和和最大减少？

【英文标题】：Does gcc-via-nvcc vectorize these sum and max reductions?gcc-via-nvcc 是否矢量化这些总和和最大减少？

【发布时间】：2013-11-25 23:30:45

【问题描述】：

我有一个 CUDA 应用程序，它当前使用推力库在 GPU 上对向量进行求和和最大缩减。我发现对于某些向量长度，如果我将向量发送回主机并计算 C++ 中的总和和最大减少量，它会快得多。

总和和最大减少量应该在主机上是可矢量化的。主机上的内存是线性/连续的，我正在使用的编译器（GCC）支持它。鉴于我看到的时间，编译器似乎正在对代码进行矢量化，但我该如何确认呢？我没有任何强制编译器优化的经验，但我知道可以使用一些编译指示语句。 （不过，你会通过谷歌搜索找到很少的信息。）另外，我宁愿不去挖掘汇编来确认，因为我不会理解它。是否有编译器设置（在 GCC 或 NVCC 中）可用于在主机上强制矢量化或确认代码正在矢量化？

我为 sum 和 max reduction 编写的函数如下。 nvcc 编译器最终会编译它，因为函数包含 CUDA 代码。

void calc_vector_max_host(double& maxval, double *const vec_h, const double *const vec_d, int len)

    //copy device vector to host
    gpuErrchk(cudaMemcpy(vec_h, vec_d, len*sizeof(double), cudaMemcpyDeviceToHost));

    //vectorized? max
    maxval = *vec_h;
    double* temp = vec_h;
    for(int i = 1; i < len; i++, temp++)
    
        if(*temp > maxval)
        
            maxval = *temp;
        
    


void calc_vector_sum_host(double& sum, double *const vec_h, const double *const vec_d, int len)

    //copy device vector to host
    gpuErrchk(cudaMemcpy(vec_h, vec_d, len*sizeof(double), cudaMemcpyDeviceToHost));

    //vectorized? sum
    sum = 0.0;
    double* temp = vec_h;
    for(int i = 0; i < len; i++, temp++)
    
        sum += *temp;
    

编辑：以下包含 gcc 自动矢量化所需的更正。还需要 cmets 中列出的编译器选项。

void calc_vector_max_host(double& maxval, double *const __restrict__ vec_h, const double *const __restrict__ vec_d, int len)

    //copy device vector to host
    gpuErrchk(cudaMemcpy(vec_h, vec_d, len*sizeof(double), cudaMemcpyDeviceToHost));

    //vectorized? max
    double local_maxval = vec_h[0];
    for(int i = 1; i < len; i++)
    
        double val = vec_h[i];
        if(val > local_maxval)
        
            local_maxval = val;
        
    
    maxval = local_maxval;


void calc_vector_sum_host(double& sum, double *const __restrict__ vec_h, const double *const vec_d, int len)

    //copy device vector to host
    gpuErrchk(cudaMemcpy(vec_h, vec_d, len*sizeof(double), cudaMemcpyDeviceToHost));

    //vectorized? sum
    double local_sum = 0.0;
    for(int i = 0; i < len; i++)
    
        local_sum += vec_h[i];
    
    sum = local_sum;

【参考方案1】：

首先，最重要的一点是 nvcc 不是编译器。它是一个编译器驱动程序——它只是使用主机编译器和几个自定义预处理工具来引导编译，这些工具可以分解实际的 GPU 代码并将其传递给 GPU 工具链。 GPU 编译器和汇编器只触及典型程序的一小部分。其余的构建直接使用主机编译器和链接器完成。

所以您发布的所有代码都是由 gcc 编译的（并且可以直接编译而无需使用 nvcc）。 nvcc 有一个选项-Xcompiler，可用于将您想要的任何选项传递给主机编译轨迹。对于矢量化，您可以传递您的 gcc 版本支持的here 中描述的任何选项。如果您愿意，也可以直接使用 SSE 风格的内在函数来简化编译器的工作

要查看您的主机代码中是否已经发生矢量化，只需使用 objdump/otool 之类的东西（取决于您使用的是 Linux 还是 OS X，您还没有说）。您将能够查看编译器发出的代码的反汇编，并且矢量化指令的存在将立即回答您的问题。

最后，nvcc 现在有很好的documentation 毫无价值，通过熟悉它，您可以找到这个问题的答案，也可能找到您曾经遇到过的有关 nvcc 的所有其他问题的答案。

【讨论】：

几个cmets，仅供参考。 gcc 选项网页的链接很有帮助。最初发布的代码毕竟不是自动矢量化的。这是由于缺少一些编译器选项。具体来说，我必须使用 -Xcompiler 标志传递所有 -ftree-vectorize -msse2 和 -ffast-math 以获得要矢量化的代码，以及 ftree-vectorize-verbose=6 以获得自动矢量化编译器输出。我正在使用 Linux 并尝试了 objdump，但没有发现程序集转储对这个问题非常有用。功能也有一些问题，已修改。