在 CUDA 中测试无穷大答案

【问题标题】：Testing for infinity in CUDA在 CUDA 中测试无穷大
【发布时间】：2010-12-03 22:04:42
【问题描述】：

在 CUDA 程序中，我最近从使用无穷大测试切换

return x==INFINITY || x==-INFINITY;

INFINITY 是从 math.h 到

return !isfinite(x);

并且对得到不同的结果感到非常惊讶。 gnu.org 建议他们实际上应该表现得相似。我错过了什么吗？ CUDA内核中不允许使用INFINITY吗？

编辑： 我刚刚发现 isinf 并注意到使用检查

return isinf(x);

给出与 INFINITY 检查相同的结果。为什么不是 isfinite(x)==!isinf(x)?

【问题讨论】：

结果有什么不同？结果如何？预期的结果是什么？
我还不确定，但似乎 isfinite 比 INFINITY 测试更严格。
@stephen canon：授予，但 Cg 文档没有提到 math.h 中的 INFINITY 宏，这是我要问的区别。显然 INFINITY 宏确实像 isinf 一样工作（请参阅我的编辑），因此它具有一定的相关性。

标签： c++ c cuda numerical

【解决方案1】：

isfinite(a) 与!isnan(a) && !isinf(a) 相同。如果x 是NaN，那么isfinite(x) 和isinf(x) 都是假的。

【讨论】：

我对 NaN (x!=x) 的检查通过了，所以根据这个测试，没有 NaN。但是，如果 Crashworks 是正确的，那么这实际上可能是解决方案。
@hannes：您使用的是什么 GPU 和 CUDA 版本？
@hannes：确保 (x!=x) 没有被您的编译器优化掉。
@Billy ONeal：手动注入 NaN 时，nan-test 可以正常工作。

【解决方案2】：

isinf() 仅检查 +INFINITY 或 -INFINITY。

!isfinite() 检查+INFINITY、-INFINITY 或NaN。

【讨论】：

【解决方案3】：

浮点比较不一定有效。例如，(1.0f + 3.0f != 2.0f + 2.0f) 有可能。 isfinite 完全有可能认为小于特定常量的值等于 INFINITE 或 -INFINITE，而您编写的是文字相等。

【讨论】：

1.0f + 3.0f != 2.0f + 2.0f 在任何具有 IEEE-754 算法的系统上（甚至在不完全符合 IEEE-754 的系统上，如一些 GPU）。浮点不是黑魔法。对于何时舍入和何时不舍入，有明确的规则。

【解决方案4】：

许多 GPU 和 SIMD 单元并不完全符合 IEEE754 标准，尤其是对于无穷大和 NaN 周围的边缘情况。就在昨晚，我注意到我使用的一个特定向量处理器声称 ∞+1 ≠ ∞ ，并且 x == x 即使对于 x &in;南。

【讨论】：

【解决方案5】：

在最近的帖子 Checking if a matrix contains nans or infinite values in CUDA 中，Robert Crovella 建议使用 isinf() 检查 CUDA 中的无限值。

下面我提供了一个使用isinf() 并利用CUDA Thrust 检查数组中无限值的示例。也许它可以作为其他用户的参考。下面的例子等价于 Matlab 的d_result=isinf(d_data);。它与我为上面引用的问题发布的示例不同，当前一个检查每个单独的元素是无限的，而另一个检查整个数组是否包含至少一个 NaN 并且等效于 Matlab 的 sum(isnan(d_data));。

#include <thrust/sequence.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust\device_vector.h>
#include <thrust\reduce.h>

#include <float.h>

// --- Operator for testing inf values
struct isinf_test { 
    __host__ __device__ bool operator()(const float a) const {
        return isinf(a);
    }
};

void main(){

    const int N = 10;

    thrust::host_vector<float> h_data(N);
    for (int i=0; i<N; i++)
        h_data[i] = rand()/RAND_MAX;

    h_data[0] = FLT_MAX/FLT_MIN;

    thrust::device_vector<float> d_data(h_data);
    thrust::device_vector<float> d_result(h_data);

    thrust::transform(d_data.begin(), d_data.end(), d_result.begin(), isinf_test());

    for (int i=0; i<N; i++) {
        float val = d_result[i];
        printf("Isinf test for element number %i equal to %f\n",i,val);
    }

    getchar();

}

【讨论】：