OpenCL - 将结构从全局复制到本地内存答案

【问题标题】：OpenCL - copying a struct from global to local memoryOpenCL - 将结构从全局复制到本地内存
【发布时间】：2012-06-18 12:16:43
【问题描述】：

我有一个 struct Primitive，它的定义如下：

typedef struct Primitive {
    float m[12];
    float invm[12];
    enum PrimitiveType type;
    int rayDensity;
    float util1;             
    float util2;                
} Primitive;

我在一个常量内存缓冲区中将这些结构的数组传递给我的内核：

__constant Primitive *objects;

作为优化练习的一部分，我想看看将结构加载到本地内存中，所以我的内核有这样的代码：

__kernel void test(int n_objects, __constant Primitives *objects) {
    local Primitive pFrom, pTo;

    for(int i = 0; i < n_objects; i++) {
        pFrom = objects[i]; 
    }

}

当我运行它时，我得到一个编译错误：

ptxas application ptx input, line 42; error: State space mismatch between instruction and address in instruction 'ld'

作为一个实验，我尝试先将结构复制到私有变量，然后复制到局部变量，如下所示：

__kernel void test(int n_objects, __constant Primitives *objects) {
    Primitive pF, Pt;
    local Primitive pFrom, pTo;

    for(int i = 0; i < n_objects; i++) {
        pF = objects[i]
        pFrom = pF; 
    }

}

现在可以编译并运行，但是对象似乎没有被深度复制到局部变量 pFrom 中。

请注意，我的代码示例纯粹是示例，为简洁起见，我删除了所有内容。当我直接从常量全局内存中使用原始结构时，我的代码也可以正常工作。

有谁知道我在这里缺少什么，肯定是深度复制或 OpenCL 地址空间的一些基本原理。

【问题讨论】：

标签： opencl gpgpu

【解决方案1】：

您需要的是async_work_group_copy function。您可以使用 wait_group_events 函数等待此异步操作完成。

希望这会有所帮助。

【讨论】：

带有 wait_group_events 的 async_work_group_copy 函数就是答案，我很惊讶这在我的任何搜索中都没有出现。奇怪的是，将结构加载到本地内存而不是使用常量会增加运行时间。也许来自常量内存的广播比我想象的要快
您是否将下一个项目的副本和处理交错？这有很大的不同。
我已经重新调整了操作顺序，以便更早地调用 async_work_group_copy 并在所有元素上执行一些工作，然后再等待屏障并使用对象。这使得运行时间减少了一小部分（大约 4%）。但是，这种方法不会随着结构数量的增加而很好地扩展，因为它会耗尽本地内存。
为了复制一个普通的浮点数组，从全局复制到本地，我尝试用 async_work_group_copy 替换一个简单的每线程复制，async_work_group_copy 非常显着，而且速度慢得惊人。我很确定我的代码是正确的，尽管你永远不知道。也许这与您使用 async_work_group_copy 看到的缓慢性能有关？也许您可以通过将其转换为浮点数/整数数组来复制您的结构？（注意，我使用的是 nvidia k520 设备）