如何强制 OpenCL 不重新对齐结构？

Posted 2023-02-21

【中文标题】如何强制 OpenCL 不重新对齐结构？

【英文标题】：How to force OpenCL not to re-align structs?

【发布时间】：2015-11-10 10:51:10

【问题描述】：

我正在构建一个内核代码，用于捕获当前 GL 场景中的三角形。 为此，我将我的顶点流与我的索引一起发送到内核代码。

这是我的内核条目的声明：

__kernel
void CaptureTriangles(
    const uint NumTriangles,
    const float16 WorldMatrix,
    __constant ushort3 *IndexDataBlock,
    __constant struct Vertex *DataBlock,
    __global struct Triangle *TriangleBuffer,
    __global uint *TriangleBufferCount)

Vertex 结构是这样定义的：

struct Vertex

    float3  position;
    float3  normal;
    float   materialIndex;

__attribute__((packed));

现在这个流是通过 GL 创建的，这就是数据的布局方式。

获取三角形时，我在内核代码中执行以下操作：

const ushort3 idx = IndexDataBlock[get_global_id(0)];
const struct Vertex v0 = DataBlock[idx.x],
                    v1 = DataBlock[idx.y],
                    v2 = DataBlock[idx.z];

但似乎 OpenCL 一直在将 Vertex 结构重新对齐到它自己的内部要求，即使它被声明为 __attribute((packed))。 所以三角形永远不会被正确捕获。

从__constant struct Vertex *DataBlock 切换到__constant float *DataBlock 并在内核代码中显式获取每个浮点数可以解决此问题。 所以这在通过浮点数读取浮点数时有效：

// __constant float *DataBlock
float4 p0 = (float4)(DataBlock[7 * idx.x + 0], DataBlock[7 * idx.x + 1], DataBlock[7 * idx.x + 2], 1.0f),
       p1 = (float4)(DataBlock[7 * idx.y + 0], DataBlock[7 * idx.y + 1], DataBlock[7 * idx.y + 2], 1.0f),
       p2 = (float4)(DataBlock[7 * idx.z + 0], DataBlock[7 * idx.z + 1], DataBlock[7 * idx.z + 2], 1.0f);

为了代码清晰，我宁愿使用 struct Vertex 语法，有没有办法让 OpenCL 不重新对齐结构？

【问题讨论】：

酷，是的，我在其他地方也看到过。所以 float3 总是与 float4 对齐，忽略任何打包属性......我想我想要的是告诉 OpenCL 如何解释一些已经在内存中布局的数据。听起来那是不可能的，感谢您难以置信的快速响应！干杯。

【参考方案1】：

对于 CL，cl_float3 和 cl_float4 的大小相等。但在您的情况下，您的 GL 代码会提供真正的 float3 值作为输出。

__attribute__((packed)) 不会解决您的问题，因为对于 CL，结构已经被打包，只是内部元素大小不同。

恐怕你必须手动解析它。