如何在 GPU（最好是 CUDA）上对两个数据容器执行关系连接？

Question

我想做什么：

在 GPU 上，我试图模仿 SQL 在关系代数中使用的约定来对 table 执行连接（例如内连接、外连接、交叉连接）。在下面的代码中，我想执行一个内部连接。想象两个 tables（容器），其中一个 table 是 Parent/Master table，另一个是子 table。 Parent to Child 连接关系是一对多（或 1 到 none，如果 Child_ParentIDs 中没有元素与 Parent_IDs 中的元素匹配）。

示例输入数据：

Parent_IDs:    [1, 2,  3,  4, 5]  ... 5 elements
Parent_Values: [0, 21, 73, 0, 91] ... 5 elements
Child_ParentIDs:   [1,   1,   1,  2,   3,   5,  5]  ... 7 elements
Child_Permanences: [120, 477, 42, 106, 143, 53, 83] ... 7 elements
Child_Values:      [0,   0,   0,  0,   0,   0,  0]  ... 7 elements

作为 SQL 查询的操作：

SELECT child.permanence * parent.value FROM child, parent WHERE child.parent_id = parent.id;

操作说明：

将Child_ParentIDs连接到Parent_IDs以访问对应的Parent_Values。用对应的Parent_Values与对应的Child_Permanences相乘，将每次运算的结果放入Child_Values.

预期输出（Child_Values 是操作期间唯一更改的向量）：

Child_ParentIDs:   [1,   1,   1,  2,    3,     5,    5]     ... 7 elements
Child_Permanences: [120, 477, 42, 106,  143,   53,   83]    ... 7 elements
Child_Values:      [0,   0,   0,  2226, 10439, 4823, 7553]  ... 7 elements

解释（以防万一）：

2226 的值是通过将 106 和 21 相乘得出的。10439 是通过将 143 和 73 相乘得出的。还要注意，所有条目都保留在子向量上（所有 7 个元素仍然存在于输出中，尽管 Child_Values 个别元素已更新）。父向量未保留在输出中（注意向量列表中缺少 ParentID 4，并且那里没有 "dummy" 占位符）。这是 "Inner Join".

的行为

我还没有开始工作的优雅解决方案的想法：

-利用 CUDA 的动态并行性。也许我在整个互联网上发现的唯一解决方案完全符合我的要求，那就是 here-part 1 and here-part 2.

-使用CUDPP的哈希运算；

-Alenka 数据库。

最后，重申我的问题：

从纯 GPU 的角度（最好是 CUDA，但 OpenCL 也可以）是否有任何可行的解决方案来完成两个独立数据容器的关系连接，以便可以通过上述方式搜索数据并并行更新元素加入？

编辑
Parent_IDs 并不总是一个序列。在运行时间内，可以删除父向量中的元素。新插入的父元素将始终附加一个 ID，该 ID 从最后一个元素的 ID 开始。话虽如此，我明白这意味着子元素可以被孤立，但我不会在这里解决这个问题的解决方案。

Answer 1

它看起来像是 Child_Permanences 的元素与 Parent_Values 的选定元素之间的简单元素乘法。通过一些限制，这可以用一个 thrust::transform.

来完成

thrust::transform(
    Child_Permanences.begin(),
    Child_Permanences.end(),
    thrust::make_permutation_iterator(
        Parent_Values.begin(),
        thrust::make_transform_iterator(Child_ParentIDs.begin(),
                                        _1 - 1)),
    Child_Values.begin(),
    _1 * _2);

您可能会注意到未使用 Parent_IDs。就是上面代码的限制。该代码假定 Parent_IDs 只能是一个 1 碱基序列。如果 Parent_IDs 是一个 0 碱基序列，或者 Child_ParentIDs 只是一个父值索引，您会发现 thrust::make_transform_iterator 不是必需的，如下所示。

Child_ParentIDs:   [0, 0, 0, 1, 2, 4, 4]

编辑

以上代码假设1）没有孤儿；和 2) Parent_IDs 是一个固定的基于 1 的序列，如 1, 2, 3, ...

前提是

没有孤儿；
Parent_IDs 是无序且唯一的；
Child_ParentIDs 未排序但不是唯一的；

并且您的 Parent_IDs 属于 int16 类型这一事实，您可以创建父值索引 table 供子元素查找，当 Parent_IDs 相当小。

假设Parent_IDs的取值范围是[1, 32767]，求解代码可以是

thrust::device_vector<int> Parent_index(32768, -1);
thrust::scatter(thrust::make_counting_iterator(0),
                thrust::make_counting_iterator(0) + Parent_IDs.size(),
                Parent_IDs.begin(),
                Parent_index.begin());
thrust::transform(
    Child_Permanences.begin(),
    Child_Permanences.end(),
    thrust::make_permutation_iterator(
        Parent_Values.begin(),
        thrust::make_permutation_iterator(
            Parent_index.begin(),
            Child_ParentIDs.begin())),
    Child_Values.begin(), _1 * _2);

请注意，每次修改父向量时都需要重新创建Parent_index。

如何在 GPU（最好是 CUDA）上对两个数据容器执行关系连接？

How to perform relational join on two data containers on GPU (preferably CUDA)?

c++

cuda

gpgpu

thrust

dynamic-parallelism