cuda内核for循环太长？

Question

经过长时间的手动调优和半分天真的调试，我终于找到了内核崩溃的原因。你不能在内核中做一个太长的for循环吗？ 这是最小完整和可验证的代码： 错误是：未指定的启动失败

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>
using namespace std;
typedef unsigned char uchar;

struct NodePointer {
    int id;
    uchar dist;
    int idintree;
    NodePointer() :id(0), dist(0), idintree(-1){}
};
struct TreeNode {
    NodePointer father;
    NodePointer children[4];
    int id;
    int childrenNum;
    int idintree;
    TreeNode() :id(0), childrenNum(0), idintree(-1){}
};

__global__ void
kernel4(int NumTN, TreeNode* tempthistree, int size)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    if (index < NumTN)
    {
        for (int i = 0; i < size; i++)
        {
            TreeNode node1 = tempthistree[i];
            printf("              node %d in tree %d, its id in tree is %d, its child num is %d\n", i, index, node1.idintree, node1.childrenNum);

        }
    }
}


int main()
{
    int n1 = 33417;
    TreeNode * testtree;
    cudaMallocManaged(&testtree, n1*sizeof(TreeNode));
    for (int i = 0; i < n1; i++)
    {
        TreeNode c;
        c.idintree = i;
        c.id = i;
        c.father.id = i - 1;
        c.father.dist = 1;
        c.childrenNum = i % 4;
        int aaa = i % 4;
        for (int j = 0; j < 4; j++)
        {
            c.children[j].dist = j;
            c.children[j].id = 1;
            c.children[j].idintree = 10;
        }
        testtree[i] = c;
    }


    kernel4 << <1, 1 >> >(4000, testtree, n1);
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess)
    {
        printf("Kernel3 error :%s\n", cudaGetErrorString(err));
        system("pause");
    }
}

Answer 1

假设您使用的是 Windows，您可能会遇到超时检测和恢复问题，这意味着您的内核执行时间过长。

看到这个答案： C++: Simple CUDA volume reconstruction code crashing

Answer 2

首先，内核调用：kernel4 << <1, 1 >> >(4000, testtree, n1)

您正在使用只有 1 个线程的 1 个块。因此，您的索引在这里仅变为 0。您正在给这个线程以下任务：CPU，请执行一个 for 循环。所以那里没有并行性。

您可以在内核中使用 for 循环，但我建议您将其放在外面。至少这是一种可能性：

__global__ void
kernel4(int NumTN, TreeNode* tempthistree, int size) 
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    for (int i=0;i<NumTN;i++)
    {
        if (index < size)
        {
            TreeNode node1 = tempthistree[index];
            printf("node %d in tree %d, its id in tree is %d, its child num is %d\n", index, i, node1.idintree, node1.childrenNum);

        }
    }
}

现在您需要更多线程，最少 n1。一个块中有 1024 个线程（现在）的限制，因此您需要更多块才能达到 n1。例如 ceil(n1/1024) 块。

此外，我不确定在 C 中，但在结构中使用数组时要小心。尝试一个有关内存传输的简短示例，看看您是否得到正确答案（我在使用 fortran 时遇到了一些问题）。

PS：还要检查你的计算能力：printf 仅在计算能力 >= 2.0 的 GPU 上可用。