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％的能力未被使用。