2011-10-03 55 views
4

我有一个调用在一个.CU文件中定义这样CUDA和模板:需要专业化声明?

template<typename T, class M> 
__global__ void compute_kernel(T* input, T* output, n) { 
    M m; 
    // compute stuff using m 
}; 

template<typename T, class M> 
void compute(T* input, T* output, int n) { 
    // ... compute blocks, threads, etc. 
    compute_kernel<T,M> <<<dim_grid, dim_block>>>(input, output, n); 
    // ... 
}; 

和一个头文件是内核(__global__)将被包括在主机代码模板化包装函数仅具有声明

template<typename T, class M> 
void compute(T* input, T* output, int n); 

但是,从具有任意模板参数的主机调用compute()时,编译失败,并且undefined reference to 'void reduce(...)'并且仅当我在代码编译时向.cu文件的末尾添加专门化声明:

template void 
compute<int, Method1<int> >(int* input, int* output, int n); 

template void 
compute<float, Method1<float> >(float* input, float* output, int n); 

template void 
compute<int, Method2<int> >(int* input, int* output, int n); 

template void 
compute<float, Method2<float> >(float* input, float* output, int n); 

那么,是否有必要专门化每个模板化的函数以使其可以从主机调用? (这是一个很大的缺点)

感谢您的意见!

+2

与任何模板一样,您必须确保整个模板定义在每个*实例化站点都可见(除非您故意提供明确的实例化)。 –

+0

@KerrekSB:我认为你是对的,这是问题所在。但是,要分离主机和设备代码(以及编译),似乎无法将整个模板定义无处不在... – bbtrb

+0

我不确定问题出在哪里。 CUDA编译器不会自动分离这两个代码路径吗?所以,只需使用函数模板的模板定义创建一个普通的头文件,并将其包含到任何地方 - 阻塞是什么? –

回答

1

这是一个C++ FAQ,不限于CUDA。

如果您在.cpp或.cu文件中有模板实现,那么当您编译该翻译单元时,编译器不可能知道您将需要的模板参数的排列方式。因此,当你链接你会得到错误。

您可以将实现放在头文件中(在这种情况下,您需要在.cu文件中实例化,因为它包含CUDA),否则您将不得不显式实例化所有需要的排列。如果你必须做很多这些,那么你可以使用宏来实例化所有的排列。

+0

谢谢,这回答它。我从来没有真正想过模板的非内联实现,并始终将它们保存在标题中。所以,只要我只从* .cu文件中调用模板函数,我就可以在任何地方包含完整的定义,并且不需要显式实例化。这已经非常有帮助,因为(最终)只有我的单元测试会从宿主代码中调用这些函数。 – bbtrb

+0

@Tom在CUDA中使用表达式模板怎么样?我已在此链接发布了一个问题(http://stackoverflow.com/questions/15628470/unresolved-externals-in-cuda-expression-template-library-under-visual-studio-201),但我没有收到回答。对于表达式模板,似乎不可能实例化所有可能的排列。你能为此提出一个解决方案吗? – JackOLantern