0

私は現在、GPU を使用し、CPU で C++ を使用して、OpenCL コードのパフォーマンスを実験しています。合計 z = x + y を計算するプログラムを作成しました。ここで、z、x、y は GPU と CPU の 2 次元配列 (行列) です。これらのプログラムをテストした後、GPU と CPU の間の PCI バスでのデータ転送が遅いため、CPU は GPU よりもこの合計を計算するのにはるかに効率的であることがわかりました。ここで、CPU よりも GPU を効率的に使用するために必要な合計数を決定したいと思います。合計 z = x + y を z = x + y + y + y + y + ... などに増やすことでこれを行う予定です。

この特定の問題の合計数を増やすだけで、GPU を CPU よりも効率的に使用することは可能でしょうか?

参考までに: nVIDIA GeForce GT 640 グラフィックス カードと i5 Intel コア CPU を使用しています。

どんな助けでも大歓迎です。

編集:

以下に、私のコードを CPU に添付しました。

int main(int argc, const char * argv[])
{

    //This value determines the size of the nxn (square array)             
    int n = 1000;

    //Allocating the memory for the nxn arrays of floats.
    float **x = (float**)malloc(sizeof(float*)*n);
    float **y = (float**)malloc(sizeof(float*)*n);
    float **z = (float**)malloc(sizeof(float*)*n);


    //Initializing the arrays.
    for(int i = 0; i<n; i++){
        x[i] = (float*)malloc(sizeof(float)*n);
        y[i] = (float*)malloc(sizeof(float)*n);
        z[i] = (float*)malloc(sizeof(float)*n);

        for(int j = 0; j<n; j++){
            x[i][j] = i+j;
            y[i][j] = i+j;

        }
    }

    for(int i = 0; i<n; i++){
        for(int j = 0; j<n; j++){

            z[i][j] = x[i][j] + y[i][j];
            for(int k = 0; k < 100; k++){
                z[i][j] += y[i][j];
            }
        }
    }

    return 0;

}

OpenCL を使用した C++ は次のとおりです (データをコピーし、GPU でカーネルを実行するために使用されます)。

int n = 1000;

for(int i = 0; i<n; i++)
    {
        //Writing the data from the host to the device
        err = clEnqueueWriteBuffer(queue, d_xx, CL_TRUE, 0, sizeof(float)*n, h_xx[i], 0, NULL, NULL);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not write to buffer d_xx" << std::endl;
            exit(1);
        }

        err = clEnqueueWriteBuffer(queue, d_yy, CL_TRUE, 0, sizeof(float)*n, h_yy[i], 0, NULL, NULL);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not write to buffer d_yy" << std::endl;
            exit(1);
        }

        //Setting the Kernel Arguments
        err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_xx);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not set kernel argument h_xx." << std::endl;
            exit(1);
        }

        err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_yy);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not set kernel argument h_yy." << std::endl;
            exit(1);
        }

        err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_zz);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not set kernel argument h_zz." << std::endl;
        }

        work_units_per_kernel = n;

        //Executing the Kernel
        err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_units_per_kernel, NULL, 0, NULL, NULL);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not execute kernel." << std::endl;
            exit(1);
        }

        //Reading the Data from the Kernel
        err = clEnqueueReadBuffer(queue, d_zz, CL_TRUE, 0, n*(sizeof(float)), h_zz[i], 0, NULL, NULL);
        if(err != CL_SUCCESS){
            std::cout << "Error: Could not read data from kernel." << std::endl;
            exit(1);
        }

    }

そして最後に、GPU で実行されるカーネル コード:

__kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc)
{

    int i = get_global_id(0);

    d_cc[i] = d_aa[i] + d_bb[i];


    for(int j = 0; j < 100; j++){
        d_cc[i] += d_bb[i];
    }


}
4

2 に答える 2

2

多くのことがあなたを遅くしています:

1- グローバル メモリの使用の乱用。各グローバル メモリ アクセスは 400 倍遅く、グローバル メモリのみを使用します (200 回の読み取り/書き込みなど)。グローバル メモリは、最初に読み取り、最後に書き込むためにのみ使用する必要があり、中間値として使用することはできません。

