CUDAに次のテンプレート__device__関数があります。
template<typename T>
__device__ void MyatomicAdd(T *address, T val){
atomicAdd(address , val);
}
T を float としてインスタンス化した場合、つまりコンパイルして正常に実行されます。
__global__ void myKernel(float *a, float b){
MyatomicAdd<float>(a,b);
}
問題なく実行されます。
double には対応していないため、この関数を特化したかったatomicAdd()ので、倍精度で実装をハンドコーディングできます。今のところ倍精度の特殊化を無視すると、単精度の特殊化とテンプレートは次のようになります。
template<typename T>
__device__ void MyatomicAdd(T *address, T val){
};
template<>
__device__ void MyatomicAdd<float>(float *address, float val){
atomicAdd(address , val);
}
今、コンパイラは、atomicAdd() が特殊化で未定義であると不平を言います。特殊化内で __syncthreads() のような CUDA 関数を使用しようとすると、同じことが当てはまります。何か案は？ありがとう。