CN113204559A - 一种gpu上的多维kd树优化方法 - Google Patents

Abstract

本发明提供一种GPU上的多维KD树优化方法，用于在GPU中加速KD索引的构建与查询过程；本发明从待划分数据集的全体出发，将传统KD树上的二等分划分操作视作一种可分配资源，并将这种资源称之为额度，从如何有效分配额度资源的角度，提出了一种GPU上优化KD树构建与查询过程的新方法和新技术；基于额度分配策略，贪心划分策略，通过提高每次排序的利用率，获得了更大的划分效率，加速了KD树构建过程，降低了KD树的深度；利用学习索引的思想优化了中间节点的参数访问方式，有效降低了批量查询过程中线程束的分化状况，优化了批量查询的性能。

Description

一种GPU上的多维KD树优化方法

技术领域

本发明属于数据库索引技术领域，具体涉及一种GPU上的多维KD树优化方法。

背景技术

多维数据划分是多维数据处理中的一个经典问题，其现实意义是将众多的多维数据划分到多个具有不同空间特征的子空间中，从而在海量数据访问中获得一定的数据过滤优势，加快数据访问的速度。K-dimensional tree(简称KD树)是一种经典的分割k维数据空间的数据结构，被广泛应用于多维，高维数据的数据划分过程中，用以加速相关应用的临近搜索和范围查询。KD树是一种轴对称的二叉空间划分树，具有场景自适应划分，低存储消耗和快速遍历等优势。因此，探究KD树的优化策略对解决各种实际问题有着重要意义。

目前，KD树在GPU上的并行构建算法往往都遵循广度优先策略。在进行每层的节点划分前，都要计算出当前层中需要进行划分的节点数量，并分配对应数量的线程和相关资源。对于KD树而言，下一层的线程数量一般是上一层线程数量的两倍。设置叶子节点可保存的最大纪录数量为LM时，KD树的理论深度是

其中，n是KD树需要划分数据集大小。

在构建后，GPU上KD树常用的批量数据查询方法有三种。三种方法在第一步都是一样的，需要为每一个查询请求分配一个线程，随后三种方法就各有不同：第一种方法中，每个查询会沿着KD树的索引结构对中间节点一层一层地进行广度优先查询，直到抵达叶子节点，并读取对应的数据；第二种方法中，不沿着KD树进行访问，而是直接对所有叶子节点进行暴力访问，直到处理所有的查询；第三种是对KD树划分出的叶子节点进行重组织，并构建一种新的过滤器，然后依据这种新的过滤路径进行叶子节点的访问与数据读取。

这三种方法中，第一种方法的每个查询至少需要执行

次的寻路操作，第二种完全舍弃了KD树带来的数据过滤优势，但可直接访问叶子节点；第三种则是第一种方法与第二种方法的混合使用，其依据叶子节点构建新的并行索引，从而使得无需执行

次寻路操作也能保持一定的过滤性，但需要保存一个额外的辅助索引。这三种优化策略的分歧在于寻路次数以及查询过程中的过滤性的权衡；由于KD树的深度与数据集大小相关，处理海量数据时，

次寻路的代价是可观的，尤其是对于GPU的单指令多线程模型。

可以认为，KD树的深度对于KD树的性能有着直接影响；而现有的KD划分方法不能降低KD索引深度的原因在于，KD树是基于二等分划分策略的划分树，也就是二叉空间划分树，这使得其深度必然与数据集大小关联，为

对于此，KDB树对KD树的扇出度进行了优化，其结合了KD树与B树的优势，通过增加扇出度，降低了KD树的深度；目前在GPU上的KDB树并行构建方法，主要通过增加参与每次划分的维度数量，提高扇出度。

另外，近些年来，一种被称为学习索引的新技术被提了出来。学习索引将查询数据的输入键值以及查询结果的实际地址视作一个训练对象(K，P)，其中，K是输入的键值，P是输入键值的查询结果的存储地址；将索引视作一个学习模型，认为可以在一些常见的机器学习模型中，比如两层的全连接神经网络，或者简单的多元线性函数，使用(K，P)进行有监督的模型训练；如果训练出的模型性能良好，那么可以在输入查询键值K后，直接输出查询键值对应的查询结果的地址P，从而获得线性的查询访问时间。目前有关KD树的学习索引方法，与前面所介绍的KD树上的第三种查询方法类似，是在构建好KD索引后，使用训练好的模型替换传统的KD索引进行查询处理，另外，这些方法并不面向GPU。

发明内容

为了提高GPU上KD树的构建速度与查询性能，缓解KD树在GPU上进行批量查询中的线程束分化状况，本发明提出了额度的概念，结合额度划分策略，贪心划分策略，以及学习索引的索引优化理论，在保证树结构在理论上平衡的同时，优化了GPU上KD树的构建速度与查询性能，包括：

步骤1：对于k个维度的数据集A，获得数据集A包含的记录数量n，并通过

计算出数据集A的额度Q，其中LM表示叶子节点可以保存的记录数量上限，α表示可调参数，α∈[0.5,1.0]，

表示向下取整；

