实例化 cuda 时模板过度递归答案

【问题标题】：Template excessive recursion at instantiation cuda实例化 cuda 时模板过度递归
【发布时间】：2014-11-01 09:59:34
【问题描述】：

我有一个模板函数，其中模板参数是一个整数。该整数用于创建不同的内核。以前在表格中手动实例化的所有可能模板（有效但丑陋）但我尝试使用提出的解决方案here。因为我有超过 800 个可能的内核，所以模板递归方法要优雅得多。我已经在我的代码的 C++ 版本上测试了模板递归，它运行良好，但 nvcc 似乎限制了我的实例化的递归。

这是我之前的丑陋模板实例化列表的简化示例，它可以正常工作（即使是 800 个内核实例化）：

// the template kernel 
template <int i> __global__ void kernel(int some_data)
{
    switch(i)
    {
    case 0:
        // do something
        break;
    case 1:
        // do some other things
        break;
    //...
    case 799:
        // do some other things
        break;
    }
}

typedef void (*kernel_pointer) (int some_data)

// the ugly huge list
kernel_pointer kernel_list[800] = {
    &kernel <0>,
    &kernel <1>,
    //...
    &kernel <799> }

int main()
{
    int kernel_index = 10;

    //the call
    kernel_pointer my_kernel = kernel_list[kernel_index];
    my_kernel<<<<1,1>>>>(the_data);        
}

这里是 nvcc 不喜欢的漂亮的模板递归。它替换了上一段代码中的列表：

#define N_KERNELS 800
template< int i> bool dispatch_init( kernel_pointer* pTable )
{
    pTable[i] = &kernel<i>;    
    return dispatch_init<i-1>( pTable );
}    
// edge case of recursion
template<> bool dispatch_init<-1>(kernel_pointer* pTable) { return true; }

// call the recursive function
const bool initialized = dispatch_init<-1>( kernel_list );

实际上，我没有一个模板参数，而是有 6 个模板参数组合起来创建了所有数百个组合。否则，一个有 800 个箱子的开关将是非常愚蠢的。有没有人有增加 nvcc 模板递归限制或其他自动方式来创建我的列表的想法？

编辑：我找到了改变实例化递归限制的 gcc 选项 ftemplate-depth 但我没有找到等效的 nvcc 选项。

【问题讨论】：

nvcc 模板递归深度限制似乎是 199。你能把你的表分成几块吗？
我也试过了，但我无法让它工作。我尝试的是创建一个递归迭代 100 次的模板函数，它调用另一个模板函数（此处为 dispatch_init）作为模板参数给出一个索引为 0 的索引，该索引由另一个模板递归迭代到 8。但我需要另一个模板参数，一步跳转之前的实例化。我的问题是为第二个模板创建递归的边缘情况，因为不允许部分模板专业化。类模板的部分特化是可能的，但我还没有测试过。
Here's 以 4 个块初始化的 800 表的示例。我承认通过 6 个自变量进行递归很困难，但您的问题似乎与此无关。
谢谢。我根据您的建议发布了答案。

标签： templates recursion cuda nvcc

【解决方案1】：

根据 Robert Crovella 提出的想法，包括将表格分成几部分来构建，这里有一个如何修复“错误”的示例：

#define N_KERNELS 850
// template kernel 
template <int i> __global__ void kernel(int a)
{
    switch(i)
    {
    case 0:
        printf("%d\n", a*i);
        break;
    case 1:
        printf("%d\n", a*i);
        break;
    //...
    case 849:
        printf("%d\n", a*i);
        break;
    }
}

typedef void (*kernel_pointer) (int);

kernel_pointer kernel_list[N_KERNELS];

// Function that instantiates all the needed kernels using recursion.
template< int i> bool dispatch_init( kernel_pointer* pTable )
{
    pTable[i] = &kernel<i>;    
    return dispatch_init<i+1>( pTable );
}    

// Edge cases of recursion made with a template specialization
template<> bool dispatch_init<199>(kernel_pointer* pTable)
{
    pTable[199] = &kernel<199>;
    return true;
}
template<> bool dispatch_init<399>(kernel_pointer* pTable) 
{
    pTable[399] = &kernel<399>;
    return true;
}
template<> bool dispatch_init<599>(kernel_pointer* pTable)
{
    pTable[599] = &kernel<599>;
    return true;
}
template<> bool dispatch_init<799>(kernel_pointer* pTable)
{
    pTable[799] = &kernel<799>;
    return true;
}
template<> bool dispatch_init<N_KERNELS>(kernel_pointer* pTable) { return true; }

// Call the recursive function few times to instantiate all the kernels without reaching the recursive instantiation limit
const bool initialized = dispatch_init<0  >( kernel_list );
const bool initialized = dispatch_init<200>( kernel_list );
const bool initialized = dispatch_init<400>( kernel_list );
const bool initialized = dispatch_init<600>( kernel_list );
const bool initialized = dispatch_init<800>( kernel_list );


int main()
{
    int kernel_index = 10;

    kernel_pointer my_kernel = kernel_list[kernel_index];
    my_kernel<<<<1,1>>>>(6);        
}

我不喜欢这个修复，但它暂时可以解决问题。如果有一天 Nvidia 开发人员来到这里，最好将选项“ftemplate-depth”添加到 nvcc，不是吗？

【讨论】：