2- あなたの N の長さは非常に短いです。CPU はわずか 1000 命令で終了しますが、GPU のすべてのレイテンシはこれよりもはるかに遅くなります。100MB のコピーは 1 バイトのコピーよりもはるかに効率的であるため、コピー操作にはオーバーヘッドがあります。

3- おそらく、CPU コードはコンパイラによって乗算に最適化されていますが、GPU コードはグローバルなどの揮発性変数にアクセスしているため、最適化できません。

4-デバイスへの/からのメモリの読み取り/書き込みは非常に高価です。これを計算に含めると、CPUが簡単に勝ちます。また、OpenCL バッファとカーネルの作成は非常に高価です。ブロッキング書き込み呼び出しも使用していることに注意してください。これは、非ブロッキング呼び出しよりもはるかに遅くなります。

于 2013-09-01T15:28:29.173 に答える
2

n = 1000*1000 の場合、コピー、操作、およびコピーバックが価値のあるポイントに達しています。DarkZero が指摘したように、グローバル メモリは最適ではないため、グローバル メモリをローカル メモリまたはスレッド メモリにキャッシュし、ローカル ワーク グループを使用できる場合、これは CPU と GPU の両方で非常に役立ちます。

カーネルから始めましょう。 d_ccグローバル メモリから 100 回参照されます。この場合の簡単な変更は、グローバル メモリをスレッド メモリにキャッシュし、最後にローカルをグローバルにコピーすることです。

 __kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc)
{

     int i = get_global_id(0);

     float t_d_cc = d_aa[i] + d_bb[i]; //make a thread only version of d_cc

     for(int j = 0; j < 100; j++){
         t_d_cc += d_bb[i];
     }

     d_cc[i] = t_d_cc; //copy the thread only back to global
} 

ハードウェアに依存する別の変更は、d_aa と d_bb をローカル メモリにキャッシュすることです。これにより、OpenCL はグローバル メモリからのバッチ コピーを利用できます。各 OpenCL デバイスにはさまざまなサイズがあり、使用できるローカル ワークグループ サイズの倍数があるため、これは少し難しい場合があります。

たとえば、i5 の最大ワークグループ サイズは 1024 で、ワークグループの倍数は 1 であるため、ローカル ワークグループは 1 から 1024 の範囲で指定できます。ATI-7970 の値はそれぞれ 256 と 64 であるため、ローカル ワークグループは64、128 などになります。これはより制限的です。

 __kernel void arraysum(__global const float *d_aa, 
                        __local float *l_d_aa,
                        __global const float *d_bb,
                        __local float *l_d_bb, 
                        __global float *d_cc,
                        __local float *l_d_cc)
{

//In this example, the global_id(1) is the number of rows and global_id(0) is the columns
//So when the kernel is called, the local work group size needs to be the size of the 
//number of columns

int i = get_global_id(1)*get_global_size(0) + get_global_id(0); //Index of the row
int j = get_local_id(0); 

l_d_aa[get_local_id(0)] = d_aa[i];
l_d_bb[get_local_id(0)] = d_bb[i];

read_mem_fence(CLK_LOCAL_MEM_FENCE);

float l_d_cc[get_local_id(0)] = l_d_aa[get_local_id(0)] + l_d_bb[get_local_id(0)]; 

for(int j = 0; j < get_global_size(0); j++){
    l_d_cc[get_local_id(0)] += l_d_bb[j];
}

d_cc[i] = l_d_cc[get_local_id(0)]; //copy the thread only back to global

}

アルゴリズムが間違っていたら申し訳ありませんが、グローバル メモリをローカル メモリにキャッシュする方法が伝わることを願っています。繰り返しになりますが、i5 では、ローカル ワークグループ サイズは 1 ~ 1024 にすることができますが、ATI7970 では列サイズが 64、128 などに制限されています。

概念的にははるかに困難ですが、このアプローチを使用すると、OpenCL のパフォーマンスははるかに向上します。

コミュニティの皆さん、お気軽にカーネルをクリーンアップしてください。

于 2013-09-01T16:45:33.717 に答える