如何优化 OpenCL 内核对三角函数的使用？答案

【问题标题】：How to optimize OpenCL kernel use of trigonometric functions?如何优化 OpenCL 内核对三角函数的使用？
【发布时间】：2020-06-24 01:30:04
【问题描述】：

我是 OpenCL 新手，我正在努力加快我的应用程序。 OpenCL 内核比使用顺序方法花费更多的时间。我正在尝试加密 4096 x 4096 图像。这是我写的内核：

__kernel void image_XOR(
__constant const unsigned int *inputImage,
__global unsigned int *outputImage,
__constant double *serpentineR,
__constant double *nonce,
__global unsigned int *signature) {
unsigned int i = get_global_id(0);

double decimalsPwr = pow(10.0, 15.0), serpentine2Pwr = pow(2.0, (*serpentineR));
unsigned int aux;

unsigned long long XORseq;
unsigned int decimals =  floor(decimalsPwr * fabs(*nonce));

XORseq = decimals ^ (unsigned long long) floor(( 1.0 / (i + 1)) * decimalsPwr);

if (i % 2 == 1) { 
    aux = floor(decimalsPwr * fabs( atan( 1.0 / tan( decimalsPwr * (double) XORseq))));
} else {
    aux = floor(decimalsPwr * fabs(sin(serpentine2Pwr * (double)XORseq) * cos(serpentine2Pwr * (double)XORseq)));
}

aux = aux << 8u; // comment if alfa chanel should be crypted as well
aux = aux >> 8u;

outputImage[i] = inputImage[i] ^ aux;

*signature = *signature ^ inputImage[i] ^ aux;}

注意：如果我注释掉这些行代码会快很多（0.5s 从 4s）

if (i % 2 == 1) { 
    aux = floor(decimalsPwr * fabs( atan( 1.0 / tan( decimalsPwr * (double) XORseq))));
} else {
    aux = floor(decimalsPwr * fabs(sin(serpentine2Pwr * (double)XORseq) * cos(serpentine2Pwr * (double)XORseq)));
}

【问题讨论】：

标签： c encryption parallel-processing opencl

【解决方案1】：

double decimalsPwr = pow(10.0, 15.0), serpentine2Pwr = pow(2.0, (*serpentineR));

serpentineR 用作标量，因此将其作为标量传递，而不是通过慢得多的全局内存。但在这里我会更进一步，不要在 GPU 端进行上述计算。只需在 CPU 中预先计算它们并传递给内核。想象一下，每次计算都必须执行 4096 x 4096 次 - 多么浪费资源！

unsigned int decimals = floor(decimalsPwr * fabs(*nonce)); - 相同，在 CPU 上预先计算并作为参数传递给内核。只需要计算一次。

另一个建议是尽可能避免在内核中使用 64 位类型。在大多数情况下，与 GPU 上的 32 位类型相比，它们要慢得多。让我们以 GeForce RTX 2060 为例。Wikipedia 声明单 fp 精度的处理能力为 5241.60 GFLOPS，但双 fp 精度仅为 163.80 GFLOPS。那是 32 倍的差异！如果降低精度不是一种选择，那么在 CPU 中执行多次 64 位计算并将结果传递给 GPU 以进行剩余计算是值得的。

【讨论】：

感谢您的回答。我理解您的建议，我会尝试相应地更改代码，但是，如问题中所述，无法预先计算的那两行代码需要很多时间。还有一个问题.. 将数组复制到 __local 内存中会有所帮助吗？如果是这样，应该怎么做，我试过了，效果不是很好。
这里使用本地内存似乎没有帮助。只是避免在全局内存上进行过多的读/写，例如我可以看到inputImage[i] 被读取了两次，但可以被读入一次寄存器值。有时__constant 比__global 慢，因此您可能需要更改并检查。
请注意，*signature = *signature ^ inputImage[i] ^ aux; 不是原子的，因此不会产生可预测的结果。
看起来signature应该是一个数组。
@doqtor 如果意图是一个标量加密签名，这对我来说就是这样。

【解决方案2】：

首先让我说我不熟悉这种加密方案，所以我的一些 cmets 可能没有用。

在我们尝试花很多时间优化它之前，您确定它会产生一致的结果吗？浮点精度在 OpenCL 中没有很好地定义，尤其是对于三角函数，所以如果您需要在另一个系统上解密（例如不同的 GPU 品牌），您是否获得了足够的精度？例如，您能否在 GPU 上加密后在 CPU 上解密图像？

除此之外，还有一些观察：

