簡單CUDA內核優化

在加快應用程序的過程中，我有一個非常簡單的內核該做類型轉換如下圖所示：簡單CUDA內核優化

__global__ void UChar2FloatKernel(float *out, unsigned char *in, int nElem){ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    if(i<nElem) 
     out[i] = (float) in[i]; 
}

全局存儲器訪問被合併，並在我的理解使用共享內存也不會有好處，因爲沒有多次讀取同一內存。是否有任何想法，如果有任何優化，可以執行加快這個內核。輸入和輸出數據已經在設備上，所以不需要主機到設備的內存拷貝。

來源

2014-01-24 Imran

你說得對，採用共享內存不會給你帶來任何beinft，因爲你還必須從全局內存加載由'in'一旦元素並將它寫入退回一次，退出。如果單個線程計算多個元素，也許你可以獲得優勢。但是你必須嘗試一下。如果你還沒有完成，你可以閱讀[「最佳實踐指南」]（http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/）。也許有一些新的提示給你。 – hubs

添加到集線器評論，嘗試使用像float4，uchar4之類的矢量數據類型，甚至讓我們知道它是否在性能上有意義。 – Meluha

您無法進一步改進。做這麼簡單的事情的核心是浪費。如果要轉換的數據將用作另一個內核的輸入，則在THAT內核上執行轉換。這將爲內核增加一些額外的計算，但會被讀取字符的I/O增益隱藏，而不是浮動。 – DarkZeros

您可以對類似代碼執行的單個最大優化是使用駐留線程並增加每個線程執行的事務數。雖然CUDA塊調度模型非常輕便，但它不是免費的，並且啓動大量包含僅執行單個內存加載和單個內存存儲的線程的塊將產生大量塊調度開銷。因此，只需啓動儘可能多的塊即可「填充」GPU的全部SM，並讓每個線程都做更多工作。

第二種顯而易見的優化是切換到128字節的內存交易的負載，這應該給你一個有形的帶寬利用增益。在費米或開普勒GPU上，這不會像第一代和第二代硬件那樣大幅提升性能。

把此共成一個簡單的基準：

__global__ 
void UChar2FloatKernel(float *out, unsigned char *in, int nElem) 
{ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    if(i<nElem) 
     out[i] = (float) in[i]; 
} 

__global__ 
void UChar2FloatKernel2(float *out, 
       const unsigned char *in, 
      int nElem) 
{ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;  
    for(; i<nElem; i+=gridDim.x*blockDim.x) { 
     out[i] = (float) in[i]; 
    } 
} 

__global__ 
void UChar2FloatKernel3(float4 *out, 
       const uchar4 *in, 
      int nElem) 
{ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;  
    for(; i<nElem; i+=gridDim.x*blockDim.x) { 
     uchar4 ival = in[i]; // 32 bit load 
     float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w); 
     out[i] = oval; // 128 bit store 
    } 
} 

int main(void) 
{ 

    const int n = 2 << 20; 
    unsigned char *a = new unsigned char[n]; 

    for(int i=0; i<n; i++) { 
     a[i] = i%255; 
    } 

    unsigned char *a_; 
    cudaMalloc((void **)&a_, sizeof(unsigned char) * size_t(n)); 
    float *b_; 
    cudaMalloc((void **)&b_, sizeof(float) * size_t(n)); 
    cudaMemset(b_, 0, sizeof(float) * size_t(n)); // warmup 

    for(int i=0; i<5; i++) 
    { 
     dim3 blocksize(512); 
     dim3 griddize(n/512); 
     UChar2FloatKernel<<<griddize, blocksize>>>(b_, a_, n); 
    } 

    for(int i=0; i<5; i++) 
    { 
     dim3 blocksize(512); 
     dim3 griddize(8); // 4 blocks per SM 
     UChar2FloatKernel2<<<griddize, blocksize>>>(b_, a_, n); 
    } 

    for(int i=0; i<5; i++) 
    { 
     dim3 blocksize(512); 
     dim3 griddize(8); // 4 blocks per SM 
     UChar2FloatKernel3<<<griddize, blocksize>>>((float4*)b_, (uchar4*)a_, n/4); 
    } 
    cudaDeviceReset(); 
    return 0; 
}

給我此小費米設備上：

>nvcc -m32 -Xptxas="-v" -arch=sm_21 cast.cu 
cast.cu 
tmpxft_000014c4_00000000-5_cast.cudafe1.gpu 
tmpxft_000014c4_00000000-10_cast.cudafe2.gpu 
cast.cu 
ptxas : info : 0 bytes gmem 
ptxas : info : Compiling entry function '_Z18UChar2FloatKernel2PfPKhi' for 'sm_2 
1' 
ptxas : info : Function properties for _Z18UChar2FloatKernel2PfPKhi 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas : info : Used 5 registers, 44 bytes cmem[0] 
ptxas : info : Compiling entry function '_Z18UChar2FloatKernel3P6float4PK6uchar4 
i' for 'sm_21' 
ptxas : info : Function properties for _Z18UChar2FloatKernel3P6float4PK6uchar4i 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas : info : Used 8 registers, 44 bytes cmem[0] 
ptxas : info : Compiling entry function '_Z17UChar2FloatKernelPfPhi' for 'sm_21' 

