在opencl内核函数中声明的全局变量的副本数量是在全局地址空间中维护的

Question

我是Opencl编程的新手.为了更好地学习opencl,花了一些时间阅读一些教程后,我开始开发一个简单的模式匹配内核函数.但我有些疑惑:

首先,我在内核函数中声明了全局变量.这是否意味着每个工作项共享每个变量的单个副本？

其次,我如何使用标准C库,尤其是 "string.h中".

   __kernel void matchPatterns_V1(__global char *strings, __global char *patterns, __global int *matchCount,
                            int strCount, int strLength, int patCount, int patLength) {


    int id = get_global_id(0);
    int rowIndex = id*strLength;
    int i, matches = 0;     

    __global char *pos = strings;
    __global char *temp = strings;
    __global char *pat = patterns;

    for(i = 0; i < patCount; i++)
    {
            temp = &strings[rowIndex];      
            pat = &patterns[i*patLength];
            while(pos != '\0') {
                    pos = StrStr(temp, pat);
                    if(pos != '\0') {
                            matches++;
                            temp = pos + patLength;
                    }
            }
    }
    matchCount[id] = matches;
    }

总而言之,每个工作项都有自己的变量'pos','temp'和'pat'的副本？

任何学习Opencl的建议都受到高度赞赏,包括对最好的书籍/教程网站的建议.

Answer 1

不,它在全局内存空间中,因此通常每个内核调用一个副本,由所有工作项共享.如果不能保证每个工作组在全局内存中都有自己独特的"项" - 或者更一般地说,内核中没有两个工作项同时写入内存中的同一位置,那么写入全局内存是危险的.因为会有竞争条件.

如果您只是从这些全局内存变量中读取数据,那当然无关紧要.

您也无法__global在内核中声明变量,因为它们只能是内核参数.如果您尝试这样做,您将从OpenCL编译器获得以下内容:

error: global variables cannot be allocated inside kernel code

有充分的理由,除非技术不可能:全局变量根本没有任何意义......我能想到的唯一可能的原因是工作项之间的沟通,这将是一种疯狂的设计模式.

至于你的问题,请在评论中参考OpenCL规范第6.5节中的这个片段:

// declares a pointer p in the __private address space that
// points to an int object in address space __global
__global int *p;

因此,与指针类型相关联的内存空间表示它们指向的变量的内存空间,而不是指针本身,它们总是__private(即每个工作项).

您不能在OpenCL中使用标准C库中的字符串操作函数,但如果您愿意,可以将它们重新编码以便在GPU上使用(其中大部分都难以实现).