在 CUDA 中测试无穷大

Question

在 CUDA 程序中，我最近从使用

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

来自 math.h 的 INFINITY 来测试无穷大，转换为使用来自 math.h 的无穷大测试，

return !isfinite(x);

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

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

return isinf(x);

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

Answer 1

isfinite(a) 与 !isnan(a) && 相同!isinf(a)。如果 x 为 NaN，则 isfinite(x) 和 isinf(x) 都为 false。

Answer 2

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

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

Answer 3

浮点比较不一定有效。例如，可能是 (1.0f + 3.0f != 2.0f + 2.0f)。 isfinite 完全有可能将小于特定常量的值视为等于 INFINITE 或 -INFINITE，而你写的是字面意义上的平等。

Answer 4

许多 GPU 和 SIMD 单元并不完全符合 IEEE754，特别是对于无穷大和 NaN 的边缘情况。就在昨晚，我注意到我使用的一个特定矢量处理器声称 ∞+1 ≠ ∞ ，并且 x == x 即使对于 x ∈ NaN 也是如此。

Answer 5

在最近的帖子中 检查是否矩阵在 CUDA 中包含 nan 或无限值，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();

}