0

私は現在、Nvidia推力ライブラリを使用してGPUで積分ヒストグラムを計算するコードを書いています。

したがって、私は常にカスタムファンクターで更新するデバイスメモリの連続ブロックを割り当てます。

問題は、デバイスメモリへの書き込みが非常に遅いことですが、読み取りは実際には問題ありません。

基本的な設定は次のとおりです。

struct HistogramCreation
{
    HistogramCreation(
    ...
    // pointer to memory
    ...
    ){}

    /// The actual summation operator
    __device__ void operator()(int index){
       .. do the calculations ..
       for(int j=0;j<30;j++){

       (1)  *_memoryPointer =  values (also using reads to such locations) ;

       }
  }
}

void foo(){

  cudaMalloc(_pointer,size);

  HistogramCreation initialCreation( ... _pointer ...);
  thrust::for_each(
    thrust::make_counting_iterator(0),
    thrust::make_counting_iterator(_imageSize),
    initialCreation);
}

（1）の書き方を次のように変更した場合>

unsigned int val = values;

パフォーマンスははるかに優れています。これは私が持っている唯一のグローバルメモリ書き込みです。

メモリ書き込みを使用すると、HDフッテージで約2秒かかります。ローカル変数を使用すると、約50ミリ秒かかるため、約40分の1になります。

なぜこれがとても遅いのですか？どうすれば改善できますか？

4

3 に答える 3

4

@OlegTitovが言ったように、グローバルメモリを使用した頻繁なロード/ストアは可能な限り回避する必要があります。それが避けられない状況がある場合、合体したメモリアクセスは、実行プロセスが遅くなりすぎないようにするのに役立ちます。ただし、ほとんどの場合、ヒストグラムの計算は、合体したアクセスを実現するのにかなり困難です。

上記のほとんどは基本的に@OlegTitovの答えを言い換えているだけですが、NVIDIACUDAで合計を見つけることについて行った調査について共有したいと思います。実際、結果は非常に興味深いものであり、他のxcuda開発者にとって役立つ情報になることを願っています。

実験は基本的に、グローバルメモリ（1スレッド）、L2キャッシュ（アトミックオペレーション-128スレッド）、およびL1キャッシュ（共有メモリ-128スレッド）を使用して、さまざまなメモリアクセスパターンで合計を見つける速度テストを実行することでした。

使用したこの実験：Kepler GTX 680、1546コア@ 1.06GHzGDDR5256ビット@3GHz

カーネルは次のとおりです。

__global__
void glob(float *h) {
    float* hist = h;
    uint sd = SEEDRND;
    uint random;
    for (int i = 0; i < NUMLOOP; i++) {
        if (i%NTHREADS==0) random = rnd(sd);
        int rind = random % NBIN;
        float randval = (float)(random % 10)*1.0f ;
        hist[rind] += randval;
    }
}

__global__
void atom(float *h) {
    float* hist = h;
    uint sd = SEEDRND;
    for (int i = threadIdx.x; i < NUMLOOP; i+=NTHREADS) {
        uint random = rnd(sd);
        int rind = random % NBIN;
    float randval = (float)(random % 10)*1.0f ;
        atomicAdd(&hist[rind], randval);
    }
}

__global__
void shm(float *h) {
    int lid = threadIdx.x;
    uint sd = SEEDRND;

    __shared__ float shm[NTHREADS][NBIN];
    for (int i = 0; i < NBIN; i++) shm[lid][i] = h[i];

    for (int i = lid; i < NUMLOOP; i+=NTHREADS) {
        uint random = rnd(sd);
        int rind = random % NBIN;
        float randval = (float)(random % 10)*1.0f ;
        shm[lid][rind] += randval;
    }

    /* reduction here */
    for (int i = 0; i < NBIN; i++) {
        __syncthreads();
        if (threadIdx.x < 64) {
            shm[threadIdx.x][i] += shm[threadIdx.x+64][i];
        }
        __syncthreads();
        if (threadIdx.x < 32) {
            shm[threadIdx.x][i] += shm[threadIdx.x+32][i];
        }
        __syncthreads();
        if (threadIdx.x < 16) {
            shm[threadIdx.x][i] += shm[threadIdx.x+16][i];
        }
        __syncthreads();
        if (threadIdx.x < 8) {
            shm[threadIdx.x][i] += shm[threadIdx.x+8][i];
        }
        __syncthreads();
        if (threadIdx.x < 4) {
            shm[threadIdx.x][i] += shm[threadIdx.x+4][i];
        }
        __syncthreads();
        if (threadIdx.x < 2) {
            shm[threadIdx.x][i] += shm[threadIdx.x+2][i];
        }
        __syncthreads();
        if (threadIdx.x == 0) {
            shm[0][i] += shm[1][i];
        }
    }

    for (int i = 0; i < NBIN; i++) h[i] = shm[0][i];
}

出力

atom:  102656.00 shm:  102656.00 glob:  102656.00
atom:  122240.00 shm:  122240.00 glob:  122240.00
... blah blah blah ...

  One Thread: 126.3919 msec
      Atomic:   7.5459 msec
      Sh_mem:   2.2207 msec