步骤2：对k个维度进行排序，并将额度Q分配给k个维度；

步骤3：在GPU上进行k层的并行循环划分，构建出KD树的主体结构；

步骤4：选取第k/2个维度作为划分维度，对待划分的节点列表List进行基于贪心策略的划分；

步骤5：对步骤3中产生的所有中间节点的参数访问进行优化。

所述步骤2包括：

步骤2.1：对数据集A的k个维度的数据进行采样，以k个维度上采样数据的均方差作为比较标准，对k个维度进行降序排序，保证均方差较大的维度占据较前的位置，排序结果记为S；

步骤2.2：将额度Q分配给k个维度，每个维度分配得到的额度不小于

并保证排序较前的维度所分配的额度必然不小于排序较后的维度所分配的额度，第i个维度所分配得到的额度保存为q_i，其中i∈[0，k-1]；

步骤2.3：对于每一个维度d_i，根据其获得的额度q_i，通过

计算出维度d_i需要划分的子段数量seg_i，维度d_i需要获得(seg_i-1)个等分点，将维度d_i的排序结果S及维度d_i对应的seg_i保存为一个结构体Pre。

所述步骤3包括：

步骤3.1：将数据集A的各个维度的数据以及对应的LM、Pre传递到GPU中，其中数据集A的各维度数据按照列存储模式进行存储，使用CUDA内置的内存申请函数cudaMallocPitch()进行内存申请，将返回的各列的间距记作pitch，将申请的全局内存空间记为M，LM和Pre保存在常量内存中，另外，申请一块位于全局内存的空间B作为划分过程中的中间数据存储介质，且满足内存空间不小于M；

步骤3.2：申请一个GPU上的全局内存空间nodeSpace[0]，其大小为一个中间节点的大小，保存有一个初始节点a；将a的地址压入待划分节点链表List，List中每个节点保存有起始地址address、地址偏移量offset两个参数，可以通过基址变址寻址的方法，使用起始地址address加上偏移量(offset+pitch*u)获取节点第u个维度数据的起始地址；设置节点a的起始地址为M，地址偏移量为0；此外，每个待划分节点还保存有当前节点的记录数量以及该节点k个维度上的范围约束，初始节点a在每个维度上的范围约束是(-∞，+∞)，对于每个非初始节点的中间节点，初始化时范围约束继承自父节点的范围约束，并在父节点划分结束之后依据划分参数改变每个子节点在划分维度上的范围约束；由于此时List中只有a一个待划分节点，设置List中的待划分节点数量n＝1；

步骤3.3：根据Pre对List进行k层循环划分，每层划分中参与划分的维度按照Pre中的维度顺序S进行依次选择，每层循环中所有待划分节点都在同一维度上进行数据划分；在第i层的并行划分中，参与划分的维度编号设置为i，i∈[0，k-1]；

步骤3.4：如果当前划分次数小于k，跳转步骤3.3.1进行下一层的并行划分，否则跳转步骤4。

所述步骤3.3中每层循环具体表述为：

步骤3.3.1：获得参与划分的维度编号i，以及对应的seg_i；根据List中的待划分节点数量n，申请大小为n*seg_i*sizeof(midNode)的全局内存nodeSpace[i+1]，用于存储划分后的n*seg_i个子节点，其中sizeof(·)表示获取某个结构体所占用内存大小，midNode表示节点的结构体，sizeof(midNode)表示获取一个节点所占用的内存大小；每个中间节点都有seg_i个节点指针，用于指向其对应的seg_i个子节点；因此，在申请全局内存nodeSpace[i+1]后，需要建立第j个待划分节点与nodeSpace[i+1]中的第(j*seg_i+s)个子节点的关联关系，其中j∈[0,n-1]，s∈[0,seg_i-1]，这个过程申请n个线程进行并行处理，每个线程处理一个待划分节点，并在构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[i+1]一直存活到KD索引的销毁；

步骤3.3.2：依次对n个待划分节点的第i维数据进行并行排序，并选取(seg_i-1)个的等分点，具体为：将每个节点的第i维数据复制到一个可用的全局空间tmp上，使用排序函数Thrust::sort()进行排序操作，并将排序后的数据等分为seg_i个子段，将每个子段在维度i上的范围、对应的记录数量保存到对应子节点的内部参数中，其中，保存子段在维度i上的范围到对应的子节点中，意味着更新子节点在第i维的范围约束；此外，每个子段的范围是半开半闭区间，具体为左开右闭；

步骤3.3.3：启动n个线程，每个线程处理seg_i个归属于同一个待划分节点的子节点；每个线程中，如果i％2＝＝0，初始化每个子节点的起始地址为B，否则，将起始地址设置为M，％表示取余；地址偏移量offset为父节点的地址偏移量；随后将子节点f的地址偏移量offset[f]更新为

f∈[0，seg_i-1],其中，subNode[t].record_num代表第t个子节点的记录数量；更新完待划分节点的所有子节点的offset后，将所有子节点的record_num重置为0；

