4

我目前正在尝试通过内存合并在 OpenCl 中转置矩阵。

我已经以“简单”的方式转换了矩阵，效果很好。当我现在尝试用内存合并做同样的事情时，我希望看到执行时间有一点改进，但​​我的实现实际上比简单的实现要慢（实现是正确的，只是效率不高）。我想我还没有完全理解如何确保水平相邻的工作项写在水平相邻的地址上。

这是我的联合实现的内核：

__kernel void MatrixTranspose(__global const float* Matrix, 
__global float* MatrixTransposed, uint Width, uint Height, __local float* block) {

    int2 GlobalID;
    GlobalID.x = get_global_id(0);
    GlobalID.y = get_global_id(1);

    int2 LocalID;
    LocalID.x = get_local_id(0);
    LocalID.y = get_local_id(1);

    block[LocalID.y*get_local_size(0) + LocalID.x] = Matrix[GlobalID.y*Width + GlobalID.x];

    barrier(CLK_LOCAL_MEM_FENCE);

    int2 groupId;
    groupId.x = get_group_id(0);
    groupId.y = get_group_id(1);
    int2 localSize;
    localSize.x = get_local_size(0);
    localSize.y = get_local_size(1);
    MatrixTransposed[Height*(LocalID.x + groupId.x*localSize.x) + Height - (LocalID.y + groupId.y*localSize.y) - 1] = block[LocalID.y*localSize.x + LocalID.x];
}

希望有人能给我建议，谢谢：）

4

2 回答 2

2

不幸的是，您将受制于设备的全局读写速度。通常你转置矩阵来做一些计算，这有助于隐藏延迟。在您的示例中，您正在读取本地内存，等待屏障，并将黑色写入全局。这只会增加使用本地内存的额外步骤和复杂性。

如果您想隐藏全局内存延迟，您应该在数据位于本地内存中时对其进行处理。

如果您只想转置矩阵，只需从全局读取并直接写入全局中的目标位置。如果您仍想尝试使用本地内存，可以查看 async_work_group_copy 。

现在我的答案。

尝试让一个工作项负责多个浮点数。如果您读取带有工作项的 4x4 区域，则可以将其转置到私有内存中。这不仅会跳过本地内存，而且会消除对屏障的需求，并将所需的工作项数量减少 16 倍。

脚步：

  • 计算 src 和 dest 全局内存地址
  • 从全局加载四个 float4 值
  • 通过相应地交换它们的 w,x,y,z 值来转置 4x4 浮点数
  • 在全局内存的新位置存储 4 个 float4 值
  • 在单独的内核中处理矩阵的边缘区域，或者在主机程序中处理具有非四维矩阵的矩阵（或填充输入矩阵以使其成为 4 的倍数）
于 2014-08-19T14:10:18.170 回答
1

从原始矩阵中读取列，将它们作为行写入本地内存以避免内存库冲突，然后将本地内存中的行存储到转置矩阵中。

在这种情况下，两个写操作很容易合并（相邻的工作项写入相邻的存储单元）。然而，读取操作并不是那么好。

顺便说一句，你的设备是什么？如果对向量操作很好，使用 vload/vstore 操作，它可能会显着提高 IO 性能。

于 2014-08-19T09:52:09.420 回答