在 CUDA 上生成决策树答案

【问题标题】：Generating decision tree on CUDA在 CUDA 上生成决策树
【发布时间】：2013-02-08 23:21:55
【问题描述】：

我想在CUDA上生成一些决策树，下面我们有伪代码（代码很原始，只是为了理解我写的内容）：

class Node
{
public :
    Node* father;
    Node** sons;
    int countSons;

    __device__ __host__ Node(Node* father)
    {
        this->father = father;
        sons = NULL;
    }
};

__global__ void GenerateSons(Node** fathers, int countFathers*, Node** sons, int* countSons)
{
    int Thread_Index = (blockDim.x * blockIdx.x) + threadIdx.x;

    if(Thread_Index < *(countFathers))
    {
        Node* Thread_Father = fathers[Thread_Index];
        Node** Thread_Sons;
        int Thread_countSons;
        //Now we are creating new sons for our Thread_Father
        /*
        * Generating Thread_Sons for Thread_Father;
        */
        Thread_Father->sons = Thread_Sons; 
        Thread_Father->countSons = Thread_countSons;

        //Wait for others
            /*I added here __syncthreads because I want to count all generated sons
             by threads
            */
            *(countSons) += Thread_countSons;
        __syncthreads();

        //Get all generated sons from whole Block and copy to sons

        if(threadIdx.x == 0)
        {
            sons = new Node*[*(countSons)];
        }
        /*I added here __syncthreads because I want to allocated array for sons
            */
        __syncthreads();

        int Thread_Offset;
        /*
        * Get correct offset for actual thread
        */
        for(int i = 0; i < Thread_countSons; i++)
            sons[Thread_Offset + i] = Thread_Sons[i];
    }
}

void main ()
{
    Node* root = new Node();
    //transfer root to kernel by cudaMalloc and cudaMemcpy
    Node* root_d = root->transfer();

    Node** fathers_d;
    /*
    * preapre array with father root and copy him to kernel
    */

    int* countFathers, countSons;
    /*
    * preapre pointer of int for kernel and for countFathers set value 1
    */

    for(int i = 0; i < LevelTree; i++)
    {
        Node** sons = NULL;
        int threadsPerBlock = 256; 
        int blocksPerGrid = (*(countFathers)/*get count of fathers*/  + threadsPerBlock - 1) / threadsPerBlock;
        GenerateSons<<<blocksPerGrid , threadsPerBlock >>>(fathers_d, countFathers, sons, countSons);
        //Wait for end of kernel call
        cudaDeviceSynchronize();

        //replace
        fathers_d = sons;
        countFathers = countSons;
    }
}

所以，它适用于 5 级（为检查器生成决策树），但在 6 级我有错误。在内核代码的某个地方，malloc 正在返回 NULL ，对我来说，这是一个信息，即 blockThreads 中的某些线程无法分配更多内存。我很确定我正在清理所有我不需要的对象，在调用内核的每一端。我想，我无法理解 CUDA 中使用内存的一些事实。如果我在线程的本地内存中创建对象并且内核结束了他的活动，那么在内核的第二次启动时，我可以看到内核第一次调用的节点是。所以我的问题是来自第一次内核调用的对象Node 存储在哪里？它们是否存储在块中线程的本地内存中？那么如果是真的，那么在每次调用我的内核函数时，我都会减少这个线程的本地内存空间？

我正在使用带有计算能力 2.1 的 GT 555m、CUDA SDK 5.0、带有 NSight 3.0 的 Visual Studio 2010 Premium

【问题讨论】：

您在内核中调用 new 而从不调用 delete。由于您使用的是____global____ void GenerateSons，我敢打赌您的设备内存不足。
好的，我的设备有 2Gb 的空间，sizeof(Node) = 28。第一次调用生成 7 个儿子，第二个 49，接下来是 379，最后一次正确调用 2769。所以我的设备生成了 3204 个儿子，提供 87Kb ???
嗯，我想知道 new 是否正在从共享内存中提取内存。我得查一下文档。
我认为“线程的寄存器”==“线程的本地内存”已填写，但我不确定。目前，我脑海中唯一的解决方案是将新儿子从内核复制到全局内存......
someValue 和 someValue2 的值是多少。您希望通过在该内核中调用两次__syncthreads() 来实现什么？

标签： c++ memory cuda binary-tree

【解决方案1】：

好吧，

我发现，内核中的new 和mallocinvoke 分配在设备上的全局内存中。我也发现了这个

默认情况下，CUDA 创建一个 8MB 的堆。

CUDA Application Design and Development, page 128

所以，我使用这种方法cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024); 将设备上的堆内存增加到 128Mb，并且程序正确生成了 6 级树（22110 个儿子），但实际上我得到了一些内存泄漏......我需要找到。

【讨论】：