CUDA性能測試

我在寫一個簡單的CUDA程序進行性能測試。
這與矢量計算無關，但僅用於簡單（並行）字符串轉換。CUDA性能測試

#include <stdio.h> 
#include <string.h> 
#include <cuda_runtime.h> 


#define UCHAR   unsigned char 
#define UINT32   unsigned long int 

#define CTX_SIZE  sizeof(aes_context) 
#define DOCU_SIZE  4096 
#define TOTAL   100000 
#define BBLOCK_SIZE  500 


UCHAR   pH_TXT[DOCU_SIZE * TOTAL]; 
UCHAR   pH_ENC[DOCU_SIZE * TOTAL]; 
UCHAR*   pD_TXT; 
UCHAR*   pD_ENC; 


__global__ 
void TEST_Encode(UCHAR *a_input, UCHAR *a_output) 
{ 
    UCHAR  *input; 
    UCHAR  *output; 

    input = &(a_input[threadIdx.x * DOCU_SIZE]); 
    output = &(a_output[threadIdx.x * DOCU_SIZE]); 

    for (int i = 0 ; i < 30 ; i++) { 
     if ((input[i] >= 'a') && (input[i] <= 'z')) { 
      output[i] = input[i] - 'a' + 'A'; 
     } 
     else { 
      output[i] = input[i]; 
     } 
    } 
} 


int main(int argc, char** argv) 
{ 
    struct cudaDeviceProp xCUDEV; 

    cudaGetDeviceProperties(&xCUDEV, 0); 


    // Prepare Source 
    memset(pH_TXT, 0x00, DOCU_SIZE * TOTAL); 

    for (int i = 0 ; i < TOTAL ; i++) { 
     strcpy((char*)pH_TXT + (i * DOCU_SIZE), "hello world, i need an apple."); 
    } 

    // Allocate vectors in device memory 
    cudaMalloc((void**)&pD_TXT, DOCU_SIZE * TOTAL); 
    cudaMalloc((void**)&pD_ENC, DOCU_SIZE * TOTAL); 

    // Copy vectors from host memory to device memory 
    cudaMemcpy(pD_TXT, pH_TXT, DOCU_SIZE * TOTAL, cudaMemcpyHostToDevice); 

    // Invoke kernel 
    int threadsPerBlock = BLOCK_SIZE; 
    int blocksPerGrid = (TOTAL + threadsPerBlock - 1)/threadsPerBlock; 

    printf("Total Task is %d\n", TOTAL); 
    printf("block size is %d\n", threadsPerBlock); 
    printf("repeat cnt is %d\n", blocksPerGrid); 

    TEST_Encode<<<blocksPerGrid, threadsPerBlock>>>(pD_TXT, pD_ENC); 

    cudaMemcpy(pH_ENC, pD_ENC, DOCU_SIZE * TOTAL, cudaMemcpyDeviceToHost); 

    // Free device memory 
    if (pD_TXT)   cudaFree(pD_TXT); 
    if (pD_ENC)   cudaFree(pD_ENC); 

    cudaDeviceReset(); 
}

當我BLOCK_SIZE值更改爲2至1000，我得到了下面的持續時間（從NVIDIA視覺分析器）

TOTAL  BLOCKS  BLOCK_SIZE Duration(ms) 
100000  50000  2   28.22 
100000  10000  10   22.223 
100000  2000  50   12.3 
100000  1000  100   9.624 
100000  500   200   10.755 
100000  250   400   29.824 
100000  200   500   39.67 
100000  100   1000  81.268

我的GPU是的GeForce GT520，最大threadsPerBlock值是1024，所以我預測當BLOCK爲1000時我會獲得最佳性能，但上表顯示不同的結果。

我不明白爲什麼持續時間不是線性的，我該如何解決這個問題。（或者我怎麼能找到優化的花紋塊值（mimimum持續時間）

來源

2012-05-15 user727062

看來2，10，50個線程不利用GPU的能力，因爲它的設計是爲了啓動更加線程。

你卡具有計算能力2.1。

每個多處理器駐留的最大線程數= 1536
每塊= 1024
每個多處理器駐留的塊的最大數目= 8
經大小= 32

這裏有兩個問題：

您嘗試佔據它將definetly是外包給當地放緩每個線程這麼多寄存器內存內存空間，如果你的塊大小增加。

以32的倍數進行你的測試，因爲這是你的卡的warp大小和許多內存操作的螺紋尺寸與經線大小的倍數進行了優化。

所以，如果你只使用約1024（1000的情況下）每個塊的線程你的GPU的33％是空閒的，因爲每個SM只能分配1塊。

如果您使用以下100％佔用大小會發生什麼？

128 = 12個塊 - >因爲只有8可以駐留每平方米的塊執行是每平方米串行化
192 = 8駐地塊每平方米
256 = 6駐地塊
512 =每sm 3個常駐塊

來源

2012-05-15 15:28:50 djmj

回答

相關問題