步骤3.3.4：启动n个线程进行数据的并行划分；每个线程处理一个待划分节点，并执行一个循环，循环次数为待划分节点中的记录数量record_num；第p次循环中，根据设定的判断方法判断第p个记录所属的子节点，并将数据插入到对应的子节点的数据空间中；所述设定的判断方法为：进行一次循环比较，循环次数是待划分节点的子节点数量，在第u次循环中，比较记录f在第i维上的数据是否满足子节点u在第i维上的范围约束，如果满足，则跳出循环，并将第f个记录插入到子节点u中，否则进行下一轮循环；对于每个子节点，每插入一个新数据，就将其保存的记录数量record_num加1；确定记录所归属的子节点后，每个记录的第p维数据的插入地址可以通过所归属的子节点的参数计算出来，具体为(address+offset+record_num+pitch*p)，其中，p∈[0，k-1]；

步骤3.3.5：清空当前List，重置List中的待划分节点数量为0；对划分过程中产生的所有子节点，即内存空间nodeSpace[i+1]中保存的所有节点进行处理，具体为：申请与nodeSpace[i+1]中节点数量相等的线程数，每个线程处理一个节点；如果节点在第i维上的左右边界相等，则将节点的节点类型设置为无效节点；如果当节点在第i维上的左右边界不相等，但记录数量超过LM，则将节点的节点类型设置为中间节点，并将该节点的地址添加到List中，List的待划分节点数量加1；否则，将节点的节点类型设置为叶子节点。

所述步骤4包括：

步骤4.1：根据List中的待划分节点数量n，申请n个线程，每个线程处理一个待划分节点；在线程编号为tid的线程中，第tid个待划分节点需要划分的子节点数量segNum_tid，计算方式为

其中，record_num是第tid个待划分节点的记录数量，

表示向上取整；

步骤4.2：申请大小为sum*sizeof(midNode)的GPU全局内存空间nodeSpace[k+1]，其中，

sum表示划分过程中产生的子节点的总数量；申请一个辅助数组Associate[]，其大小为n，

tid∈[0，n-1]；

步骤4.3：将n个待划分节点上的数据进行并行划分操作；

步骤4.4：划分结束后，申请与nodeSpace[k+1]中节点数量相等的线程数，每个线程处理一个节点，并行处理所有划分子节点的节点类型；节点类型处理结束后表示数据划分过程结束；所述每个线程处理一个节点，并行处理所有划分子节点的节点类型，具体表示为：如果节点在第k/2维上的左右边界相等，则将节点的节点类型设置为无效节点；否则设置为叶子节点。

所述步骤4.3包括：

步骤4.3.1：申请n个线程进行并行处理，每个线程处理一个待划分节点，第tid个线程中，第tid个待划分节点需要与nodeSpace[k+1]中的第(Associate[tid]+s)个子节点建立关联关系，其中tid∈[0,n-1]，s∈[0,segNum_tid-1]；构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[k+1]一直存活到KD索引的销毁；

步骤4.3.2：按照步骤3.3.2至步骤3.3.4执行；其中，参与划分的维度编号为k/2；并且进行数据处理时，需要将待划分节点tid的子节点数量seg_i替换成对应待划分节点的segNum_tid进行实际处理；步骤3.3.3初始化每个子节点的起始地址时，参与计算的i＝k。

所述步骤5包括：

步骤5.1：根据步骤3中k层中间节点的总数n，申请n个线程，每个线程处理一个中间节点的参数访问优化；

步骤5.2：在线程编号为tid的线程中，线程首先读取待划分节点的等分点数量num_tid和对应的等分点数值，并以一次线性函数y＝ax+b为模型对等分点构成的键值对进行线性函数逼近，使得对线性函数输入查询键值时，可以输出距离查询键值最近的等分点的相对存储地址，其中，a、b均为浮点数，初始化为0.0；

步骤5.3：在线程编号为tid的线程中，使用等分点数值计算出

其中，tmp[t]代表节点的第t个等分点参数，如果累加过程中，出现tmp[t+1]等于tmp[t]，导致

出现无效值的情况，则在累加过程中舍弃这个无效值；将第

个等分点的键值对传入y＝ax+b，计算出b。

本发明的有益效果：

本发明提出了一种GPU上的多维KD树优化方法，提出了额度的概念，并以额度的概念从数据集全局考虑KD树的划分本质，将传统KD树上的二等分划分操作视作一种可分配资源，并从这个新角度出发，提出了一种GPU上优化KD树构建与查询过程的新方法和新技术；另外，从降低GPU上KD树划分时间出发，在保证叶子节点在各维度均有良好约束的前提上，通过提高每次排序的利用率，获得了更大的划分效率，加速了KD树的构建；最后，利用学习索引的思想优化了中间节点的等分点参数的访问方式，有效降低了线程束的分化状况；本专利方法基于GPU硬件特性，通过减少KD树的深度，降低线程束分化程度，提高了KD树在GPU上的查询性能。

附图说明

图1为本发明中一种GPU上的多维KD树的优化方法流程图。

图2为本发明中的进行额度分配时的参数配置图。

图3为本发明中不同层的节点的组织结构示意图。

图4为本发明中并行划分过程中节点参数调整示意图。

