使用 cuda 并行执行乘法和加法答案

【问题标题】：Perform both multiplication and addition in parallel using cuda使用 cuda 并行执行乘法和加法
【发布时间】：2018-05-09 07:41:04
【问题描述】：

我有两个向量，我想并行计算这两个向量的点积。我能够并行地对每个元素进行乘法运算，然后并行地进行加法运算。以下是我尝试过的代码。

但我想同时进行乘法和加法。这意味着即使其他元素尚未进行乘法运算，也应添加已执行乘法运算的元素。希望你明白我所说的。

#include<stdio.h>
#include<cuda.h>
__global__ void dotproduct(int *a,int *b,int *c,int N)
{
    int k=N;
    int i=threadIdx.x;                      
    c[i]=a[i]*b[i];                         
    if(N%2==1)                          
        N=N+1;                          
    __syncthreads();                        
    while(i<(N/2))
    {
        if((i+1)*2<=k)
        {
            c[i]=c[i*2]+c[i*2+1];               
        }
        else
        {
            c[i]=c[k-1];
        }
        k=N/2;
        N=N/2;
        if(N%2==1&&N!=1)                            
            N=N+1;
        __syncthreads();                    //wait for all the threads to synchronize               
    }
}

int main()
{
    int N=10;                           //vector size
    int a[N],b[N],c[N];                     
    int *dev_a,*dev_b,*dev_c;                   
    cudaMalloc((void**)&dev_a,N*sizeof(int));           
    cudaMalloc((void**)&dev_b,N*sizeof(int));           
    cudaMalloc((void**)&dev_c,N*sizeof(int));           
    for(int i=0;i<N;i++)
    {
        a[i]=rand()%10;                     
        b[i]=rand()%10;                     
    }
    cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);   
    cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);   
    dotproduct<<<1,N>>>(dev_a,dev_b,dev_c,N);           
    cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);   
    for(int i=0;i<N;i++)
    {
        printf("%d,%d\n",a[i],b[i]);                

    }
    printf("the answer is : %d in GPU\n",c[0]);         

    cudaFree(dev_a);                        
    cudaFree(dev_b);                        
    cudaFree(dev_c);                        
    cudaThreadExit();                       
    return 0;

}

【问题讨论】：

标签： parallel-processing cuda dot-product

【解决方案1】：

我认为并行进行乘法和加法是没有意义的 - 所有乘法都将花费相同的时间，并且尝试同时运行不同的指令会降低性能。但是可以优化对乘法结果求和的部分。

许多人需要使用原子指令或随机指令 - 阅读以下内容以获得很好的解释：https://devblogs.nvidia.com/faster-parallel-reductions-kepler/

如果这不是练习而是真正的任务，我建议你使用 cuBLAS，它内置了这个功能：https://docs.nvidia.com/cuda/cublas/index.html#cublas-lt-t-gt-dot

【讨论】：

【解决方案2】：

采用您最喜欢的归约（= 向量值的累加），并对其进行修改，以便从两个输入向量执行两次读取，而不是每次读取单个输入向量，使用相同的索引，并将结果相乘。

这将保持尽可能多的并行性，因为您可以使用缩减内核来实现；如果您的读取以前是内存合并的，那么它们现在也将合并。对于非常长的向量，您的每时间单位输出元素的吞吐量应该几乎是之前的一半。

【讨论】：

能否如您所说的那样修改代码。
是的，我可以。但你也可以。此外，您的代码在多对向量之间执行点积 - 每个线程块一个 - 并且对于较大的 N 值是有问题的。

【解决方案3】：

我认为还有一个特殊的 PTX ISA 指令可以立即进行点积（ dp4a.atype.btype d, a, b, c; ）。不费吹灰之力，您就可以尝试编写一个小的内联 PTX 汇编函数。检查文档。

【讨论】：