如何使用CUDA传递Tree结构

Question

我创建了一个简单的C树,我打算在GPU上移植它.

我树的结构如下:

typedef struct node{

    short int final; // 2 byte
    char number; // 1 byte 
    struct node *child[2]; // 8 * NUM_SIZE byte

}node;

我现在想知道如何在CUDA上移植代码,或者如何在CUDA中使用结构,我应该在设备上创建树吗？或者在主机上创建树并将其传递给设备？

在设备上创建树似乎是我的最佳答案,虽然我不确定如何在设备上使用结构.

谢谢

Answer 1

您可以使用三种基本方法来解决此类问题:

1. 使用主机API为树分配全局内存,并在主机上构建树,然后将该树复制到设备.
2. 使用主机API为树分配全局内存,并在设备上构建树
3. 为设备上的树分配运行时堆内存,并在设备上构建树

哪个最合适取决于您的使用案例以及您使用的GPU.每个都有缺点 - 例如(1)需要在主机存储器中保存设备树的副本,(2)可能需要在设备代码中很难实现的同步,以及(3)受到设备性能有限的影响在支持它的那些设备上运行时堆内存分配,并将生成一个无法使用主机API直接访问的树

最后要说的是,基于指针的树通常不具备高性能或非常适合GPU计算,您可能希望在使用诸如您在其中描述的结构之类的结构之前考虑替代数据结构和算法.题.

编辑:

显然所有这些选项对你来说都太难了,所以这里有一个绝对琐碎的例子,你可以看看它是如何完成的.首先是代码:

#include <iostream>

struct __align__(8) node{
    char val; 
    struct node *child;
};

__global__
void kernel(node * tree, char *out, int n)
{
    node *p = tree;
    int i=0;
    while(p->val != 0) {
    out[i++] = p->val;
    p = p->child;
    }
}

int main(void)
{
    const int n = 15;
    char data[n] = "tietamattomana";
    node tree[n]; 

    node * tree_d;
    char * output_d;
    cudaMalloc((void **)&tree_d, n * sizeof(node));
    cudaMalloc((void **)&output_d, n * sizeof(char));

    node * p = tree_d;
    for(int i=0; i<n; i++) {
        tree[i].val = data[i];
        tree[i].child = (++p);
    }

    cudaMemcpy(tree_d, tree, n * sizeof(node), cudaMemcpyHostToDevice);
    kernel<<<1,1>>>(tree_d, output_d, n);

    char output[n];
    cudaMemcpy(output, output_d, n * sizeof(char), cudaMemcpyDeviceToHost);
    for(int i=0; i<n; i++) {
        std::cout << output[i];
    }
    std::cout << std::endl;

    return 0;
}

在这里,我刚刚在主机上填充了一个简单的链表,并将其复制到设备中,每个列表节点都从主机字符串中保存一个值.单个GPU线程从头到尾遍历列表,读出每个节点的值并将其存储在输出数组中.为确认一切正常,主机将输出数组从GPU复制回来并回显输出数组的内容,即:

$ nvcc -arch=sm_30 -Xptxas="-v" tree.cu
ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z6kernelP4nodePci' for 'sm_30'
ptxas info    : Function properties for _Z6kernelP4nodePci
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 10 registers, 340 bytes cmem[0]

$ ./a.out 
tietamattomana

也许这至少可以帮助你开始实现你想要实现的目标,并且可以阐明我提到的其他可能性如何在实践中发挥作用以及我提到的缺点将发挥作用.