图5为本发明中并行划分过程中数据空间调整示意图。

图6为本发明中中间节点的等分点参数及其偏移量映射关系示意图。

图7为本发明中中间节点等分点参数访问优化示意图。

具体实施方式

结合附图对本发明做进一步描述。本发明提出的一种GPU上的多维KD树优化方法，是一种在GPU中加速多维数据上的KD树构建与查询过程的方法。本发明方法利用CPU-GPU异构的模型，充分发挥GPU的并行计算能力，结合使用额度分配策略，贪心划分策略，以及学习索引优化策略，有效提升了KD树在GPU上的构建与查询速度。

如图1所示，一种GPU上的多维KD树优化方法，在预处理数据集后，首先基于额度分配策略对数据集进行k层的并行划分，随后再使用贪心划分策略对划分结果进行一次并行划分处理，最后，使用学习索引优化中间节点的参数访问方式，包括：

步骤1：对于k个维度的数据集A，获得数据集A包含的记录数量n，并通过

计算出数据集A的额度Q，其中LM表示叶子节点可以保存的记录数量上限，

表示向下取整，α表示可调参数，α∈[0.5,1.0]；α用于根据不同数据集调整计算出的额度，默认值为1。额度在理论上等价于：将数据集A视作一个单维数据集，并对其进行二分划分，假设划分维度上的每个数据不重复，那么当叶子节点保存的记录数量在区间[LM,2*LM)之中时的二分树的深度；

步骤2：对k个维度进行排序，并将额度Q分配给k个维度，包括：

步骤2.1：对数据集A的k个维度的数据进行采样，以k个维度上采样数据的均方差作为比较标准，对k个维度进行降序排序，保证均方差较大的维度占据较前的位置，排序结果记为S；较好的采样方法一般需要随着数据集大小和类型变化，这里默认的采样法是按照1％的比例对各维度的数据进行等距抽样；

步骤2.2：将额度Q分配给k个维度，每个维度分配得到的额度不小于

并保证排序较前的维度所分配的额度必然不小于排序较后的维度所分配的额度，第i个维度所分配得到的额度保存为q_i，其中i∈[0,k-1]；如图2所示展示了三维数据集的一种额度分配过程：当数据集大小为1124，叶子节点的LM为1时，通过计算

得到额度10；当维度为3时，我们先平均为每个维度分配3个划分额度，最后一个划分额度分配给均方差最大的维度x，因此，最终额度分配结果为{(x,4),(y,3),(z，3)}。

步骤2.3：对于每一个维度d_i，根据其获得的额度q_i，通过

计算出维度d_i需要划分的子段数量seg_i，维度d_i需要获得(seg_i-1)个等分点，将维度d_i的排序结果S及维度d_i对应的seg_i保存为一个结构体Pre；图2中展示了维度分配的额度以及其应当划分的子段数量的对应关系；

步骤3：基于结构体Pre，在GPU上进行k层的并行循环划分，构建出KD树的主体结构，包括：

步骤3.1：将数据集A的各个维度的数据以及对应的LM、Pre传递到GPU中，其中数据集A的各维度数据按照列存储模式进行实际存储，并保存在全局内存中，使用CUDA内置的内存申请函数cudaMallocPitch()进行内存申请，将返回的各列的间距记作pitch，将申请的全局内存空间记为M，LM和Pre保存在常量内存中，另外，申请一块位于全局内存的空间B作为划分过程中的中间数据存储介质，且满足内存空间不小于M；

步骤3.2：申请一个GPU上的全局内存空间nodeSpace[0]，其大小为一个中间节点的大小，保存有一个初始节点a；将a的地址压入待划分节点链表List，List中每个节点保存有起始地址address、地址偏移量offset两个参数，可以通过基址变址寻址的方法，使用起始地址address加上偏移量(offset+pitch*u)获取节点第u个维度数据的起始地址；设置节点a的起始地址为M，地址偏移量为0；此外，每个待划分节点还保存有当前节点的记录数量以及该节点k个维度上的范围约束，初始节点a在每个维度上的范围约束是(-∞，+∞)，对于每个非初始节点的中间节点，初始化时范围约束继承自父节点的范围约束，并在父节点划分结束之后依据划分参数改变每个子节点在划分维度上的范围约束；由于此时List中只有a一个待划分节点，设置List中的待划分节点数量n＝1；

划分过程中，每个中间节点都不保存实际数据，只保存了用于计算节点所属数据的起始地址的相关参数address和offset；另外，待划分节点链表List本质上是一块较大的连续全局内存，用于保存待划分节点的地址；在第i层划分中，List保存着第i层的所有待划分的中间节点地址，在第i层划分结束后，更新List，使其保存着第i+1层的所有待划分的中间节点地址。

步骤3.3：根据Pre对List进行k层循环划分，每层划分中参与划分的维度按照Pre中的维度顺序S进行依次选择，每层循环中所有待划分节点都在同一维度上进行数据划分；在第i层的并行划分中，参与划分的维度编号设置为i，i∈[0，k-1]；从图4中可以看出，子节点0，子节点1，子节点2都在维度y上进行划分；每层循环具体表述为：

