实例化cuda时模板过度递归

Question

我有一个模板函数，其中模板参数是整数。该整数用于创建不同的内核。以前，所有可能的模板都是在表格中手动实例化的（可以工作，但很难看），但我尝试使用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个case的switch就太蠢了。 有谁知道增加 nvcc 模板递归限制或其他自动方式来创建我的列表？

编辑：我发现 gcc 选项 ftemplate-depth 可以更改实例化递归限制，但我还没有找到等效的 nvcc 选项。

Answer 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 开发人员来到这里，向 nvcc 添加“ftemplate-深度”选项应该是一个好主意，不是吗？

Answer 2

至少在 CUDA 12 中（可能已经在早期版本中），nvcc 编译器具有

-ftemplate-depth

命令行开关。

参见：https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#ftemplate-深度-limit-ftemplate-深度