ptxas : info : Function properties for _Z17UChar2FloatKernelPfPhi 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas : info : Used 3 registers, 44 bytes cmem[0] 
tmpxft_000014c4_00000000-5_cast.cudafe1.cpp 
tmpxft_000014c4_00000000-15_cast.ii 

>nvprof a.exe 
======== NVPROF is profiling a.exe... 
======== Command: a.exe 
======== Profiling result: 
Time(%)  Time Calls  Avg  Min  Max Name 
    40.20 6.61ms  5 1.32ms 1.32ms 1.32ms UChar2FloatKernel(float*, unsigned char*, int) 
    29.43 4.84ms  5 968.32us 966.53us 969.46us UChar2FloatKernel2(float*, unsigned char const *, int) 
    26.35 4.33ms  5 867.00us 866.26us 868.10us UChar2FloatKernel3(float4*, uchar4 const *, int) 
    4.02 661.34us  1 661.34us 661.34us 661.34us [CUDA memset]

在後者的兩個內核，僅使用8個塊給出大的加速與4096個塊相比，這確認了每個線程有多個工作項是提高這種內存限制的低指令內核性能的最佳方法。

來源

2014-01-24 19:30:37 talonmies

+1，優秀的答案！ – JackOLantern

你有沒有關於你上面提到的那個線程調度開銷的任何參考？我不太明白的是單個內存讀/寫如何影響線程調度。而且，線程調度只能做一次，不是嗎？ –

@GregKasapidis：我沒有看到我在那個答案中提到了線程調度開銷。 – talonmies

您可以通過const __restrict__限定符修飾輸入數組，以通知編譯器數據是隻讀的而不是其他指針的別名。通過這種方式，編譯器將檢測到訪問是統一的，並且可以通過使用其中一個只讀緩存（常量緩存或計算能力≥3.5，只讀數據緩存稱爲紋理緩存）來優化訪問。

您還可以通過__restrict__限定符來修飾輸出數組以建議編譯器進行其他優化。

最後，DarkZeros的推薦值得關注。

來源

2014-01-24 17:09:26 JackOLantern

你最好寫一個你的代碼的矢量化版本，立刻寫入float4。如果nElem碰巧是4倍數的邊界，這應該是非常簡單的，否則，你可能需要介意一個殘留物。

來源

2014-01-24 19:04:54 sramij

這是一個cpu版本的函數和4個gpu內核。 3個內核來自@talonmies答案，我添加了僅僅利用矢量數據類型的kernel2。

// cpu version for comparison 
void UChar2Float(unsigned char *a, float *b, const int n){ 
    for(int i=0;i<n;i++) 
     b[i] = (float)a[i]; 
} 

__global__ void UChar2FloatKernel1(float *out, const unsigned char *in, int nElem){ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    if(i<nElem)  out[i] = (float) in[i]; 
} 

__global__ void UChar2FloatKernel2(float4 *out, const uchar4 *in, int nElem){ 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    if(i<nElem) { 
     uchar4 ival = in[i]; // 32 bit load 
     float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w); 
     out[i] = oval; // 128 bit store 
    } 
} 

__global__ void UChar2FloatKernel3(float *out, const unsigned char *in, int nElem) { 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for(; i<nElem; i+=gridDim.x*blockDim.x) 
    { 
     out[i] = (float) in[i]; 
    } 
} 

__global__ void UChar2FloatKernel4(float4 *out, const uchar4 *in, int nElem) { 
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for(; i<nElem; i+=gridDim.x*blockDim.x) 
    { 
     uchar4 ival = in[i]; // 32 bit load 
     float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w); 
     out[i] = oval; // 128 bit store 
    } 
}

在我的GeForce GT 640，這裏是時序結果：

simpleKernel (cpu):   0.101463 seconds. 
simpleKernel 1 (gpu):  0.007845 seconds. 
simpleKernel 2 (gpu):  0.004914 seconds. 
simpleKernel 3 (gpu):  0.005461 seconds. 
simpleKernel 4 (gpu):  0.005461 seconds.

所以我們可以看到kernel2只利用向量類型，就是贏家。我已經爲（32 * 1024 * 768）元素做了這些測試。nvprof輸出也顯示如下：

Time(%)  Time  Calls  Avg  Min  Max Name 
91.68% 442.45ms   4 110.61ms 107.43ms 119.51ms [CUDA memcpy DtoH] 
3.76% 18.125ms   1 18.125ms 18.125ms 18.125ms [CUDA memcpy HtoD] 
1.43% 6.8959ms   1 6.8959ms 6.8959ms 6.8959ms UChar2FloatKernel1(float*, unsigned char const *, int) 
1.10% 5.3315ms   1 5.3315ms 5.3315ms 5.3315ms UChar2FloatKernel3(float*, unsigned char const *, int) 
1.04% 5.0184ms   1 5.0184ms 5.0184ms 5.0184ms UChar2FloatKernel4(float4*, uchar4 const *, int) 
0.99% 4.7816ms   1 4.7816ms 4.7816ms 4.7816ms UChar2FloatKernel2(float4*, uchar4 const *, int)

來源

2014-01-25 10:46:57 Imran

簡單CUDA內核優化

回答

相關問題