正如@doqtor 已经指出的那样，您有一堆值不会因工作项而异，因此请预先计算并直接传递它们。
decimals = floor(decimalsPwr * fabs(*nonce)); 看起来会溢出，假设你的nonce 接近1.0。这是故意的吗？
由于 GPU 通常以锁步方式调度线程，因此您希望相邻的工作项遵循相同的条件路径。这与您对if (i % 2 == 1) 所做的相反。我建议重新安排跨工作项的计算，使 32 或 64 个相邻的工作项遵循一条路径，而下一组遵循另一条路径。
例如，i = (i & 0xffffff80) | ((i & 0x1) << 6) | ((i & 0x3e) >> 1) 应按该顺序处理项目 0, 2, 4, 6, … 124, 126, 1, 3, 5, … 125, 127, 128, 130, …。（除非您专门处理其他情况，否则它将需要 128 个输入像素的整数倍。）
这应该可以防止所有线程执行所有可能的计算并丢弃其中的一半。
您应该能够使用一些三角恒等式来简化计算。例如，sin(θ)cos(θ) = sin(2θ)/2。这将节省您评估 cos 和 sin 的时间。
正如我在@doqtor 答案的评论中已经提到的，*signature = *signature ^ inputImage[i] ^ aux; 不是原子的，因此它不会产生可预测的结果。请改用atomic_xor()。（您可能希望收集一个工作组的签名，并且只有该组中的一个工作项更新全局签名，因为原子全局操作会带来性能损失。）
由于您受到严重的三角约束，并且许多 GPU 上的三角函数不使用 FPU 的相同部分作为乘法和加法，您可能希望尝试使用某些三角函数的显式实现，以便它们可以更好地并行化。根据您需要的精度，您还可以尝试使用查找表。

【讨论】：

感谢您的回答。你能解释一下常规or和atomic_xor之间的区别吗？另外，这只是我的应用程序的一部分。我会尽量利用POCL，它可以保证代码在各种平台上的可移植性，至少在一定程度上。我刚开始使用 OpenCL，因此，我不知道我可能遇到的所有问题。很抱歉：D
另外，i = (i & 0xffffff80) | ((i & 0x1) << 6) | ((i & 0x3e) >> 1) 应该替换 i % 2 == 1?我不明白所有这些值的来源。
*signature = *signature ^ a 执行以下操作： 1. 读取signature 指向的值 2. 与a 异或 3. 将结果写入signature 指向的内存位置。如果*signature 的内存在步骤 (1.) 和 (3.) 之间发生了变化，那么这些变化将被简单地覆盖。原子操作避免了这种情况。一般来说，您可能应该阅读 OpenCL 中的内存同步，这个主题的内容比我在 SO 评论中所能涵盖的要多得多。
i 中位的重新排列将在i = get_global_id(0); 之后立即完成。这个想法是我们颠倒了最低位和接下来的 6 个高位的顺序，以便工作项 0 到 63 的i 的低位始终为 0，即前 64 个工作项处理偶数，接下来的 64 项处理奇数。这意味着第一组 64 人将全部使用else，第二组将使用if。我建议编写一个小程序来打印数字 0 到 255 在执行此操作之前和之后，以了解我的意思。
6 位（因此 64 项）是任意选择，您可以尝试不同的值，因为它也会影响内存访问模式。

【解决方案3】：

只是一些观察...

unsigned long long XORseq;

AFAIK，OpenCL 中没有这种类型。它可能有效，但它是不可移植的代码。此外，OpenCL 类型不是 C 类型；与 C 类型不同，“unsigned long”在 OpenCL 中始终为 64 位。 IOW，OpenCL“无符号长”== C中的uint64_t；其他类型也类似。

此外，OpenCL 中的三角函数对范围和允许的结果误差有预定义要求（在 ULP 中）。这些要求非常严格，以至于大多数 GPU（尤其是消费者）根本没有硬件来使用单个硬件指令来计算它，因为（对于游戏）你 99.999% 的时间都不需要它，它只会占用硅。 GPU 确实有硬件 sin/cos 指令，但这些指令的精度和范围要有限得多——对于图形来说就足够了。您可以在 OpenCL 中将它们与“native_cos”和“native_sin”一起使用。 “完整”的 sin/cos（这是您的代码所具有的）是使用一些例程计算的 - 所以它很慢。此外，使用双精度而不是浮点数的代码会进一步减慢速度（在消费级 GPU 上减慢 8-32 倍）。

我不确定您为什么决定使用双三角法进行加密，但我怀疑它在消费级 GPU 上会非常快。

另外还有一个问题——不同的 OpenCL 实现（AMD、Nvidia、Intel 等）可以为 trigo 函数返回不同的结果； OpenCL 仅指定最大 ULP 错误。所以例如如果你尝试在 nvidia GPU opencl 上加密然后在 intel CPU opencl 上解密，你不一定会得到原来的。

【讨论】：