CUDA内核与Julia集CPU版本的性能下降

Question

我从the book“CUDA通过示例学习CUDA”。第4章中有一个生成Julia分形的演示。该演示展示了CPU和GPU版本。我决定为两种情况增加一段时间来查看执行速度，并且令我惊喜的是发现CPU版本比GPU执行速度快3倍。

CPU朱生成总时间：

745毫秒。

GPU朱代总时间：

2456毫秒。

那么是怎么回事？很明显，至少从CUDA内核代码来看，执行是平行的，分布在1000个块中，每个块计算1000x1000分辨率最终图像的像素。

这里是执行的源代码：

#define N 10 
#define DIM 1000 
typedef unsigned char byte; 

struct cuComplex { 
    float r; 
    float i; 
    __host__ __device__ cuComplex(float a, float b) : r(a), i(b) {} 
    __host__ __device__ float magnitude2(void) { 
      return r * r + i * i; 
    } 
    __host__ __device__ cuComplex operator*(const cuComplex& a) { 
     return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i); 
    } 
    __host__ __device__ cuComplex operator+(const cuComplex& a) { 
     return cuComplex(r+a.r, i+a.i); 
    } 
}; 

__device__ int juliaGPU(int x , int y){ 
    const float scale =1.3; 
    float jx = scale * (float)(DIM/2 -x)/(DIM/2); 
    float jy= scale *(float)(DIM/2 -y)/(DIM/2); 

    cuComplex c(-0.8 ,0.156); 
    cuComplex a(jx ,jy); 
    int i = 0; 
    for(i=0; i <200;i++){ 
     a = a * a +c; 
     if(a.magnitude2() >1000){ 

      return 0; 
     } 
    } 
    return 1; 

} 

__global__ void kernelGPU(byte *ptr){ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    int offset =x + y * gridDim.x; 

    int juliaValue =juliaGPU(x , y); 
    ptr[offset * 4 + 0]=255 * juliaValue; 
    ptr[offset * 4 + 1]=0; 
    ptr[offset * 4 + 2]=0; 
    ptr[offset * 4 + 3]=255 ; 
} 


struct DataBlock { 
    unsigned char *dev_bitmap; 
}; 
void juliaGPUTestSample(){ 
DataBlock data; 
CPUBitmap bitmap(DIM,DIM); 
byte *dev_bitmap; //memory on GPU 
HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap , bitmap.image_size())); 
data.dev_bitmap =dev_bitmap; 
dim3 grid(DIM,DIM); 
int starTime=glutGet(GLUT_ELAPSED_TIME); 

kernelGPU<<<grid ,1 >>>(dev_bitmap); 
HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr() , dev_bitmap ,bitmap.image_size() ,cudaMemcpyDeviceToHost)); 
int endTime=glutGet(GLUT_ELAPSED_TIME)-starTime; 
printf("Total time %d\n:" ,endTime); 
HANDLE_ERROR(cudaFree(dev_bitmap)); 

bitmap.display_and_exit(); 
} 

int main(void){ 
juliaGPUTestSample(); 
return 1; 

} 

这里是CPU版本：

///的 “cuComplex” 结构是从上方一样。

int julia (int x , int y){ 

const float scale = 1.3; 
float jx = scale * (float)(DIM/2 -x)/(DIM/2); 
float jy = scale * (float)(DIM/2 -y)/(DIM/2); 

cuComplex c(-0.8 ,0.156); 
cuComplex a(jx ,jy); 

int i = 0; 
for(i=0; i <200;i++){ 

    a = a * a +c; 
    if(a.magnitude2() >1000){ 

     return 0; 
    } 
} 

return 1; 

} 

void kernel(unsigned char *ptr){ 

for(int y = 0 ; y <DIM ;++y){ 
    for(int x = 0 ; x <DIM ; ++x){ 
     int offset =x + y * DIM; 
     int juliaValue = julia(x , y); 

     ptr[offset * 4 + 0 ] = juliaValue * 125; 
     ptr[offset * 4 + 1 ] = juliaValue * x; 
     ptr[offset * 4 + 2 ] = juliaValue * y; 
     ptr[offset * 4 + 3 ] = 255 ; 
    } 
} 

} 
void juliaCPUTestSample(){ 

CPUBitmap bitmap(DIM ,DIM); 
unsigned char *ptr = bitmap.get_ptr(); 
int starTime=glutGet(GLUT_ELAPSED_TIME); 

kernel(ptr); 

int endTime=glutGet(GLUT_ELAPSED_TIME)-starTime; 
printf("Total time %d\n:" ,endTime); 
bitmap.display_and_exit(); 

} 

更新 - 系统配置：

64位Windows 7

CPU - 英特尔酷睿i7 -3770CPU 3.40GHz，16GB RAM

GPU - 的NVIDIA Quadro 4000

Answer 1

其他有noticed this。

首先，当谈到CPU和GPU之间的perf比较时，提到系统配置（包括hw平台和软件）是一个不错的主意。例如，我将代码运行在带有核心i7 2.60GHz四核CPU和一个运行RHEL 6.2和cuda 5.0的quadro1000M GPU的惠普笔记本电脑上，我得到了GPU的438分和CPU的441分。其次，更重要的是，该书中的茱莉亚样本是CUDA编码的一个相对较早的例子，所以它并不是真正面向最大性能，而是为了说明迄今为止已经讨论过的概念。该书以及其他各种CUDA教程资料首先在块级中引入使用CUDA的并行编程。这样做的指示是在这里：

kernelGPU<<<grid ,1 >>>(dev_bitmap); 

的内核启动参数<<<grid, 1>>>表明，一些数量（grid，这是百万在这种情况下总的块）的块将被启动，与具有每个块的栅格单线程。例如，与使用每个线程块的完整线程启动网格相比，这会立即将Fermi级GPU的功耗降低1/32倍。费米级GPU中的每个SM都有32个线程处理器，全部以锁步执行。如果你启动一个只有16个线程的块，那么16个线程处理器将执行你的代码，而其他16个线程处理器将不会执行任何操作（即没有用处）。因此，仅包含1个线程的线程块将仅使用32个线程处理器中的1个，另外31个为空闲。

因此，这个特定的代码示例不能很好地利用GPU的全部并行功能。鉴于本书对CUDA概念的阐述相对较早，这是可以理解的;我不相信这是作者意图将此代码作为基准或用作如何在GPU上编写快速代码的合法代表。

考虑到1/32的这个因素，在你的系统上CPU的速度只有3倍，而在我的系统上，CPU和GPU具有相当的吞吐量（不管这些是特别高性能的CUDA GPU ，最有可能的）我认为这显示了GPU的光线很好。 GPU正在进行这场战斗，其中约97％的能力未被使用。