35

atomicAdd()CUDA 4.0 以降の一部として明示的に for doubles が実装されていないのはなぜですか?

CUDA プログラミング ガイド 4.1の付録 F ページ 97 から、atomicAdd の次のバージョンが実装されました。

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val)

同じページに続いて、プロジェクトで使い始めたばかりの double 用のatomicAdd の小さな実装を次に示します。

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                             (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));
    } while (assumed != old);
    return __longlong_as_double(old);
}

上記のコードを CUDA の一部として定義しないのはなぜですか?

4

1 に答える 1

41

atomicAdd()編集: CUDA 8 の時点で、SM_6X (Pascal) GPU のハードウェア サポートを使用して、倍精度が CUDA に実装されています。

atomicAdd現在、ハードウェアでサポートされているCUDA デバイスはありませんdouble。atomicCASご指摘のとおり、64 ビット整数の観点から実装できますが、そのためには自明ではないパフォーマンス コストがかかります。

したがって、CUDA ソフトウェア チームは、正しい実装を CUDA 標準ライブラリの一部にするのではなく、開発者向けのオプションとして文書化することを選択しました。このようにして、開発者は、理解していないパフォーマンス コストを無意識のうちにオプトインすることはありません。

余談ですが、この質問を「建設的ではない」として閉じるべきではないと思います。それは完全に有効な質問だと思います、+1。

于 2012-09-27T23:59:23.287 に答える