具有内存合并功能的 OpenCl 矩阵转置

Question

I\xc2\xb4m 目前正在尝试通过内存合并在 OpenCl 中转置矩阵。

\n\n

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

\n\n

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

\n\n

__kernel void MatrixTranspose(__global const float* Matrix, \n__global float* MatrixTransposed, uint Width, uint Height, __local float* block) {\n\n    int2 GlobalID;\n    GlobalID.x = get_global_id(0);\n    GlobalID.y = get_global_id(1);\n\n    int2 LocalID;\n    LocalID.x = get_local_id(0);\n    LocalID.y = get_local_id(1);\n\n    block[LocalID.y*get_local_size(0) + LocalID.x] = Matrix[GlobalID.y*Width + GlobalID.x];\n\n    barrier(CLK_LOCAL_MEM_FENCE);\n\n    int2 groupId;\n    groupId.x = get_group_id(0);\n    groupId.y = get_group_id(1);\n    int2 localSize;\n    localSize.x = get_local_size(0);\n    localSize.y = get_local_size(1);\n    MatrixTransposed[Height*(LocalID.x + groupId.x*localSize.x) + Height - (LocalID.y + groupId.y*localSize.y) - 1] = block[LocalID.y*localSize.x + LocalID.x];\n}\n

\n\n

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

\n

Answer 1

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

如果您想隐藏全局内存延迟，则应该在数据位于本地内存中时对数据执行某些操作。

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

现在我的答案。

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

脚步：

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