Ответ, данный Робертом Кровеллой, является правильным в 5-м пункте, ошибка заключалась в использовании init в каждом рекурсивном вызове, но я хочу уточнить кое-что, что может быть полезно для других новичков с CUDA.
Я использовал эту переменную, потому что когда я пытался запустить дочернее ядро, передавая локальную переменную, я всегда получал исключение: Error: a pointer to local memory cannot be passed to a launch as an argument.
Поскольку я C# эксперт-разработчик, я не использовал к использованию указателей (Ref выполняет для этого низкоуровневую работу), поэтому я подумал, что в программировании на CUDA / c нет способа.
Как показывает Роберт в своем коде, это возможно копирование указателя с помощью memalloc для использования его в качестве ссылочного аргумента.
Вот ядро, упрощенное как пример глубокой рекурсии.
__device__ char init[6] = { "12345" };
__global__ void Recursive(int depth, const char* route) {
// up to depth 6
if (depth == 5) return;
//declaration for a referable argument (point 6)
char* newroute = (char*)malloc(6);
memcpy(newroute, route, 5);
int o = 0;
int newlen = 0;
for (int i = 0; i < (6 - depth); ++i)
{
if (i != threadIdx.x)
{
newroute[i - o] = route[i];
newlen++;
}
else
{
o = 1;
}
}
printf("%s\n", newroute);
Recursive <<<1, newlen>>>(depth + 1, newroute);
}
__global__ void RecursiveCount() {
Recursive <<<1, 5>>>(0, init);
}
Я не добавляю основной вызов потому что я использую ManagedCUDA для C#, но, как говорит Роберт, можно выяснить, каков вызов RecursiveCount.
Об окончании массивов char с помощью /0 ... извините, но я не не знаю точно, в чем выгода; этот код работает без них.