ubuntu 11.04 tesla c2050 cuda sdk 4.0の環境でdoubleのアトミック演算が正しく動作しません オプションは -arch=compute_20を指定してます。
本家のプログラミングガイドに書いてある、以下のコードを実行してます。 __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); }
引数のadrees*には共有メモリのアドレスを渡してます。 一つのスレッドのみが共有メモリにアクセスする場合は正しく加算は行われますが、二つ以上のスレッドが衝突する場合は 正しく動作しません。printfで見てみたら、二つ目のスレッドがループを抜けれずに何度も共有メモリの値に 加算しているようでした。上の文をdouble→float,long long → int に書き換えた場合は正しく動作しました。 また、atomicCASのところを二つのprintf、 printf("%d %f %f\n",threadIdx.x,__longlong_as_double(assumed) ,__longlong_as_double(old)); ではさむとなぜか正しくループを抜け、正しくsharedメモリに値を書き込みます。
