OpenCL (Cuda) 中的元素操作

Posted 2023-03-23

【中文标题】OpenCL (Cuda) 中的元素操作

【英文标题】：Elementwise operations in OpenCL (Cuda)

【发布时间】：2011-08-28 01:52:00

【问题描述】：

我为两个矩阵的元素乘法构建了一个内核，但至少在我的配置中，我的 OpenCL 内核只有在每个矩阵大于 2GB 时才会更快。所以我想知道，是因为我的幼稚内核（见下文）还是因为元素操作的性质，这意味着元素操作不会从使用 GPU 中获益。

感谢您的意见！

内核：

KERNEL_CODE = """
// elementwise multiplication: C = A .* B.
__kernel void matrixMul(
        __global float* C,
        __global float* A,
        __global float* B,
        int width, int height)

    // ID
    int x = get_global_id(0);
    int y = get_global_id(1);

    // Multiplying
    C[y * height + x ] = A[y * height + x] * B[y * height + x];

"""

附言我读到一些专家认为，CUDA 与 OpenCL 的差异太大，无法在同一个问题中回答两者，因此可以随意将其从标题和标签中删除。

【问题讨论】：

检查在启动内核时浪费了多少时间。你把这个和什么比较？在 cpu 上进行相同的操作？

【参考方案1】：

这种操作有 N 个 FLOP，但有 3N 个内存事务，所以它完全受内存带宽限制。没有数据重用的余地，因此比参考 CPU 版本加速的上限是 GPU 与 CPU 带宽的比率。这个数字很少超过 10 倍，并且可以很快被将数据移入和移出 GPU 内存的成本所侵蚀。一般来说，这种操作最好与其他 O(N) 操作“融合”以提高性能。您通常不会只在单个内核中计算 Hadamard 乘积，而是将其作为一个内核中一系列 O(N) 操作的一部分。所以，不，这不是一个很好的加速候选者，即使内核是最优的。

你的内核肯定不是。您为每个 FLOP 执行 3 次 IOP，这是一个巨大的惩罚。你绝对可以做一些事情来改善这一点，但什么事情将完全取决于它将在什么样的硬件上运行。

【参考方案2】：

关于元素操作：这取决于设备。例如 NVidia GPU 使用标量处理器（带有标量指令），不需要矢量化。相反，ATI 具有 5d（或 4d）VLIW 处理器，对于这些处理器来说，矢量化至关重要。但是，有时它可以由编译器执行，而不是直接在代码中使用矢量数据类型，但在针对 ATI 的 GPU 进行优化时，这是首先要做的事情。

尽管如此，正如 talonmies 所指出的，上述算法几乎不受内存带宽限制，您不能等待单独使用 GPU 来获得太多加速。

【参考方案3】：

您发布的内核至少应该和 CPU 一样快。但是您根本没有使用合并的内存访问！

这会扼杀你的表现。

但是，正如@talonmies 所说。这对于 GPU 来说不是一个好案例。你把所有的时间都浪费在记忆副本上。