步骤3.3.1：获得参与划分的维度编号i，以及对应的seg_i；根据List中的待划分节点数量n，申请大小为n*seg_i*sizeof(midNode)的全局内存nodeSpace[i+1]，用于存储划分后的n*seg_i个子节点，其中sizeof(·)表示获取某个结构体所占用内存大小，midNode表示节点的结构体，sizeof(midNode)表示获取一个节点所占用的内存大小；每个中间节点都有seg_i个节点指针，用于指向其对应的seg_i个子节点；因此，在申请全局内存nodeSpace[i+1]后，需要建立第j个待划分节点与nodeSpace[i+1]中的第(j*seg_i+s)个子节点的关联关系，其中j∈[0,n-1]，s∈[0，seg_i-1]，这个过程申请n个线程进行并行处理，每个线程处理一个待划分节点，并在构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[i+1]一直存活到KD索引的销毁；图3中展示了当每个待划分节点产生两个子节点时，根节点进行两次划分后的节点组织示意图；在进行第1次并行划分时，List中只有第0层中的一个待划分节点，划分出的子节点是第1层的2个节点，划分结束后，List中的待划分节点更新为第1层中的两个2个待划分节点；图3中，各层的节点主要是按照层来存储的，但层与层之间各个节点，通过设置对应的指针保持着KD树中父子节点的关联性。

步骤3.3.2：依次对n个待划分节点的第i维数据进行并行排序，并选取(seg_i-1)个的等分点，具体为：将每个节点的第i维数据复制到一个可用的全局空间tmp上，使用排序函数Thrust::sort()进行排序操作，并将排序后的数据等分为seg_i个子段，将每个子段在维度i上的范围、对应的记录数量保存到对应子节点的内部参数中，其中，保存子段在维度i上的范围到对应的子节点中，意味着更新子节点在第i维的范围约束；此外，每个子段的范围是半开半闭区间，具体为左开右闭；

步骤3.3.3：启动n个线程，每个线程处理seg_i个归属于同一个待划分节点的子节点；每个线程中，如果i％2＝＝0，初始化每个子节点的起始地址为B，否则，将起始地址设置为M，％表示取余；地址偏移量offset为父节点的地址偏移量；随后将子节点f的地址偏移量offset[f]更新为

f∈[0，seg_i-1],其中，subNode[t].record_num代表第t个子节点的记录数量；更新完待划分节点的所有子节点的offset后，将所有子节点的record_num重置为0；

图4展示了根节点作为待划分节点，进行第一次划分时，其划分出的子节点的部分参数的设置，图中只展示了3个子节点的参数配置；其中，根节点的参数在划分前与划分后都是一样的，并不存在变化，但划分结束后，理论上根节点不持有任何数据，因此，虽然划分后根节点仍可以通过起始地址address、地址偏移量offset访问到数据，但是该操作是无意义的；图4主要展示子节点三个参数的变化，分别是范围约束的变化：由于根节点的待划分维度是x，因此在划分后，所有子节点在维度x上的维度约束都产生了变化；数据空间的变化：由于根节点的起始地址是M，所以子节点的起始地址是空间B，父节点与子节点的起始地址空间始终是不同的；偏移量offset的变化：各个子节点的偏移量首先继承父节点的偏移量0，然后加上编号较前的节点保存的记录数量，比如节点1的偏移量就等于父节点的偏移量0加上节点0的记录数量64；

步骤3.3.4：启动n个线程进行数据的并行划分；每个线程处理一个待划分节点，并执行一个循环，循环次数为待划分节点中的记录数量record_num；第p次循环中，根据设定的判断方法判断第p个记录所属的子节点，并将数据插入到对应的子节点的数据空间中；所述设定的判断方法为：进行一次循环比较，循环次数是待划分节点的子节点数量，在第u次循环中，比较记录f在第i维上的数据是否满足子节点u在第i维上的范围约束，如果满足，则跳出循环，并将第f个记录插入到子节点u中，否则进行下一轮循环；对于每个子节点，每插入一个新数据，就将其保存的记录数量record_num加1；确定记录所归属的子节点后，每个记录的第p维数据的插入地址可以通过所归属的子节点的参数计算出来，具体为(address+offset+record_num+pitch*p)，其中，p∈[0，k-1]；

所有线程中，每个待划分点持有一个线程参与该节点的子节点划分；每个线程执行对应待划分节点的子节点划分，并通过计算数据的插入地址保证数据插入的准确性；图5在每个待划分节点产生两个子节点的情况下，展示了如何利用空间M、B，以及节点偏移量，保证待划分节点的多个子节点中执行数据插入的准确性：在第一次划分中，首先，由于节点0的数据空间是M，因此将子节点的数据空间设置为B，然后将待划分节点0的数据复制到到临时空间tmp中，进行排序并获得中值点，然后计算出节点1应当保存有的数据num0，并将子节点0的偏移量设置为0，子节点1的偏移量设置为num0，从而使得节点1与节点2数据插入的起始地址归属于同一空间B，并且两节点的数据插入的起始地址之间的距离恰好是节点1应当保存有的记录数量，从而保证节点1和节点2的数据空间不会互相覆盖，实现数据插入的正确性；这种双空间划分策略用于在并行划分过程中，优化GPU中动态申请子节点的数据空间的代价；其中，图5中的线程数量指执行划分操作的线程数量，其数量与待划分节点的数量相等；此外，图5中的划分展示的是单个维度上的数据插入过程，但由于数据空间M，B都是按照列存储的形式组织数据的，因此，也可以将图5视作k个维度上每个维度的数据插入过程。

