CUDAは型テンプレートをサポートしており、同じコードで複数の型を処理する必要がある場合にカーネルコードを実装する最も効率的な方法であることは間違いありません。
簡単な例として、単純なBLASAXPYタイプのカーネルを考えてみましょう。
template<typename Real>
__global__ void axpy(const Real *x, Real *y, const int n, const Real a)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
for(; tid<n; tid += stride) {
Real yval = y[tid];
yval += a * x[tid];
y[tid] = yval;
}
}
このテンプレート化されたカーネルは、一般性を失うことなく、倍精度と単精度の両方でインスタンス化できます。
template axpy<float>(const float *, float *, const int, const float);
template axpy<double>(const double *, double *, const int, const double);
CUDAツールキットのすべての最新バージョンに同梱されているスラストテンプレートライブラリは、タイプに依存しないアルゴリズムを実装するためにこの機能を広範囲に使用します。