简单加法示例：共享内存版本的缩减执行速度低于全局内存

Question

我已经实现了两个版本的add。两者的加法概念完全相同。唯一的区别是在一个代码中（下面的第一个代码）我使用全局内存，而第二个代码使用共享内存。正如在几个地方提到的那样，共享内存版本应该更快，但就我而言，全局内存版本更快。 请告诉我哪里出错了。注意：我有一个cc 2.1的GPU。因此，对于共享内存，我有32家银行。由于我在示例中仅使用了16个整数，所以我的代码不应该有银行冲突。 请让我知道这是否正确。

全球版本

#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; 
} 

Answer 1

有很多错在这里。首先，一些一般性评论：

• 您正在对16个数字进行缩减，这是一个可笑的 小输入大小。 CUDA在主机 和设备端都有很多固定开销。您为设备执行的工作量太小，您所测量的只是那些开销，而不是GPU的执行时间。您看到 的两个代码之间的差异可能仅仅是由于在共享内存版本为 的情况下增加了设置开销。当然与代码本身无关。如果您想测量代码的实际性能，那么您为该代码提供的工作量必须足够大，以确保执行时间远大于设置时间。请放心，即使在一个小型GPU上，对于这个问题你也有5个数量级的工作量。
• 你已经提到银行冲突，但这是你正在使用的架构上的稻草人。与旧的硬件相比，费米拥有完全不同的共享内存布局，并且在银行冲突方面只有相对较小的问题。在这种情况下肯定没什么可担心的。

至于实际减少代码本身：

• 如果你不能想出办法对于要在单个内核启动 降至输入数组每个线程一个部分和，那么你真的还没有想到 足够的问题。您在“全球” 和“共享”版本中的当前方法极其低效。并行压缩是一个解决的问题，CUDA SDK附带一份关于优化和降低GPU性能的优秀白皮书。你应该在做任何事之前阅读它。
• 一旦您达到每个线程有一个部分总和的点，您想要执行共享内存减少每块以减少每个块发出一个部分总和。这将只需要两个内核启动 来计算完整的减少。
• 您的“共享”版本有一个缓冲区溢出，应该导致运行时错误 。在启动 时间指定的动态共享内存大小是字节，而不是单词。如果你的代码有错误检查，你 已经找到了这个。费米拥有出色的共享内存保护，如果您尝试在静态或动态分配的内容之外编写 ，则会产生运行时错误。