步骤3.3.5：清空当前List，重置List中的待划分节点数量为0；对划分过程中产生的所有子节点，即内存空间nodeSpace[i+1]中保存的所有节点进行处理，具体为：申请与nodeSpace[i+1]中节点数量相等的线程数，每个线程处理一个节点；如果节点在第i维上的左右边界相等，则将节点的节点类型设置为无效节点；如果当节点在第i维上的左右边界不相等，但记录数量超过LM，则将节点的节点类型设置为中间节点，并将该节点的地址添加到List中，List的待划分节点数量加1；否则，将节点的节点类型设置为叶子节点；

步骤3.4：如果当前划分次数小于k，跳转步骤3.3.1进行下一层的并行划分，否则跳转步骤4；

步骤4：选取第k/2个维度作为划分维度，对待划分的节点列表List进行基于贪心策略的划分，包括：

步骤4.1：根据List中的待划分节点数量n，申请n个线程，每个线程处理一个待划分节点；在线程编号为tid的线程中，第tid个待划分节点需要划分的子节点数量segNum_tid，计算方式为

其中，record_num是第tid个待划分节点的记录数量，

表示向上取整；

步骤4.2：申请大小为sum*sizeof(midNode)的GPU全局内存空间nodeSpace[k+1]，其中，

sum表示划分过程中产生的子节点的总数量；申请一个辅助数组Associate[]，其大小为n，

tid∈[0，n-1]；

步骤4.3：将n个待划分节点上的数据进行并行划分操作，包括：

步骤4.3.1：申请n个线程进行并行处理，每个线程处理一个待划分节点，第tid个线程中，第tid个待划分节点需要与nodeSpace[k+1]中的第(Associate[tid]+s)个子节点建立关联关系，其中tid∈[0,n-1]，s∈[0,segNum_tid-1]；构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[k+1]一直存活到KD索引的销毁；

步骤4.3.2：按照步骤3.3.2至步骤3.3.4执行；其中，参与划分的维度编号为k/2；并且进行数据处理时，需要将待划分节点tid的子节点数量seg_i替换成对应待划分节点的segNum_tid进行实际处理；步骤3.3.3初始化每个子节点的起始地址时，参与计算的i＝k；

步骤4.4：划分结束后，申请与nodeSpace[k+1]中节点数量相等的线程数，每个线程处理一个节点，并行处理所有划分子节点的节点类型；节点类型处理结束后表示数据划分过程结束；所述每个线程处理一个节点，并行处理所有划分子节点的节点类型，具体表示为：如果节点在第k/2维上的左右边界相等，则将节点的节点类型设置为无效节点；否则设置为叶子节点；

步骤5：对步骤3中产生的所有中间节点的参数访问进行优化，包括：

步骤5.1：根据步骤3中k层中间节点的总数n，申请n个线程，每个线程处理一个中间节点的参数访问优化；

步骤5.2：在线程编号为tid的线程中，线程首先读取待划分节点的等分点数量num_tid和对应的等分点数值，并以一次线性函数y＝ax+b为模型对等分点构成的键值对(等分点值、相对存储地址)进行线性函数逼近，使得对线性函数输入查询键值时，可以输出距离查询键值最近的等分点的相对存储地址，，其中，a、b均为浮点数，初始化为0.0；

步骤5.3：在线程编号为tid的线程中，使用等分点数值计算出

其中，tmp[t]代表节点的第t个等分点参数，如果累加过程中，出现tmp[t+1]等于tmp[t]，导致

出现无效值的情况，则在累加过程中舍弃这个无效值；将第

个等分点的键值对传入y＝ax+b，计算出b。

图6展示了某节点内部的等分点参数K及其对应的位移P的(K，P)键值对的分布情况，该节点有4个等分点，也就是存在5个子节点；其中，等分点参数{1，11，25，48}的偏移地址分别为{0，1，2，3}；可以看出，(K，P)键值对的分布是一个单调递增的一维分段函数，且每个分段都是等值函数，这种分布简单且已知的(K，P)键值对，对于有监督的模型训练是较为容易的。因此，默认训练模型是比较简单的线性模型y＝ax+b；考虑到并行地为每个中间节点进行有监督的训练是较为复杂的，本发明使用了一种预设的计算公式替代训练过程，用于加速节点内部参数访问的模型参数计算过程；如果等分点参数过多，使用y＝ax+b简单模型不能有效计算，可以将y＝ax+b替换为其他更为复杂的模型，比如一元多次模型，或者单输入单输出的简单神经网络模型，进行有监督的训练，从而达到对模型性能的要求；此外，要求同一层的所有中间节点的参数优化模型都采用同一个模型，这样可以优化该层进行批量查询时的线程束分化情况。图7中展示了某个训练成功的节点参数访问模型F(x)，对于输入13，其输出1；1的实际意义是等分点保存数组的下标，所以我们将13与11对比，发现比11大，然后与25对比，发现比25小，从而结束比较，定位到下一层。

