OpenCL void 指针算术 - 奇怪的行为

Question

我编写了一个 OpenCL 内核，它使用 opencl-opengl 互操作性来读取顶点和索引，但这可能并不重要，因为我只是在做简单的指针相加，以便通过索引获取特定的顶点。

uint pos = (index + base)*stride;

这里我正在计算以字节为单位的绝对位置，在我的示例中，pos 是 28,643,328，步幅为 28，index = 0 和 base = 1,022,976。嗯，这似乎是正确的。

不幸的是，我不能直接使用vload3，因为偏移量参数不是以字节为单位的绝对地址计算的。所以我只是将pos添加到指针void* vertices_gl

void* new_addr = vertices_gl+pos;

new_addr 在我的示例中 = 0x2f90000，这就是奇怪的部分开始的地方，

vertices_gl = 0x303f000

---

结果 (new_addr) 应该是 0x4B90000 (0x303f000 + 28,643,328)

我不明白为什么地址 vertices_gl 减少了 716,800 (0xAF000)

---

我的目标是 GPU：AMD Radeon HD5830

Ps：对于那些想知道的人，我正在使用 printf 来获取这些值 :)（无法让 CodeXL 工作）

Answer 1

void* 指针没有指针算法。使用char* 指针执行逐字节指针计算。

或者比这更好：使用指针指向的 real 类型，并且不要乘以偏移量。只需编写 vertex[index+base] 假设 vertex 指向包含 28 字节数据的类型。

性能考虑：将您的顶点属性与 2 的幂对齐，以实现合并的内存访问。这意味着，在每个顶点条目之后添加 4 个字节的填充。要自动执行此操作，如果您的属性都是浮点值，请使用 float8 作为顶点类型。我假设您使用位置和普通数据或类似的数据，因此编写一个自定义结构以方便且不言自明的方式封装这两个向量可能是个好主意：

// Defining a type for the vertex data. This is 32 bytes large.
// You can share this code in a header for inclusion in both OpenCL and C / C++!
typedef struct {
    float4 pos;
    float4 normal;
} VertexData;

// Example kernel
__kernel void computeNormalKernel(__global VertexData *vertex, uint base) {
    uint index = get_global_id(0);
    VertexData thisVertex = vertex[index+base];   // It can't be simpler!
    thisVertex.normal = computeNormal(...);       // Like you'd do it in C / C++!
    vertex[index+base] = thisVertex;              // Of couse also when writing
}

注意：如果您只是将其中一个float4s 更改为float3，则此代码不适用于您的步幅为 28，因为float3 也消耗 4 个浮点数的内存。但是你可以这样写，不会加padding（但是注意这会惩罚内存访问带宽）：

typedef struct {
    float pos[4];
    float normal[3];  // Assuming you want 3 floats here
} VertexData;