memory-management - OpenCLの問題

Question

私は OpenCL で問題に直面しています。誰かが原因についてヒントを得られることを願っています。以下は、問題を軽減したバージョンのプログラムです。サイズ 4000 の入力 int 配列があります。カーネルでスキャンを実行しています。明らかに、これを並行して行う良い方法がありますが、問題を再現するために、1 つのスレッドのみが計算全体を実行します。スキャン前は、入力 (result_mask) の値は 0 または 1 のみです。

__kernel void
sel_a(__global db_tuple * input,
      __global int * result_mask,
      __global int * result_count,
      const unsigned int max_id)
{
// update mask based on input in parallel

mem_fence(CLK_GLOBAL_MEM_FENCE);

if(gid == 0)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}
}

予想される結果は、最初に 0 以外の値を持ち、結果マスクに 5 しかなかった要素の数です。しかし、そうではありません。出力は次のようになります。

...
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 
5 5 5 5 5 5 5 5 5 0 0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 0 1 0 0 0 1 0 0 0 1 0 0 0 0 0 0 0 0 0 0 
0 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 5 5 5 5 5 5 5 5 5 5 5 
5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5
...

この 80 要素のブロックは、約 1 秒後にどこかで取得します。3200要素。常に同じ位置ではありませんが、常に同じ要素数 (80) です。さらに奇妙なことに、最初の行を if(gid == 2000) に変更すると、問題はなくなります。しかし、スレッド ID をいじってみたところ、問題が解決したのではなく、移動しただけだという結論に達しました。スレッド 1425 を使用すると、半分の時間で問題が発生し、問題が発生すると、バグのあるブロックが配列の最後にあります。したがって、0と1がない場合、ブロックはさらに「移動」したと思います。さらに興味深いことに、入力サイズを 5000 に増やすと、出力はすべて 0 で構成されます。さらに、次のコードは機能しません。

if(gid == 0)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}
if(gid == 3999)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}

一方、

if(gid == 3999)
{
    int i, c = 0;
    for(i = 0; i < max_id; i++)
    {
        if(result_mask[i]!=0)
        {
            c++;
            result_mask[i] = 5;
        }
        else
        {
            result_mask[i] = 5;
        }
    }
    *result_count = c;
}

動作します (やはり、より大きな入力では動作しない可能性があります)。デバイスの詳細は次のとおりです。

Device name: GeForce 9600M GT
Device vendor: NVIDIA
    Clock frequency:        1250 MHz
    Max compute units:      4
    Global memory size:     256 MB
    Local memory size:.     16 KB
    Max memory allocation size: 128 MB
    Max work group size:        512

明らかに、私はここで何か大きなものを見逃しています。私が最初に考えたのは、80 個の要素のブロックが別の「スレッド」によってオーバーライドされるメモリの競合であるということでした。しかし、考えれば考えるほど意味がわからなくなってきます。

ヒントをいただければ幸いです。ありがとう。

編集:応答が遅くなって申し訳ありません。そのため、問題を再現するためにコードを最小限に抑えてコードを修正しました。以下は、プログラムの C コードです。

#include <stdio.h>
#include <stdlib.h>

#include <OpenCL/openCL.h>

#define INPUTSIZE (200)

typedef struct tag_openCL
{
    cl_device_id        device;

    cl_context          ctx;
    cl_command_queue    queue;
    cl_program          program;
} openCL;