Claims (7)

1.一种GPU上的多维KD树优化方法，其特征在于，包括：

步骤1：对于k个维度的数据集A，获得数据集A包含的记录数量n，并通过

计算出数据集A的额度Q，其中LM表示叶子节点可以保存的记录数量上限，α表示可调参数，α∈[0.5，1.0]，

表示向下取整；

步骤2：对k个维度进行排序，并将额度Q分配给k个维度；

步骤3：在GPU上进行k层的并行循环划分，构建出KD树的主体结构；

步骤4：选取第k/2个维度作为划分维度，对待划分的节点列表List进行基于贪心策略的划分；

步骤5：对步骤3中产生的所有中间节点的参数访问进行优化。

2.根据权利要求1所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤2包括：

步骤2.1：对数据集A的k个维度的数据进行采样，以k个维度上采样数据的均方差作为比较标准，对k个维度进行降序排序，保证均方差较大的维度占据较前的位置，排序结果记为S；

步骤2.2：将额度Q分配给k个维度，每个维度分配得到的额度不小于

并保证排序较前的维度所分配的额度必然不小于排序较后的维度所分配的额度，第i个维度所分配得到的额度保存为q_i，其中i∈[0，k-1]；

步骤2.3：对于每一个维度d_i，根据其获得的额度q_i，通过

计算出维度d_i需要划分的子段数量seg_i，维度d_i需要获得(seg_i-1)个等分点，将维度d_i的排序结果S及维度d_i对应的seg_i保存为一个结构体Pre。

3.根据权利要求1所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤3包括：

步骤3.1：将数据集A的各个维度的数据以及对应的LM、Pre传递到GPU中，其中数据集A的各维度数据按照列存储模式进行存储，使用CUDA内置的内存申请函数cudaMallocPitch()进行内存申请，将返回的各列的间距记作pitch，将申请的全局内存空间记为M，LM和Pre保存在常量内存中，另外，申请一块位于全局内存的空间B作为划分过程中的中间数据存储介质，且满足内存空间不小于M；

步骤3.2：申请一个GPU上的全局内存空间nodeSpace[0]，其大小为一个中间节点的大小，保存有一个初始节点a；将a的地址压入待划分节点链表List，List中每个节点保存有起始地址address、地址偏移量offset两个参数，可以通过基址变址寻址的方法，使用起始地址address加上偏移量(offset+pitch*u)获取节点第u个维度数据的起始地址；设置节点a的起始地址为M，地址偏移量为0；此外，每个待划分节点还保存有当前节点的记录数量以及该节点k个维度上的范围约束，初始节点a在每个维度上的范围约束是(-∞，+∞)，对于每个非初始节点的中间节点，初始化时范围约束继承自父节点的范围约束，并在父节点划分结束之后依据划分参数改变每个子节点在划分维度上的范围约束；由于此时List中只有a一个待划分节点，设置List中的待划分节点数量n＝1；

步骤3.3：根据Pre对List进行k层循环划分，每层划分中参与划分的维度按照Pre中的维度顺序S进行依次选择，每层循环中所有待划分节点都在同一维度上进行数据划分；在第i层的并行划分中，参与划分的维度编号设置为i，i∈[0，k-1]；

步骤3.4：如果当前划分次数小于k，跳转步骤3.3.1进行下一层的并行划分，否则跳转步骤4。

4.根据权利要求3所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤3.3中每层循环具体表述为：

步骤3.3.1：获得参与划分的维度编号i，以及对应的seg_i；根据List中的待划分节点数量n，申请大小为n*seg_i*sizeof(midNode)的全局内存nodeSpace[i+1]，用于存储划分后的n*seg_i个子节点，其中sizeof(·)表示获取某个结构体所占用内存大小，midNode表示节点的结构体，sizeof(midNode)表示获取一个节点所占用的内存大小；每个中间节点都有seg_i个节点指针，用于指向其对应的seg_i个子节点；因此，在申请全局内存nodeSpace[i+1]后，需要建立第j个待划分节点与nodeSpace[i+1]中的第(j*seg_i+s)个子节点的关联关系，其中j∈[0，n-1]，s∈[0，seg_i-1]，这个过程申请n个线程进行并行处理，每个线程处理一个待划分节点，并在构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[i+1]一直存活到KD索引的销毁；

步骤3.3.2：依次对n个待划分节点的第i维数据进行并行排序，并选取(seg_i-1)个的等分点，具体为：将每个节点的第i维数据复制到一个可用的全局空间tmp上，使用排序函数Thrust：：sort()进行排序操作，并将排序后的数据等分为seg_i个子段，将每个子段在维度i上的范围、对应的记录数量保存到对应子节点的内部参数中，其中，保存子段在维度i上的范围到对应的子节点中，意味着更新子节点在第i维的范围约束；此外，每个子段的范围是半开半闭区间，具体为左开右闭；

