1

CUDA Cの数学関数の実装（cuda/math_function.h）にacosfは、次のパッセージが含まれています。

if (__float_as_int(a) < 0) {
  t1 = CUDART_PI_F - t1;
}

ここで、およびaはt1、以前は数学定数Piに近い数値に設定されていました。条件付き（if-clause）が何をテストしているのか、それと同等のCまたは関数/マクロは何であるのかを理解しようとしています。私はの実装を探しましたが、成功しませんでした。NVIDIANVCCに組み込まれているマクロまたは関数のようです。NVCCが上記のパッセージから生成するPTXを見てください。floatsCUDART_PI_Ffloat__float_as_int(a)__float_as_int()__float_as_int()

    .reg .u32 %r<4>;
    .reg .f32 %f<46>;
    .reg .pred %p<4>;
    // ...
    mov.b32         %r1, %f1;
    mov.s32         %r2, 0;
    setp.lt.s32     %p2, %r1, %r2;
    selp.f32        %f44, %f43, %f41, %p2;

丸めで__float_as_int()はないfloatことが明らかになります。int（これにより。が生成されcvt.s32.f32ます。）代わりにfloat %f1、ビットコピー（b32）として%r1（注意：%r1はu32（unsigned int）!!）に割り当てられ、 （signed int、紛らわしい!!）で%r1あるかのように比較されます。 （誰の値は）。s32%r20

私にはこれは少し奇妙に見えます。しかし、明らかにそれは正しいです。

誰かが何が起こっているのかを説明できますか、特に__float_as_int()否定的であるためのif-clauseテストのコンテキストで何が起こっているのかを説明できますか（<0）？..そしてif節および/または__float_as_int()marcoと同等のCを提供しますか？

4

1 に答える 1

4

__float_as_intとして再解釈floatしますint。最上位ビットがオンになっているときですint。<0これはfloat、符号ビットがオンになっていることも意味しますが、数値が負であることを正確に意味するわけではありません（たとえば、「負のゼロ」になる可能性があります）。をチェックしてから、であるかどうかをチェックする方が速い場合floatがあり< 0.0ます。

C関数は次のようになります。

int __float_as_int(float in) {
     union fi { int i; float f; } conv;
     conv.f = in;
     return conv.i;
}

他のバージョンでは、このヘッダー__cuda___signbitfが代わりに使用されます。

于 2012-12-10T13:48:34.167 に答える