int main(void)
{
    int err;
    openCL* cl_ctx = malloc(sizeof(openCL));

    if(!cl_ctx)
        exit(1);

    err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &cl_ctx->device, NULL);

    cl_ctx->ctx = clCreateContext(0, 1, &cl_ctx->device, clLogMessagesToStdoutAPPLE, NULL, &err);

    cl_ctx->queue = clCreateCommandQueue(cl_ctx->ctx, cl_ctx->device, CL_QUEUE_PROFILING_ENABLE, &err);

    printf("Successfully created context and queue for openCL device. \n");

    /* Build program */

    char * kernel_source = "__kernel void \
sel(__global int * input, \
    __global int * result_mask, \
    const unsigned int max_id) \
{ \
    int gid = get_global_id(0); \
    \
    result_mask[gid] = input[gid] % 2 == 0; \
    result_mask[gid] &= (input[gid] + 1) % 3 == 0; \
    \
    if(gid == 0) { \
        int i; \
        for(i = 0; i < max_id; i++) { \
            if(result_mask[i]) { \
                result_mask[i] = 5; \
            } \
            else { \
                result_mask[i] = 5; \
            } \
        } \
    } \
}";

    cl_program prog = clCreateProgramWithSource(cl_ctx->ctx, 1, (const char**)&kernel_source, NULL, &err);
    cl_ctx->program = prog;

    err = clBuildProgram(cl_ctx->program, 0, NULL, NULL, NULL, NULL);

    cl_kernel kernel = clCreateKernel(cl_ctx->program, "sel", &err);

    /* create dummy input data */
    int * input = calloc(sizeof(int), INPUTSIZE);
    int k;
    for(k = 0; k < INPUTSIZE; k++)
    {
        input[k] = abs((k % 5) - (k % 3))+ k % 2;
    }

    cl_mem source, intermediate;

    unsigned int problem_size = INPUTSIZE;

    source = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);
    clEnqueueWriteBuffer(cl_ctx->queue, source, CL_TRUE, 0, problem_size * sizeof(int), (void*) input, 0, NULL, NULL);

    intermediate = clCreateBuffer(cl_ctx->ctx, CL_MEM_READ_WRITE, problem_size * sizeof(int), NULL, NULL);

    int arg = 0;
    clSetKernelArg(kernel, arg++, sizeof(cl_mem), &source);
    clSetKernelArg(kernel, arg++, sizeof(cl_mem), &intermediate);
    clSetKernelArg(kernel, arg++, sizeof(unsigned int), &problem_size);

    size_t global_work_size = problem_size;
    size_t local_work_size = 1;
    clEnqueueNDRangeKernel(cl_ctx->queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

    clFinish(cl_ctx->queue);

    // read results
    int * result = calloc(sizeof(int), problem_size );
    clEnqueueReadBuffer(cl_ctx->queue, intermediate, CL_TRUE, 0, problem_size * sizeof(int), result, 0, NULL, NULL);
    clFinish(cl_ctx->queue);


    int j;
    for(j=1; j<=problem_size; j++)
    {
        printf("%i \t", result[j-1]);
        if(j%10 ==0 && j>0)
            printf("\n");
    }

    return EXIT_SUCCESS;
}

結果はまだ非決定論的です。出力のランダムな位置で 0 と 1 を取得します。ローカルワークグループのサイズが 1 の場合、それらは配列の前半にあり、サイズが 2 の場合は配列の後半にあり、サイズが 4 の場合は後半にあり、200 要素の場合は問題ないように見えますが、ここでも 0 と 1 があります。問題のサイズは 400 です。さらに、グローバルワークグループのサイズが 1 の場合、すべて正常に動作します。つまり、2 つのカーネルを使用すると、1 つは [problem size] のグローバルワークグループサイズで並列計算を実行し、もう 1 つはグローバルワークグループサイズ 1 で並列計算を実行すると、すべてがうまく機能します。繰り返しますが、これがそれを行う方法ではないことを完全に認識しています（カーネルがそのようなシーケンシャルコードを実行している）が、何かが欠けているように見えるので、なぜそれが機能しないのか知りたい.

ありがとう、ヴァシル

score 1 · Accepted Answer

あなたの OpenCL コードは非常に単純で、結果は非常に奇妙です。問題はセットアップ部分にあると思います。バッファの作成、EnqueueNDRange の呼び出しなど。セットアップ部分を投稿していただけますか? 問題はそこにあると思います。

編集:あなたのコードを見てテストした後、最初はあなたの問題を完全に理解していないことに気付きました。あなたがマスクの更新部分にコメントしたように、私の心はその行を取り除きました。1回目で正解できたはずです。

問題は、異なるワークグループを同期できないことです。CLK_GLOBAL_MEM_FENCE は、ワークグループのメモリ順序付けアクセスに影響を与えます (グローバルメモリへの書き込みが読み戻しの前に行われることを確認します)。問題の本当の解決策は、コードを 2 つの呼び出しで実行することです。最初にマスクを並行して更新し、最初の呼び出しが終了したときに実行される別のカーネルで残りのものを実行します。続行する前に操作全体を完了する必要があるため、コマンドキューレベルでバリアを使用する必要があります。他に方法はありません。

仕様からの逐語：

OpenCL には 2 つの同期ドメインがあります。

単一のワークグループ内のワークアイテム

単一のコンテキストでコマンドキューにエンキューされたコマンド

memory-management - OpenCLの問題

1 に答える 1

Related

Reference