模板是C++的一个重要特征，它可以让我们简化代码，同时使代码更整洁。CUDA中也支持模板，这给我们编写cuda程序带来了方便。不过cuda4.0之前和之后使用模板的方法不一样，这给我们带来了少许困难。在cuda4.0之前，模板的使用和C++中无区别，使用非常方便，在此不做过多介绍。不过在cuda4.0之后，由于编译器的升级，导致之前的模板使用方法不再有效，我们需要重新设计代码。
如果按照之前的方式编写代码，如下面简单示例：
template <type T>
__global__ void foo(T *odata, T* idata)
{
extern __shared__ T sdata[];
// ... do stuff with odata, idata, and sdata
}
foo<int><<<blocks, threads, mem>>>(d_odata, d_idata);
foo<float><<<blocks, threads, mem>>>(d_odata, d_idata);
编译之后会遇到下述错误：“declaration is incompatible with previous "sdata" (declared at line 3)extern __declspec(__shared__) T sdata[];”。原因就是在编译的时候，cuda会为上述两次调用生成相同的代码，因而会出现变量重复定义的问题。
解决方法如下：将模板类实例化。首先新建一个头文件，SharedMem.h，内容如下：
#include <cutil_inline.h>
template <class T>
class SharedMem
{
public:
T* getPointer() { return NULL; };
};// specialization for int
template <>
class SharedMem <int>
{
public:
__device__ int* getPointer() { extern __shared__ int s_int[]; return s_int; }
};// specialization for float
template <>
class SharedMem <float>
{
public:
__device__ float* getPointer() { extern __shared__ float s_float[]; return s_float; }
};
上述代码实际上就是将定义共享内存的代码单独拿出来，然后放在类中实现。上述代码需要注意以下几个方面：
1. 因为在定义共享内存时用到关键字__shared__，所以我们要将函数定义成cuda函数。在函数前面需要加相应关键字，但不能是__global__，因为它要求返回void类型，所以只能是__device__；
2. 包含cuda程序相应的头文件，否则编译不通过；
完成上述头文件的编写，在具体调用过程中代码如下：
template<class T>
__global__ void foo( T* g_idata, T* g_odata)
{
// shared memory the size is determined by the host application SharedMem<T> shared;
T* sdata = shared.getPointer(); // .. the rest of the code remains unchanged!
}
这样我们就可以在cuda中正常使用模板了。
参考网页：
1. cuda中应用模板函数
