openCL支持向量作为内核参数吗？

问题描述 投票：0回答：1

我一直在想办法把这段代码改写成openCL内核形式。转换并不难(摆脱glm类型和位掩码)，但我被卡住的部分是如何将 _triangles, _uvs, _indices和 _normals 到内核。在openCL中是否有任何内置的向量功能？

如果没有任何向量支持的话，我看到的唯一选择是传递4个类型为 float3 我需要为3个变量分别返回(_triangles, _uvs和 _normals)和2个数组的 float3 对于 _indices. 然后在CPU中把数组转换回向量，并把它们缩小以适应。我不太确定把这么多的内存缓冲区传给内核是一个有效的方法，因为那将会有14个数组传给内核并从内核返回。我的其他解决方案在并行化的时候就不行了。有没有办法简化这个解决方案，或者更好的纯粹的解决方案？

我遇到问题的函数是 _addRectangle 和 _createMesh 是函数，它将在内核中结合。

void Chunk::_addRectangle(glm::vec3 center, glm::vec3 height, glm::vec3 width, unsigned tex_num, cl_uint LOD)
{
    glm::vec3 corner1 = center - (height / 2.0) - (width / 2.0);
    glm::vec3 corner2 = center - (height / 2.0) + (width / 2.0);
    glm::vec3 corner3 = center + (height / 2.0) + (width / 2.0);
    glm::vec3 corner4 = center + (height / 2.0) - (width / 2.0);

    glm::vec3 normal = glm::cross(height, width);

    glm::vec2 uv1;
    glm::vec2 uv2;
    glm::vec2 uv3;
    glm::vec2 uv4;

    if (fabs(normal[1]) == 1.0)
    {
        uv1 = glm::vec2(1.0 / _tex_atlas_width, 1);
        uv2 = glm::vec2(1.0 / _tex_atlas_width, 0);
        uv3 = glm::vec2(0, 0);
        uv4 = glm::vec2(0, 1);
    }
    else
    {
        uv1 = glm::vec2(1.0 / _tex_atlas_width, 1);
        uv2 = glm::vec2(1.0 / _tex_atlas_width, 0);
        uv3 = glm::vec2(0, 0);
        uv4 = glm::vec2(0, 1);
    }

    float add = (1.0 / double(_tex_atlas_width)) * tex_num;
    uv1.x += add;
    uv2.x += add;
    uv3.x += add;
    uv4.x += add;

    // triangle 1
    _triangles.push_back(corner3);
    _triangles.push_back(corner2);
    _triangles.push_back(corner1);

    _normals.push_back(normal);
    _normals.push_back(normal);
    _normals.push_back(normal);

    _uvs.push_back(uv1);
    _uvs.push_back(uv2);
    _uvs.push_back(uv3);

    _indices.push_back(glm::ivec3(nrOfIndices + 0, nrOfIndices + 1, nrOfIndices + 2));

    // triangle 2 

    _triangles.push_back(corner4);
    _normals.push_back(normal);
    _uvs.push_back(uv4);


    _indices.push_back(glm::ivec3(nrOfIndices + 2, nrOfIndices + 3, nrOfIndices + 0));
    nrOfIndices += 4;

}

void Chunk::_createMesh(glm::ivec3 pos, int landmap_flags[96 * 96 * 96], cl_int LOD)
{
    std::byte* faces = new std::byte[chunkSize / LOD * chunkSize / LOD * chunkSize / LOD];

    int index = 0;

    // a index conversion from a single index array to a 3d array
    // landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] is

    for (int x = LOD; x < chunkSize + LOD; x += LOD) {
        for (int y = LOD; y < chunkSize + LOD; y += LOD) {
            for (int z = LOD; z < chunkSize + LOD; z += LOD) {
                x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD));
                faces[index] = (std::byte)0;
                if (landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                {
                    index++;
                    continue;
                }
                if (landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] != BLOCK::AIR)
                {
                    if (landmap_flags[(x - LOD) + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::South;
                    if (landmap_flags[(x + LOD) + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::North;
                    if (landmap_flags[x + (y - LOD) * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::Down;
                    if (landmap_flags[x + (y + LOD) * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::Up;
                    if (landmap_flags[x + y * (chunkSize + (2 * LOD)) + (z - LOD) * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::West;
                    if (landmap_flags[x + y * (chunkSize + (2 * LOD)) + (z + LOD) * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))] == BLOCK::AIR)
                        faces[index] |= (std::byte)Direction::East;
                }

                if (faces[index] == (std::byte)0)
                    continue;

                if ((faces[index] & (std::byte)Direction::North) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2) + (float(LOD) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2)),
                        glm::vec3(0, LOD, 0),
                        glm::vec3(0, 0, -LOD),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                if ((faces[index] & (std::byte)Direction::East) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2) + (float(LOD) / 2)),
                        glm::vec3(0, LOD, 0),
                        glm::vec3(LOD, 0, 0),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                if ((faces[index] & (std::byte)Direction::South) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2) - (float(LOD) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2)),
                        glm::vec3(0, LOD, 0),
                        glm::vec3(0, 0, LOD),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                if ((faces[index] & (std::byte)Direction::West) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2) - (float(LOD) / 2)),
                        glm::vec3(0, LOD, 0),
                        glm::vec3(-LOD, 0, 0),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                if ((faces[index] & (std::byte)Direction::Up) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2) + (float(LOD) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2)),
                        glm::vec3(LOD, 0, 0),
                        glm::vec3(0, 0, LOD),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                if ((faces[index] & (std::byte)Direction::Down) != (std::byte)0)
                {
                    _addRectangle(
                        glm::vec3( x + ((chunkSize - 1) * pos.x) - (chunkSize / 2) - (float(LOD - 1) / 2),
                                   y + ((chunkSize - 1) * pos.y) - (chunkSize / 2) - (float(LOD - 1) / 2) - (float(LOD) / 2),
                                   z + ((chunkSize - 1) * pos.z) - (chunkSize / 2) - (float(LOD - 1) / 2)),
                        glm::vec3(LOD, 0, 0),
                        glm::vec3(0, 0, -LOD),
                        landmap_flags[x + y * (chunkSize + (2 * LOD)) + z * (chunkSize + (2 * LOD)) * (chunkSize + (2 * LOD))],
                        LOD);
                }
                index++;
            }
        }
    }

    delete[]faces;
}

谢谢

编辑: 一个可能的更有效的存储数据的方法是用几个float4类型，例如...。

const uint n = get_global_id(0);

float4 triangles{1, 2, 3, 4}; // calculated values for each vertex

//(float4 list[size];) from constructor
list[n] = triangles;
c++ kernel opencl opencl-c
1个回答
2
投票

在 OpenCL 有一些矢量类型，如 float4，等等。更多内容可以阅读 此处.没有容器像 std::vector 在 c++ 所以必须使用C式数组传递数据。

从问题中的部分代码来看。_triangles, _uvs, _indices和 _normals 将会充满结果，所以需要分配适当的缓冲区并传递给内核以存储结果，并在内核完成工作后将其读回。

传递14个数组应该不是问题，只要内核的计算量足够大，从代码来看可能是这样的，因为有两个嵌套循环。但是看起来很多事情都要看是多大的 chunkSize 和 LOD 变量。你需要尝试一下，看看它的性能如何。

将数据复制到 std::vector - 就用 memcpy.

最新问题
© www.soinside.com 2019 - 2024. All rights reserved.