openCL 辅助函数调用不起作用

Question

我正在编写一个 openCL 内核来为益智游戏执行基于蛮力的 AI，但我的内核代码和/或它调用的辅助函数存在问题。这是我的内核代码（我确信输入正确地传递到这里）：60 是 clEnqueueNDRangeKernel 设置的全局工作大小。

内核的输入如下：

__global char * in //用于测试目的的虚拟输入

__global char * board_in, // 一个包含 60 个板的大字符数组

__global int *lookup, // 一个数组，我用来快速获取得分动作的分数

输出：

__global char * out, //用于测试的虚拟输出

__global int * score_out, //60 个分数的数组：每个棋盘一个

__global int * row_out, // 一个 60 行的数组：每个评估的板一个

__global int * col_out // 60 列的数组：...

__kernel void helloworld(__global char* in,
                    __global char* board_in,
                    __global int* lookup,
                    __global char* out, 
                    __global int * score_out,
                    __global int * row_out,
                    __global int * col_out)
{

    int num = get_global_id(0);
    char workingBoard[72];
    int scoreMat[64];
//set up the array for each thread to use
    for(int k=0; k< 72; k++)
    {
        workingBoard[k] = board_in[num*BOARDSIZE+k];
    }
// Make a copy of the score matrix for each thread to use
    for(int j=0; j<64; j++)
    {
        scoreMat[j] = lookup[j];
    }
    int s=0;
    int r=0;
    int c=0;
    findBestMove(workingBoard,scoreMat,&s,1,&r,&c);
    col_out[num] = ?????????
    score_out[num] = ???????????
    row_out[num] = ???????????????
}

findBestMove 函数的工作原理是这样的（它经过了很好的测试。我已经在 CPU 实现中使用了一段时间）： 它需要一个 Board（字符数组）、一个得分查找数组、一个指向移动得分的指针、当前深度以及指向行和列的指针。它应该设置分数，行和列。它调用我在同一个文档中定义的其他函数。

如果我在 CPU 上运行这段代码 sn-p，我会得到正确的输出：

// workerBoard and lookuparr are set previous to this to be the same as what
//the kernel thread is supposed to have
int s=0;
int r=0;
int c=0;
findBestMove(workerBoard,lookuparr,&s,1,&r,&c);
cout<<s<<","<<r<<","<<c<<endl;

当我运行我的内核代码时，我不会让它通过函数调用。该函数与内核在同一文档中定义，并且不使用动态内存、函数指针、递归或全局内存（在内核参数之外）。我确实使用了一些#define 语句。

我想设置 ????我的内核部分是 r、c 和 s，但如前所述，我没有到达那里。我是否犯了任何严重错误（注意：内核通过了我的代码检查器和 AMD 的内核分析器）。另外，我对 openCL 还很陌生，所以也欢迎任何提示。如果我可以提供更多信息来帮助回答这个问题，请告诉我！

Answer 1

根据您的 cmets 似乎问题出在您的 findBestMove 函数中。顺便说一句，如果你有一个无限循环，看门狗会在某一时刻触发，你的驱动程序很可能会崩溃，导致黑屏或死机。

因此，我建议您在函数中注释所有代码，并为 r、s、c 变量分配一个选定的值，例如使用 get_global_id 函数处理这些特定变量的工作项 id。当然要换？？？与：

  col_out[num] = c;
  score_out[num] = s;
  row_out[num] = r;

如果你得到正确的值，开始调试你的函数，你会确定你的问题出在函数中。

既然您在这里询问了一些提示，我认为这会提高性能（一旦您修复了错误 :)）：不要将私有内存用于 scoreMat 数组，而是使用本地内存。这样做可以避免让每个线程一遍又一遍地访问全局内存中的相同数据（这很慢）。要从全局内存中获取数据到本地内存，您可以使用async_work_group_copy 函数。

所以你的情况是这样的：

local int scoreMat[64];
event_t ev = async_work_group_copy(lookup, scoreMat, 64, 0); 
// Wait to make sure everything is copied   
wait_group_events (1, &ev);

您可能需要更改更多代码以考虑您现在使用的本地内存。基本上它的工作方式与全局方式相同（从访问的角度来看），但速度要快得多。

请注意，与您所拥有的不同之处在于只会制作一份副本，而不是 60 个（工作项的数量）。这一次，您从全局获取的数据可以从工作组中的所有工作项访问。在每个工作项都有自己的副本之前。强调在工作组中这一事实也很重要。但由于您只使用 60 个工作项，您很可能只有一个工作组。