2

我创建了一个简单的 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 中使用结构，我应该在设备上创建树吗？还是在主机上创建树并将其传递给设备？

在设备上创建树似乎对我来说是最好的答案，尽管我不确定如何在设备上使用结构。

谢谢

4

1 回答 1

7

对于此类问题，您可以使用三种基本方法：

  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

也许这至少会让你开始走上你想要实现的目标，并且可能会阐明我提到的其他可能性在实践中如何发挥作用，以及我提到的缺点将在哪里发挥作用。

于 2013-08-26T12:19:40.853 回答