c++ 简单加法示例:还原的共享内存版本执行速度慢于全局内存

qacovj5a  于 2023-05-20  发布在  其他
关注(0)|答案(1)|浏览(153)

我已经实现了add的两个版本。两者中加法的概念是完全相同的。唯一的区别是,在一个代码(下面的第一个)中,我使用全局内存,而在第二个代码中,我使用共享内存。正如在几个地方提到的,共享内存版本应该更快,但对于我的情况,全局内存版本更快。请告诉我我错在哪里注意:我有一个CC 2.1的GPU。因此,对于共享内存,我有32个银行。由于我在示例中只使用了16个int,因此我的代码应该没有内存冲突。请告诉我是否正确
全球版本

#include<stdio.h>
__global__ void reductionGlobal(int* in, int sizeArray, int offset){

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if(tid < sizeArray ){
        if(tid % (offset * 2 ) == 0){
            in[tid] += in[tid+offset];
        }
        
    }

}
int main(){
    int size = 16; // size of present input array. Changes after every loop iteration
    int cidata[] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16};

    int* gidata;
    cudaMalloc((void**)&gidata, size* sizeof(int));
    cudaMemcpy(gidata,cidata, size * sizeof(int), cudaMemcpyHostToDevice);
    int offset = 1; 
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);
    while(offset < size){
        //use kernel launches to synchronize between different block. syncthreads() will not work
        reductionGlobal<<<4,4>>>(gidata,size,offset);
        offset *=2;
        
    }
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime , start, stop);
    printf("time is %f ms", elapsedTime);
    int* output = (int*)malloc( size * sizeof(int));
    cudaMemcpy(output, gidata, size * sizeof(int), cudaMemcpyDeviceToHost);
    printf("The sum of the array using only global memory is %d\n",output[0]);
    getchar();
    return 0;
}

共享内存版本:

#include<stdio.h>

__global__ void computeAddShared(int *in , int *out, int sizeInput){
    extern __shared__ float temp[];

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int ltid = threadIdx.x;
    temp[ltid] = 0;
    while(tid < sizeInput){
        temp[ltid] += in[tid];
        tid+=gridDim.x * blockDim.x; // to handle array of any size
    }
    __syncthreads();
    int offset = 1;
    while(offset < blockDim.x){
        if(ltid % (offset * 2) == 0){
            temp[ltid] = temp[ltid] + temp[ltid + offset];
        }
        __syncthreads();
        offset*=2;
    }
    if(ltid == 0){
        out[blockIdx.x] = temp[0];
    }
    
}

int main(){
    
    int size = 16; // size of present input array. Changes after every loop iteration
    int cidata[] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16};
    
    int* gidata;
    int* godata;
    cudaMalloc((void**)&gidata, size* sizeof(int));
    cudaMemcpy(gidata,cidata, size * sizeof(int), cudaMemcpyHostToDevice);
    int TPB  = 4;
    int blocks = 10; //to get things kicked off
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);
    while(blocks != 1 ){
        if(size < TPB){
            TPB  = size; // size is 2^sth
        }
        blocks  = (size+ TPB -1 ) / TPB;
        cudaMalloc((void**)&godata, blocks * sizeof(int));
        computeAddShared<<<blocks, TPB,TPB>>>(gidata, godata,size);
        cudaFree(gidata);
        gidata = godata;
        size = blocks;
    }
    
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime , start, stop);
    printf("time is %f ms", elapsedTime);
    int *output = (int*)malloc(sizeof(int));
    cudaMemcpy(output, gidata, sizeof(int), cudaMemcpyDeviceToHost);
    //Cant free either earlier as both point to same location
    cudaFree(godata);
    cudaFree(gidata);
    printf("The sum of the array is %d\n", output[0]);
    getchar();
    return 0;
}
yvfmudvl

yvfmudvl1#

这里有很多错误。首先是一些一般性意见:

  • 您正在对16个数字执行约简,这是一个非常小的输入大小。CUDA在主机和设备端都有很多固定开销。你给设备的工作量是如此之小,以至于你所测量的只是这些开销,而不是GPU的执行时间。您看到的两个代码之间的差异可能只是由于共享内存版本增加了安装开销。当然,这与代码本身无关。如果你想测量代码的实际性能,你给予代码的工作量必须足够大,以确保执行时间远远大于设置时间。请放心,你有大约5个数量级太少的工作,这一点,即使在一个小的GPU。
  • 您提到了银行冲突,但这是您正在使用的架构上的一个稻草人。费米有一个完全不同的共享内存布局相比,旧的硬件,只有一个相对较小的问题,银行冲突。在这种情况下当然没什么好担心的。

至于实际的缩减代码本身:

  • 如果你不能在一次内核启动中将输入数组减少到每个线程一个部分和,那么你真的没有充分考虑这个问题。您当前的方法在“全局”和“共享”版本中的效率都非常低。并行约简是一个解决的问题,CUDA SDK附带了一份关于GPU上的优化和约简性能的优秀白色。你应该在做其他事情之前先读一读。
  • 一旦达到了每个线程都有一个部分和的程度,就需要执行共享内存减少 * 每个块 *,以使减少每个块产生一个部分和。这将只需要两次内核启动来计算完整的约简。
  • 您的“共享”版本有一个缓冲区溢出,这将导致运行时错误。在启动时指定的动态共享内存大小以 * 字节 * 为单位,而不是字。如果你的代码有错误检查,你应该已经发现了。Fermi具有出色的共享内存保护,如果您尝试在静态或动态分配的内存之外进行写入,它将生成运行时错误。

相关问题