2

これは、ある種の並列リダクション/極値カーネルの一部です。それでもクラッシュする最小限のコードに減らしましたclBuildProgram（実際にはクラッシュし、エラーコードを返すだけではないことに注意してください）：

EDIT :の代わりにlocal_valueis が宣言されている場合にも発生するようです。globallocal

EDIT2 / SOLUTION :問題は、無限ループがあったことでした。remaining_items >>= 1の代わりに書くべきでしremaining_items >> 1た。回答で述べられているように、コンパイル/最適化エラーに関しては、nvidiaコンパイラはあまり堅牢ではないようです。

kernel void testkernel(local float *local_value)
{
    size_t thread_id = get_local_id(0);

    int remaining_items = 1024;

    while (remaining_items > 1)
    {
        // throw away the right half of the threads
        remaining_items >> 1; // <-- SPOTTED THE BUG
        if (thread_id > remaining_items)
        {
            return;
        }

        // look for a greater value in the right half of the memory space
        int right_index = thread_id + remaining_items;
        float right_value = local_value[right_index];
        if (right_value > local_value[thread_id])
        {
            local_value[thread_id] = right_value;
        }

        barrier(CLK_GLOBAL_MEM_FENCE);
    }
}

行を削除しreturn;たりlocal_value[thread_id] = right_value;、clBuildProgram を正常に終了させたりすることができます。

この問題は、すべてのコンピューター (NVIDIA GTX 560、GT 555M、GT 540M、すべて Fermi 2.1 アーキテクチャー) で再現できます。x64 または x86 ライブラリを使用している場合、NVIDIA CUDA Toolkit SDK バージョン 4.0、4.1、および 4.2 で明らかです。

何が問題なのか誰にも分かりますか?

ローカル（別名共有）メモリが自動的に想定される可能性はあります(WORK_GROUP_SIZE) * siezof(its_base_type)か？上記の行が削除されたときに機能する理由は、これで説明できます。

再現用の最小限のホスト コード (C99 互換):

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

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

#define RETURN_THROW(expression) do { cl_int ret = expression; if (ret) { printf(#expression " FAILED: %d\n" , ret); exit(1); } } while (0)
#define REF_THROW(expression) do { cl_int ret; expression; if (ret) { printf(#expression " FAILED: %d\n" , ret); exit(1); } } while (0)

int main(int argc, char **argv)
{
    // Load the kernel source code into the array source_str
    FILE *fp;

    fp = fopen("testkernel.cl", "rb");
    if (!fp)
    {
        fprintf(stderr, "Failed to load kernel.\n");
        exit(1);
    }
    fseek(fp, 0, SEEK_END);
    int filesize = ftell(fp);
    rewind(fp);
    char *source_str = (char*)calloc(filesize, sizeof(char));
    size_t bytes_read = fread(source_str, 1, filesize, fp);
    source_str[bytes_read] = 0;
    fclose(fp);

    // Get platform information
    cl_uint num_platforms;
    RETURN_THROW(clGetPlatformIDs(0, NULL, &num_platforms));

    cl_platform_id *platform_ids = (cl_platform_id *)calloc(num_platforms, sizeof(cl_platform_id));
    RETURN_THROW(clGetPlatformIDs(num_platforms, platform_ids, NULL));

    cl_device_id selected_device_id = NULL;

    printf("available platforms:\n");
    for (cl_uint i = 0; i < num_platforms; i++)
    {
        char platform_name[50];
        RETURN_THROW(clGetPlatformInfo(platform_ids[i], CL_PLATFORM_NAME, 50, platform_name, NULL));
        printf("%s\n", platform_name);

        // get devices for this platform
        cl_uint num_devices;
        RETURN_THROW(clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices));

        cl_device_id *device_ids = (cl_device_id *)calloc(num_devices, sizeof(cl_device_id));
        RETURN_THROW(clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_GPU, num_devices, device_ids, NULL));

        // select first nvidia device
        if (strstr(platform_name, "NVIDIA"))        // ADAPT THIS ACCORDINGLY
        {
            selected_device_id = device_ids[0];
        }
    }

    if (selected_device_id == NULL)
    {
        printf("No NVIDIA device found\n");
        exit(1);
    }

    // Create an OpenCL context
    cl_context context;
    REF_THROW(context = clCreateContext(NULL, 1, &selected_device_id, NULL, NULL, &ret));

    // Create a program from the kernel source
    cl_program program;
    REF_THROW(program = clCreateProgramWithSource(context, 1, (const char **)&source_str, NULL, &ret));

    // Build the program
    cl_int ret = clBuildProgram(program, 1, &selected_device_id, NULL, NULL, NULL);
    if (ret)
    {
        printf("BUILD ERROR\n");
        // build error - get build log and display it
        size_t build_log_size;
        ret = clGetProgramBuildInfo(program, selected_device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size);
        char *build_log = new char[build_log_size];
        ret = clGetProgramBuildInfo(program, selected_device_id, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL);
        printf("%s\n", build_log);
        exit(1);
    }

    printf("build finished successfully\n");
    return 0;
}
4

1 に答える 1

1

私の経験では、nvidia コンパイラはビルド エラーの処理に関してはあまり堅牢ではないため、おそらくどこかでコンパイル エラーが発生しています。

あなたの問題は確かにreturn、またはそれ以上の との組み合わせだと思いますbarrier。バリアに関するopencl仕様によると：

プロセッサ上でカーネルを実行するワークグループ内のすべてのワークアイテムは、バリアを超えて実行を継続できるようになる前に、この関数を実行する必要があります。この関数は、カーネルを実行しているワークグループ内のすべてのワークアイテムで遭遇する必要があります。

バリアが条件ステートメント内にある場合、いずれかの作業項目が条件ステートメントに入ってバリアを実行する場合、すべての作業項目は条件ステートメントに入らなければなりません。

barrer がループ内にある場合、すべての作業項目は、バリアを超えて実行を継続できるようになる前に、ループの反復ごとにバリアを実行する必要があります。

したがって、あなたの問題はおそらく、バリアに到達する前に多くのスレッドが返され、このコードが無効になることだと思います。たぶん、次のようなことを試してみてください：

kernel void testkernel(local float *local_value) {
    size_t thread_id = get_local_id(0);
    int remaining_items = 1024;
    while (remaining_items > 1) {
        remaining_items >>= 1;// throw away the right half of the threads
        if (thread_id <= remaining_items) {
             // look for a greater value in the right half of the memory space
             int right_index = thread_id + remaining_items;
             float right_value = local_value[right_index];
             if (right_value > local_value[thread_id])
                 local_value[thread_id] = right_value;
        }
        barrier(CLK_GLOBAL_MEM_FENCE);
    }
}

編集：さらに、コメントに記載されているように、無限ループの生成を回避するために、remaining_items>>=1代わりにする必要があります。remaining_items>>1

于 2012-07-02T12:33:17.563 に答える