これらのカーネル間の比率は57：17：1です。ここでは多くのことを分析できますが、L1またはL2メモリスペースを使用すると、プログラム全体の10倍以上の速度が常に得られるという意味ではありません。

そして、これがメインと他の機能です：

#include <iostream>
#include <cstdlib>
#include <cstdio>
using namespace std;

#define NUMLOOP 1000000
#define NBIN 36
#define SEEDRND 1

#define NTHREADS 128
#define NBLOCKS 1

__device__ uint rnd(uint & seed) {
#if LONG_MAX > (16807*2147483647)
    int const a    = 16807;
    int const m    = 2147483647;
    seed = (long(seed * a))%m;
    return seed;
#else
    double const a    = 16807;
    double const m    = 2147483647;

    double temp = seed * a;
    seed = (int) (temp - m * floor(temp/m));
    return seed;
#endif
}

... the above kernels ...

int main()
{
    float *h_hist, *h_hist2, *h_hist3, *d_hist, *d_hist2,
    *d_hist3;
    h_hist = (float*)malloc(NBIN * sizeof(float));
    h_hist2 = (float*)malloc(NBIN * sizeof(float));
    h_hist3 = (float*)malloc(NBIN * sizeof(float));
    cudaMalloc((void**)&d_hist, NBIN * sizeof(float));
    cudaMalloc((void**)&d_hist2, NBIN * sizeof(float));
    cudaMalloc((void**)&d_hist3, NBIN * sizeof(float));

    for (int i = 0; i < NBIN; i++) h_hist[i] = 0.0f;
    cudaMemcpy(d_hist, h_hist, NBIN * sizeof(float),
    cudaMemcpyHostToDevice);
    cudaMemcpy(d_hist2, h_hist, NBIN * sizeof(float),
    cudaMemcpyHostToDevice);
    cudaMemcpy(d_hist3, h_hist, NBIN * sizeof(float),
    cudaMemcpyHostToDevice);

    cudaEvent_t start, end;
    float elapsed = 0, elapsed2 = 0, elapsed3;
    cudaEventCreate(&start);
    cudaEventCreate(&end);

    cudaEventRecord(start, 0);

    atom<<<NBLOCKS, NTHREADS>>>(d_hist);
    cudaThreadSynchronize();

    cudaEventRecord(end, 0);
    cudaEventSynchronize(start);
    cudaEventSynchronize(end);
    cudaEventElapsedTime(&elapsed, start, end);

    cudaEventRecord(start, 0);

    shm<<<NBLOCKS, NTHREADS>>>(d_hist2);
    cudaThreadSynchronize();

    cudaEventRecord(end, 0);
    cudaEventSynchronize(start);
    cudaEventSynchronize(end);
    cudaEventElapsedTime(&elapsed2, start, end);

    cudaEventRecord(start, 0);

    glob<<<1, 1>>>(d_hist3);
    cudaThreadSynchronize();

    cudaEventRecord(end, 0);
    cudaEventSynchronize(start);
    cudaEventSynchronize(end);
    cudaEventElapsedTime(&elapsed3, start, end);

    cudaMemcpy(h_hist, d_hist, NBIN * sizeof(float),
    cudaMemcpyDeviceToHost);
    cudaMemcpy(h_hist2, d_hist2, NBIN * sizeof(float),
    cudaMemcpyDeviceToHost);
    cudaMemcpy(h_hist3, d_hist3, NBIN * sizeof(float),
    cudaMemcpyDeviceToHost);

    /* print output */
    for (int i = 0; i < NBIN; i++) {
        printf("atom: %10.2f shm: %10.2f glob:
    %10.2f¥n",h_hist[i],h_hist2[i],h_hist3[i]);
    }

    printf("%12s: %8.4f msec¥n", "One Thread", elapsed3);
    printf("%12s: %8.4f msec¥n", "Atomic", elapsed);
    printf("%12s: %8.4f msec¥n", "Sh_mem", elapsed2);

    return 0;
}
于 2012-12-21T17:54:45.143 に答える
1

変更を行った後、NVCC は多くのコードを最適化する可能性があることに注意してください。NVCC は、グローバル メモリへの書き込みが行われていないことを検出し、「不要な」コードを削除するだけです。したがって、この高速化は、グローバル ライター自体から出てくるものではない可能性があります。

実際のコード (グローバル書き込みを含むコード) でプロファイラーを使用して、アラインされていないアクセスやその他のパフォーマンスの問題があるかどうかを確認することをお勧めします。

于 2012-12-21T22:35:45.227 に答える
1

GPU コードを記述するときは、グローバル メモリへの読み書きを避ける必要があります。グローバル メモリは GPU で非常に低速です。それがハードウェアの特徴です。できる唯一のことは、グローバル メモリ内の隣接するアドレスで隣接する踏み板を読み取り/書き込みにすることです。これにより、合体が発生し、プロセスが高速化されます。ただし、一般的には、データを 1 回読み取り、処理して、1 回書き出す必要があります。

于 2012-12-21T14:21:13.497 に答える