步骤3.3.3：启动n个线程，每个线程处理seg_i个归属于同一个待划分节点的子节点；每个线程中，如果i％2＝＝0，初始化每个子节点的起始地址为B，否则，将起始地址设置为M，％表示取余；地址偏移量offset为父节点的地址偏移量；随后将子节点f的地址偏移量offset[f]更新为

f∈[0，seg_i-1]，其中，subNode[t].record_hum代表第t个子节点的记录数量；更新完待划分节点的所有子节点的offset后，将所有子节点的record_hum重置为0；

步骤3.3.4：启动n个线程进行数据的并行划分；每个线程处理一个待划分节点，并执行一个循环，循环次数为待划分节点中的记录数量record_num；第p次循环中，根据设定的判断方法判断第p个记录所属的子节点，并将数据插入到对应的子节点的数据空间中；所述设定的判断方法为：进行一次循环比较，循环次数是待划分节点的子节点数量，在第u次循环中，比较记录f在第i维上的数据是否满足子节点u在第i维上的范围约束，如果满足，则跳出循环，并将第f个记录插入到子节点u中，否则进行下一轮循环；对于每个子节点，每插入一个新数据，就将其保存的记录数量record-num加1；确定记录所归属的子节点后，每个记录的第p维数据的插入地址可以通过所归属的子节点的参数计算出来，具体为(address+offset+record_hum+pitch*p)，其中，p∈[0，k-1]；

步骤3.3.5：清空当前List，重置List中的待划分节点数量为0；对划分过程中产生的所有子节点，即内存空间nodeSpace[i+1]中保存的所有节点进行处理，具体为：申请与nodeSpace[i+1]中节点数量相等的线程数，每个线程处理一个节点；如果节点在第i维上的左右边界相等，则将节点的节点类型设置为无效节点；如果当节点在第i维上的左右边界不相等，但记录数量超过LM，则将节点的节点类型设置为中间节点，并将该节点的地址添加到List中，List的待划分节点数量加1；否则，将节点的节点类型设置为叶子节点。

5.根据权利要求1所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤4包括：

步骤4.1：根据List中的待划分节点数量n，申请n个线程，每个线程处理一个待划分节点；在线程编号为tid的线程中，第tid个待划分节点需要划分的子节点数量segNum_tid，计算方式为

其中，record_hum是第tid个待划分节点的记录数量，

表示向上取整；

步骤4.2：申请大小为sum*sizeof(midNode)的GPU全局内存空间nodeSpace[k+1]，其中，

sum表示划分过程中产生的子节点的总数量；申请一个辅助数组Associate[]，其大小为n，

tid∈[0，n-1]；

步骤4.3：将n个待划分节点上的数据进行并行划分操作；

步骤4.4：划分结束后，申请与nodeSpace[k+1]中节点数量相等的线程数，每个线程处理一个节点，并行处理所有划分子节点的节点类型；节点类型处理结束后表示数据划分过程结束；所述每个线程处理一个节点，并行处理所有划分子节点的节点类型，具体表示为：如果节点在第k/2维上的左右边界相等，则将节点的节点类型设置为无效节点；否则设置为叶子节点。

6.根据权利要求5所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤4.3包括：

步骤4.3.1：申请n个线程进行并行处理，每个线程处理一个待划分节点，第tid个线程中，第tid个待划分节点需要与nodeSpace[k+1]中的第(Associate[tid]+s)个子节点建立关联关系，其中tid∈[0，n-1]，s∈[0，segNum_tid-1]；构建好父节点与子节点的关联关系后，通过继承父节点范围约束的方法，初始化子节点的范围约束；nodeSpace[k+1]一直存活到KD索引的销毁；

步骤4.3.2：按照步骤3.3.2至步骤3.3.4执行；其中，参与划分的维度编号为k/2；并且进行数据处理时，需要将待划分节点tid的子节点数量seg_i替换成对应待划分节点的segNum_tid进行实际处理；步骤3.3.3初始化每个子节点的起始地址时，参与计算的i＝k。

7.根据权利要求1所述的一种GPU上的多维KD树优化方法，其特征在于，所述步骤5包括：

步骤5.1：根据步骤3中k层中间节点的总数n，申请n个线程，每个线程处理一个中间节点的参数访问优化；

步骤5.2：在线程编号为tid的线程中，线程首先读取待划分节点的等分点数量num_tid和对应的等分点数值，并以一次线性函数y＝ax+b为模型对等分点构成的键值对进行线性函数逼近，使得对线性函数输入查询键值时，可以输出距离查询键值最近的等分点的相对存储地址，其中，a、b均为浮点数，初始化为0.0；

步骤5.3：在线程编号为tid的线程中，使用等分点数值计算出

其中，tmp[t]代表节点的第t个等分点参数，如果累加过程中，出现tmp[t+1]等于tmp[t]，导致

出现无效值的情况，则在累加过程中舍弃这个无效值；将第

个等分点的键值对传入y＝ax+b，计算出b。

Images