-1

次の方法で様式化できるデータ処理タスクがあります。私はdata(~1-10GB) と、これといくつかの (double) inputsummaryに基づいて (~1MB) を生成する関数を持っています。の ~1000 の値に対してこれを取得する必要があり、これは GPU にとって完璧なタスクのように見えました。繰り返しますが、入力はすべてのスレッドで同じであり、直線的に読み取られますが、各スレッドは独自の を生成する必要があります。関数は、異なる に対して個別に実行されます。dataxsummaryxdatasummaryx

ただし、オン CPU のすべての値をブルート ワンスレッドで循環させてxも、K520 よりも 3 倍悪いパフォーマンスしか得られません。これがメモリ集約型のタスクであることは理解していますが (スレッドは彼のランダムな部分にアクセスして書き込む必要がありますsummary)、GPU が最初の 1000 倍の利点を失う可能性があることを理解するのにまだ苦労しています。dataメモリを使用して to フィードをチャンクでフィードしようとしましたが__constant__(すべてのスレッドで同じ入力であるため)、目に見える改善はありませんでした。nvprof によって報告される典型的なブロック実行時間は 10 ～ 30 秒です。

このタスクに適した最適化についての洞察をいただければ幸いです。

編集: 以下は、問題を再現するサンプル コードです。g++ (5 秒のレポート実行時間) と nvcc (7 秒のレポート実行時間) の両方でコンパイルできます。プロファイリング結果は以下の通り

==23844== プロファイリング結果:
Time(%) Time Calls Avg Min Max Name
98.86% 4.68899s 1 4.68899s 4.68899s 4.68899s Kernel(Observation*, int*, Info**)
1.09% 51.480ms 4 12.870ms 1.9200us 50.426ms [CUDA memcpy HtoD]
0.06% 2.6634ms 800 3.3290us 3.2950us 5.1200us [CUDA memcpy DtoD]
0.00% 4.3200us 1 4.3200us 4.3200us 4.3200us [CUDA memcpy DtoH] 

#include <iostream>
#include <fstream>
#include <cstdlib>
#include <ctime>
#include <cstring>

#define MAX_OBS 1000000
#define MAX_BUCKETS 1000

using namespace std;

// Cross-arch defines
#ifndef __CUDACC__

#define GPU_FUNCTION

#define cudaSuccess 0

typedef int cudaError_t;

struct dim3
{
    int x;
    int y;
    int z;
} blockIdx, threadIdx;

enum cudaMemcpyKind
{
    cudaMemcpyHostToDevice = 0,
    cudaMemcpyDeviceToHost = 1, 
    cudaMemcpyDeviceToDevice = 2
};

cudaError_t cudaMalloc(void ** Dst, size_t bytes)
{
    return !(*Dst = malloc(bytes));
}

cudaError_t cudaMemcpy(void * Dst, const void * Src, size_t bytes, cudaMemcpyKind kind)
{
    return !memcpy(Dst, Src, bytes);
}

#else
#define GPU_FUNCTION __global__
#endif

// Basic observation structure as stored on disk
struct Observation
{
    double core[20];
};

struct Info
{
    int left;
    int right;
};

GPU_FUNCTION void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[threadIdx.x * 40 + threadIdx.y];

    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_obs[i].core[threadIdx.x] < (threadIdx.x + 1) * threadIdx.y)
            summary[d_bucket[i]].left++;
        else
            summary[d_bucket[i]].right++;
    }
}

int main()
{
    srand((unsigned int)time(NULL));

    // Generate dummy observations
    Observation * obs = new Observation [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        for (int j = 0; j < 20; j++)
            obs[i].core[j] = (double)rand() / RAND_MAX;

    // Attribute observations to one of the buckets
    int * bucket = new int [MAX_OBS];
    for (int i = 0; i < MAX_OBS; i++)
        bucket[i] = rand() % MAX_BUCKETS;

    Info summary[MAX_BUCKETS];
    for (int i = 0; i < MAX_BUCKETS; i++)
        summary[i].left = summary[i].right = 0;

    time_t start;
    time(&start);

    // Init device objects
    Observation * d_obs;                    
    int * d_bucket; 
    Info * d_summary;
    Info ** d_summaries;

    cudaMalloc((void**)&d_obs, MAX_OBS * sizeof(Observation));
    cudaMemcpy(d_obs, obs, MAX_OBS * sizeof(Observation), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_bucket, MAX_OBS * sizeof(int));
    cudaMemcpy(d_bucket, bucket, MAX_OBS * sizeof(int), cudaMemcpyHostToDevice);
    cudaMalloc((void**)&d_summary, MAX_BUCKETS * sizeof(Info));
    cudaMemcpy(d_summary, summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyHostToDevice);

    Info ** tmp_summaries = new Info * [20 * 40];
    for (int k = 0; k < 20 * 40; k++)           
        cudaMalloc((void**)&tmp_summaries[k], MAX_BUCKETS * sizeof(Info));
    cudaMalloc((void**)&d_summaries, 20 * 40 * sizeof(Info*));
    cudaMemcpy(d_summaries, tmp_summaries, 20 * 40 * sizeof(Info*), cudaMemcpyHostToDevice);
    for (int k = 0; k < 20 * 40; k++)
        cudaMemcpy(tmp_summaries[k], d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToDevice);

#ifdef __CUDACC__
    Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);
#else
    for (int k = 0; k < 20 * 40; k++)
    {
        threadIdx.x = k / 40;
        threadIdx.y = k % 40;
        Kernel(d_obs, d_bucket, d_summaries);
    }
#endif      

    cudaMemcpy(summary, d_summary, MAX_BUCKETS * sizeof(Info), cudaMemcpyDeviceToHost);

    time_t end;
    time(&end);
    cout << "Finished calculations in " << difftime(end, start) << "s" << endl;
    cin.get();
    return 0;
}

