4

全体として、単純な平均化を使用して RGB 画像をグレースケールに変換する非常に単純な OpenCL カーネルを作成しました。

背景:

  1. イメージは、24 ビットのパディングされていないメモリ ブロックとして、マップされたメモリに格納されます。
  2. 出力配列は固定メモリ ( でマップclEnqueueMapBuffer) に格納され、8 bpp です。
  3. clCreateBufferデバイス ( ) には 2 つのバッファーが割り当てられていますclWriteBuffer。clReadBuffer

これを 1280x960 の画像で実行しています。アルゴリズムのシリアル バージョンは平均 60 ミリ秒、OpenCL カーネルは平均 200 ミリ秒です!!! 私は何か間違ったことをしていますが、どのように進めるべきか、何を最適化すべきかわかりません。(カーネル呼び出しなしで読み取り/書き込みのタイミングを計ると、アルゴリズムは 15 ミリ秒で実行されます)

カーネルのセットアップ (サイズと引数) とカーネルを添付しています

編集: それで、内部でグローバルメモリアクセスを行わないさらに愚かなカーネルを書きましたが、それはわずか150ミリ秒でした...これはまだ途方もなく遅いです。グローバルメモリの読み取りを台無しにしているのではないかと思いました.4バイトに揃える必要がありますか？いいえ...

編集 2:カーネルからすべてのパラメーターを削除すると、速度が大幅に向上しました...混乱しているのでclEnqueueWriteBuffer、カーネルはホスト->デバイスおよびデバイス->ホストからのメモリ転送を行うべきではないと思いました... .

編集 3:理解しましたが、理由はまだわかりません。誰かがそれを説明できるなら、私は彼らに正しい答えを与えて喜んでいます. 問題は、カスタム構造体を値で渡すことでした。cl_memそれらにグローバルメモリの場所を割り当てて、それらを渡す必要があるようです

カーネルコール:

//Copy input to device
result = clEnqueueWriteBuffer(handles->queue, d_input_data, CL_TRUE, 0, h_input.widthStep*h_input.height, (void *)input->imageData, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to write to input buffer on device!")) return 0;

//Set kernel arguments
result = clSetKernelArg(handles->current_kernel, 0, sizeof(OpenCLImage), (void *)&h_input);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set input struct.")) return 0;
result = clSetKernelArg(handles->current_kernel, 1, sizeof(cl_mem), (void *)&d_input_data);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set input data.")) return 0;
result = clSetKernelArg(handles->current_kernel, 2, sizeof(OpenCLImage), (void *)&h_output);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set output struct.")) return 0;
result = clSetKernelArg(handles->current_kernel, 3, sizeof(cl_mem), (void *)&d_output_data);
if(check_result(result, "opencl_rgb_to_gray", "Failed to set output data.")) return 0;

//Determine run parameters
global_work_size[0] = input->width;//(unsigned int)((input->width / (float)local_work_size[0]) + 0.5);
global_work_size[1] = input->height;//(unsigned int)((input->height/ (float)local_work_size[1]) + 0.5);

printf("Global Work Group Size: %d %d\n", global_work_size[0], global_work_size[1]);

//Call kernel
result = clEnqueueNDRangeKernel(handles->queue, handles->current_kernel, 2, 0, global_work_size, local_work_size, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to run kernel!")) return 0;

result = clFinish(handles->queue);
if(check_result(result, "opencl_rgb_to_gray", "Failed to finish!")) return 0;

//Copy output
result = clEnqueueReadBuffer(handles->queue, d_output_data, CL_TRUE, 0, h_output.widthStep*h_output.height, (void *)output->imageData, 0, 0, 0);
if(check_result(result, "opencl_rgb_to_gray", "Failed to write to output buffer on device!")) return 0;

カーネル：

typedef struct OpenCLImage_t
{
    int width;
    int widthStep;
    int height;
    int channels;
} OpenCLImage;

__kernel void opencl_rgb_kernel(OpenCLImage input, __global unsigned char*  input_data, OpenCLImage output, __global unsigned char * output_data)
{
    int pixel_x = get_global_id(0);
    int pixel_y = get_global_id(1);
    unsigned char * cur_in_pixel, *cur_out_pixel;
    float avg = 0;

    cur_in_pixel = (unsigned char *)(input_data + pixel_y*input.widthStep + pixel_x * input.channels);
    cur_out_pixel = (unsigned char *)(output_data + pixel_y*output.widthStep + pixel_x * output.channels);

    avg += cur_in_pixel[0];
    avg += cur_in_pixel[1];
    avg+= cur_in_pixel[2];
    avg /=3.0f;

    if(avg > 255.0)
        avg = 255.0;
    else if(avg < 0)
        avg = 0;

    *cur_out_pixel = avg;
}
4

2 に答える 2

5

作成されるすべてのスレッドに値をコピーするオーバーヘッドが、時間の原因として考えられます。グローバルメモリに関しては、他の場合には参照で十分です。SDK実装者だけが正確に答えることができます.. :)

于 2013-04-04T06:04:23.137 に答える
0

メモリ呼び出しを結合するために、[64, 1, 1] のような local_work_size を試すことができます。(64 は 1280 の割り算であることに注意してください)。

前述のように、より多くの情報を取得するにはプロファイラーを使用する必要があります。nvidia カードを使用していますか? 次に、openCL プロファイラーが含まれている CUDA 4 (5 ではありません) をダウンロードします。

あなたのパフォーマンスは最適とは程遠いものでなければなりません。ローカル ワーク サイズ、グローバル ワーク サイズを変更し、トレッドごとに 2 つまたは 4 つのピクセルを処理してみてください。治療前にピクセルを保存する方法を変更できますか? 次に、メモリアクセスをより効果的に結合するために、ツリー配列の構造体を分割します。

Tou は、GPU 作業でメモリ転送を隠すことができます。近くにあるプロファイラーを使用すると、より簡単に行うことができます。

于 2013-04-06T10:11:57.347 に答える