Шаблон чрезмерной рекурсии при создании экземпляра 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, который сочетает в себе, чтобы создать все сотни комбинаций. Else, коммутатор с 800 корпусами был бы действительно глупым. У кого-нибудь есть идея увеличить предел рекурсии nvcc или другой автоматический способ создания моего списка?

Редактировать: Я нашел параметр gcc option ftemplate-depth, который изменяет предел рекурсии, но я не нашел эквивалентную опцию nvcc.

Answer 1

От идеи, предложенной Робертом 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, нет?