編集 2: 私は厳しい散在メモリ アクセスを並列化することによってコードを作り直してみました。簡単に言うと、私の新しいカーネルは次のようになります

__global__ void Kernel(Observation * d_obs, 
                         int * d_bucket,
                         double * values,
                         Info ** d_summaries)
{
    Info * summary = d_summaries[blockIdx.x * 40 + blockIdx.y];

    __shared__ Info working_summary[1024];
    working_summary[threadIdx.x] = summary[threadIdx.x];
    __syncthreads();

    for (int i = 0; i < MAX_OBS; i++)
    {
        if (d_bucket[i] != threadIdx.x) continue;
        if (d_obs[i].core[blockIdx.x] < values[blockIdx.y])
            working_summary[threadIdx.x].left++;
        else
            working_summary[threadIdx.x].right++;
    }
    __syncthreads();

    summary[threadIdx.x] = working_summary[threadIdx.x];
}

これには --- に 18 秒、 --- に 172 秒かかります<<<dim(20, 40, 1), 1000>>>。これ<<<dim(20,40,10), 1000>>>は、単一の CPU スレッドよりも悪く、並列タスクの数が直線的に増加します。

4

1 に答える 1

2

使用している K520 ボードには 2 つの GPU があり、それぞれに 8 つのストリーミング マルチプロセッサがあり、GPU あたり最大 160 GB/秒の帯域幅を備えていると思います。上記のコードでは、この帯域幅によって制限されるはずであり、GPU ごとに少なくとも 100 GB/s を取得することを検討する必要があります (ただし、最初は単一の GPU をターゲットにします)。打てないかもしれないし、倒せるかもしれないが、狙うには良いターゲットだ。

ブロック数

最初に行うことは、起動パラメーターを修正することです。この行：

Kernel<<<1, dim3(20, 40, 1)>>>(d_obs, d_bucket, d_summaries);

800 スレッドの 1 つの CUDA ブロックを起動していることを意味します。これは、GPU にとって十分な並列処理にはほど遠いものです。少なくともストリーミング マルチプロセッサと同じ数 (つまり 8)、できればそれ以上 (つまり 100 以上) のブロックが必要です。これにより、パフォーマンスが大幅に向上します。800 通りの並列処理は、GPU には十分ではありません。

分散書き込み

GPU は、アクセス パターンの影響をかなり受けやすい場合があります。次のコード:

summary[d_bucket[i]].left++;

summary に散在する 4 バイトの書き込みを行います。散在するメモリ トランザクションは GPU でコストがかかるため、メモリにバインドされたコードで適切なパフォーマンスを得るには避ける必要があります。この場合、私たちはそれについて何ができますか？私の意見では、解決策は並列処理を追加することです。スレッドごとに要約を作成する代わりに、ブロックごとに要約を作成します。各スレッドは範囲 のサブセットで動作し、 に0...MAX_OBSあるブロック全体の要約配列をインクリメントできますshared memory。カーネルの最後で、結果をグローバル メモリに書き戻すことができます。幸いなことに、これにより、上記の並列処理の欠如の問題も解決されます!

次は何？

この時点で、改善の余地がどれくらいあるかを測定する方法を考え出す必要があります。ピーク帯域幅にどれだけ近づいているかを調べたいと思うでしょう (移動する必要があるデータと実際に移動しているデータの両方を考慮するのが最善だと思います)。 それでも大幅にずれている場合は、可能であれば、メモリアクセスを削減し、アクセスをさらに最適化します。

于 2014-12-14T17:21:08.943 に答える