CUDA版本比CPU版本慢嗎？

我正在CUDA中編寫一個圖像子採樣器，並使用這些線程來執行平均操作。但是，如果我在不調用內核的情況下執行此操作，與實際調用CUDA內核相比，運行速度要快得多。圖像大小現在爲1280x1024。內核調用是否需要大量時間，或者我的實現有問題嗎？CUDA版本比CPU版本慢嗎？

P.S我試着只調用內核（代碼被移除），它與內核的代碼幾乎是同一時間。另外我的內核調用代碼運行大約350毫秒，而內核調用運行時間接近1000毫秒。

__global__ void subsampler(int *r_d,int *g_d,int *b_d, int height,int width,int *f_r,int*f_g,int*f_b){ 
     int id=blockIdx.x * blockDim.x*blockDim.y+ threadIdx.y*blockDim.x+threadIdx.x+blockIdx.y*gridDim.x*blockDim.x*blockDim.y; 
     if (id<height*width/4){ 
     f_r[id]=(r_d[4*id]+r_d[4*id+1]+r_d[4*id+2]+r_d[4*id+3])/4; 
     f_g[id]=(g_d[4*id]+g_d[4*id+1]+g_d[4*id+2]+g_d[4*id+3])/4; 
     f_b[id]=(b_d[4*id]+b_d[4*id+1]+b_d[4*id+2]+b_d[4*id+3])/4; 
     } 
     }

我定義blockSizeX和blockSizeY爲1和1（我試圖使它們4,16），但不知何故，這是最快的

dim3 blockSize(blocksizeX,blocksizeY); 
    int new_width=img_width/2; 
    int new_height=img_height/2; 

    int n_blocks_x=new_width/blocksizeX+(new_width/blocksizeY == 0 ?0:1); 
    int n_blocks_y=new_height/blocksizeX+(new_height/blocksizeY == 0 ?0:1); 
    dim3 gridSize(n_blocks_x,n_blocks_y);

，然後我打電話跟gridSize，BLOCKSIZE內核。

來源

2011-02-05 Manish

線程數/塊數？你爲什麼不指定線程的數量，這樣你可以擺脫if（）？ – 2011-02-05 21:44:02

我編輯上面的線程/塊。我不知道我該如何擺脫'如果'，如果它傷害了性能（因爲我測量性能刪除那塊，並調用空內核，它幾乎需要相同的時間） – Manish 2011-02-06 00:14:02

這可能是因爲內核執行得不好，或者可能是將數據移入或移出GPU卡的開銷正在減少任何計算上的好處。嘗試單獨對內核進行基準測試（無需CPU <-> GPU內存傳輸），查看內核佔用您的總時間以及內存傳輸的總時間。然後，您可以根據這些測量結果來決定是否需要在內核上做更多的工作。

來源

2011-02-05 19:24:56

雖然我不知道你的硬件正在運行這一個，你應該能夠讓這個內核接近1000 fps的執行，而不是1000毫秒/幀:)

建議1：如果該處理通過OpenGL/DirectX或類似方式與可視化進行任何交互，只需將其作爲着色器進行處理即可處理網格/塊大小，內存佈局等的所有細節。如果你真的需要在CUDA中自己實現，請繼續閱讀：

首先，我假設你在每個方向上對1280x1024圖像進行二次採樣，得到640x512的圖像。生成圖像中的每個像素是源圖像中四個像素的平均值。圖像有三個通道，RGB。

問題1：你真的想32位每通道還是你想RGB888（每通道8位）？ RGB888相當普遍 - 我會認爲這就是你的意思。

問題2：您的數據實際上是平面的，還是您從交錯格式中提取？ RGB888是一種交錯格式，其中像素以RGBRGBRGB形式出現在內存中。我會編寫你的內核來處理它的原始格式的圖像。我會假設你的數據實際上是平面的，所以你有三架飛機，R8，G8和B8。

要做的第一件事就是考慮內存佈局。您將需要一個線程爲目標圖像中的每個像素。假定子採樣的存儲器訪問模式不合並，您將需要將像素數據讀入共享內存。考慮一個32x8線程的塊大小。這允許每個塊以40 * 8 * 4像素讀取，或者在3bpp讀取3072個字節。實際上，您將讀取的內容略多於此，以保持負載合併，每塊總共4096字節。這現在給你：

dim3 block(32, 8); 
dim3 grid(1280/2/32, 1024/2/8); // 20x64 blocks of 256 threads

現在來有趣的部分：做共享內存。你的內核可能看起來像這樣：

__global__ void subsample(uchar* r, uchar* g, uchar* b, // in 
          uchar* ro, uchar* go, uchar* bo) // out 
{ 
    /* Global offset into output pixel arrays */ 
    int gid = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x; 

    /* Global offset into input pixel arrays */ 
    int gidin = gid * 2; 

    __shared__ uchar* rc[1024]; 
    __shared__ uchar* gc[1024]; 
    __shared__ uchar* bc[1024]; 

    /* Read r, g, and b, into shmem cache */ 
    ((int*)rc)[threadIdx.x] = ((int*)r)[gidin + threadIdx.x]; 
    ((int*)gc)[threadIdx.x] = ((int*)g)[gidin + threadIdx.x]; 
    ((int*)bc)[threadIdx.x] = ((int*)b)[gidin + threadIdx.x]; 

    __syncthreads(); 

    /* Shared memory for output */ 
    __shared__ uchar* roc[256]; 
    __shared__ uchar* goc[256]; 
    __shared__ uchar* boc[256]; 

    /* Do the subsampling, one pixel per thread. Store into the output shared memory */ 

    ... 

    __syncthreads(); 

    /* Finally, write the result to global memory with coalesced stores */ 
    if (threadIdx.x < 64) { 
     ((int*)ro)[gid + threadIdx.x] = ((int*)roc)[threadIdx.x]; 
    } else if (threadIdx.x < 128) { 
     ((int*)go)[gid + threadIdx.x-64] = ((int*)goc)[threadIdx.x-64]; 
    } else if (threadIdx.x < 192) { 
     ((int*)bo)[gid + threadIdx.x-128] = ((int*)boc)[threadIdx.x-128]; 
    } 
}

嗚！很多東西在那裏，遺憾的代碼轉儲。一些要記住的原則：

1）當您使用合併加載/存儲時，內存很快。這意味着每個線程在一個32的warp中，每個線程訪問32個字節。如果32字節索引與warp中的線索索引相匹配，則所有32個訪問都被放入一個128個事務中。這就是你如何獲得GPU的100GB/s帶寬。

2）進行二次採樣時的存儲器訪問模式不合並，因爲它依賴於原始存儲器不具備的二維空間局部性。（也可以爲此使用紋理內存......）通過將輸入存儲在共享內存中，然後進行處理，可最大限度地降低對計算性能的影響。

我希望這會有所幫助 - 如果您願意，我可以在某些部分回覆更多細節。

來源

2012-04-14 23:09:47

CUDA版本比CPU版本慢嗎？

回答

相關問題