1

パフォーマンステスト用の簡単な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 Visual Profilerから）

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はGeForceGT520で、最大threadsPerBlock値は1024なので、BLOCKが1000のときに最高のパフォーマンスが得られると予測しましたが、上の表は異なる結果を示しています。

継続時間が線形でない理由と、この問題を解決するにはどうすればよいか理解できません。（または、最適化されたブロック値（最小継続時間）を見つけるにはどうすればよいですか？

4

1 に答える 1

3

その設計ははるかに多くのスレッドを開始することであるため、2、10、50スレッドはGPUの機能を利用していないようです。

カードには計算機能2.1があります。

  • マルチプロセッサあたりの常駐スレッドの最大数=1536
  • ブロックあたりの最大スレッド数=1024
  • マルチプロセッサあたりの常駐ブロックの最大数=8
  • ワープサイズ=32

2つの問題があります：

1.1。

ブロックサイズが大きくなると、スレッドごとに非常に多くのレジスタメモリを占有しようとするため、ローカルメモリスペースを遅くするためにアウトソーシングされます。

2.2。

これはカードのワープサイズであり、多くのメモリ操作はワープサイズの倍数のスレッドサイズに最適化されているため、32の倍数でテストを実行します。

したがって、ブロックごとに約1024（この場合は1000）スレッドのみを使用する場合、SMごとに割り当てることができるブロックは1つだけなので、GPUの33％がアイドル状態になります。

次の100％占有サイズを使用するとどうなりますか？

  • 128 =12ブロック->smごとに常駐できるのは8つだけなので、ブロックの実行はシリアル化されます
  • 192=1smあたり8つの常駐ブロック
  • 256=1smあたり6つの常駐ブロック
  • 512=1smあたり3つの常駐ブロック
于 2012-05-15T15:28:50.353 に答える