OpenCL 在 GPU 上分配结构会导致垃圾答案

【问题标题】：OpenCL allocating struct on GPU results in garbageOpenCL 在 GPU 上分配结构会导致垃圾
【发布时间】：2018-12-10 11:14:29
【问题描述】：

我在 C++ 主机代码中定义了以下结构：

    struct __declspec(align(16)) MyNode {
    cl_uchar        mData;
    cl_int3         mPos;
};

在 OpenCL 中：

struct  __attribute__((aligned(16))) MyNode {
    uchar mData;
    int3 mPos;
};

现在来自主机代码我正在调用：

MyNode node= {0};
node.mPos.x = 1;
node.mPos.y = 2;
node.mPos.z = 3;

cl_mem clnode_mem = clCreateBuffer(
        mOpenCLctx, 
        CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, 
        sizeof(MyNode), 
        &node, 
        &err);

clSetKernelArg(m_kernel_test, 0, sizeof(cl_mem), &clnode_mem));

然后我在 GPU 上调用一个测试内核，定义为：

__kernel void test(__global MyNode* node)
{
    printf("pos = %d,%d,%d\n",
        node->mPos.x,
        node->mPos.y,
        node->mPos.z);
}

但我看到的输出是pos = 0,0,0，如果我从主机和设备结构定义中删除mData 成员，那么它打印正确。这里发生了什么？

我正在使用具有 x64 配置的 VS2015 进行构建，并在 NVIDIA GPU 上运行 OpenCL 1.2。

【问题讨论】：

标签： c++ opencl gpgpu gpu

【解决方案1】：

似乎可能是对齐问题。 int3 可能有 16 字节对齐，而 C 类型 cl_int3 可能只有 4 字节对齐。所以内核会在偏移量 16 处寻找它，而 CPU 在偏移量 4 处提供它。

一个好的完整性检查是在主机和内核上打印sizeof(MyNode)。如果这些不同，我建议为结构字段尝试显式数量的填充或对齐注释。

【讨论】：

【解决方案2】：

我只能建议不要在 OpenCL 中使用结构体，这不仅是因为对齐和填充问题，而且结构体数组的内存访问会变慢，因为通常在任何时间点都只访问一个字段而留下很大内存缓存中的漏洞。

既然这只是一个单一的结构，为什么不简单地将它作为内核的两个参数呢？我知道这是更多的代码，但恐怕这只是理智的做法。

【讨论】：