CN102955686B

CN102955686B - 一种N‑body问题近程作用计算在GPU结构的优化映射方法

Info

Publication number: CN102955686B
Application number: CN201110238072.6A
Authority: CN
Inventors: 雷咏梅; 单健晨; 朱劲师; 方韵
Original assignee: Shanghai Redneurons Co Ltd; University of Shanghai for Science and Technology
Current assignee: Shanghai Redneurons Co Ltd; University of Shanghai for Science and Technology
Priority date: 2011-08-18
Filing date: 2011-08-18
Publication date: 2017-04-05
Anticipated expiration: 2031-08-18
Also published as: CN102955686A

Abstract

本发明涉及一种N‑body问题近程作用计算在GPU结构的优化映射方法，用于计算粒子受到本盒子中其他粒子以及邻居盒子中粒子的近程作用，所述方法包括建立数组Neighbor List，依次存储编号大于当前盒子编号的邻居盒子信息；建立数组acc记录全局存储器中粒子所受其他粒子近程作用的计算结果；建立多个block与所有盒子的一一对应关系，使每个block处理与其对应的盒子中粒子近程作用的计算。本发明提出的通过高速存取的片上共享存储来复用数据，解决了CPU算法访存的瓶颈，利用牛顿第三定律减少了算法的计算量，总体时间得到了进一步优化，实现近程计算部分在GPU结构上的高效执行，提高了计算效率。

Description

一种N-body问题近程作用计算在GPU结构的优化映射方法

技术领域

本发明涉及GPU并行计算领域，尤其涉及一种N-body问题近程作用计算在GPU(Graphic Processing Unit，图形处理器)结构的优化映射方法。

背景技术

随着技术的发展，众核GPU已经成为目前重要的处理器。传统上GPU只用于处理3D渲染任务，而其它大多数的任务都交给了CPU处理。CPU作为一种通用处理器，其体系结构设计必须兼顾各种任务的需要，因此CPU中大多数晶体管被用于制造庞大的缓存和复杂的控制逻辑，而运算单元占用面积则并不多；相反地，GPU由于图形渲染的并行特性与生俱来具有大量运算单元，非常适合计算密集型的大规模数据并行计算。但由于传统GPU硬件架构的限制很难有效利用其资源进行通用计算，为此，NVIDIA公司推出了CUDA(Compute UnifiedDevice Architecture，通用并行计算架构)完全扭转了局面，由于其功能强大、应用领域广泛，推动了GPU通用计算的发展。但如何高效地在GPU架构上实现算法是研究人员当前面临的巨大挑战。

N-body问题是高性能计算领域最具代表性、最有影响力以及最有挑战性的问题之一，具有广泛的应用领域，并因其应用的重要性和实际计算的复杂性一直是高性能计算的主要热点之一。N-body问题主要应用在天体物理学、分子动力学等领域，多体问题是计算场中的多个粒子之间的相互作用及其运动轨道，是最具普适性的力学问题之一。当粒子为宏观的天体时，天体多体模拟计算是当前研究星系以及宇宙结构形成的主要途径。当粒子为微观的分子、原子时，多体问题即表现为分子动力学问题，由于分子动力学可以预测纳米尺度材料动力学特性，因此在物理、化学、生物、医药、新材料设计等领域有着广泛的应用。N-body问题主要解决有限空间中每个粒子所受其它粒子作用的计算，算法复杂度为O(N²)，由于当空间中粒子数很大时计算量非常巨大，因此为了减少计算量，提出了树形算法，其中最主要的Fast Multipole Method(快速多极算法，以下简称FMM算法)将算法复杂度降低到了O(N)。树形算法的核心思想是通过对空间进行划分将粒子所受作用分类为近程作用和远程作用，近程作用直接计算而远程作用则近似计算。

目前关于N-body的并行计算主要采用基于MPI编程模型的CPU集群，存在大量的通信和负载的不均衡造成的性能瓶颈，而以GRAPE为代表的专用机虽然取得较高的性能但并未对整个树形算法加速，并且专用机本身应用的局限性提高了设备的成本。

GPU加速平台提供解决上述问题的途径，但现有对求解N-body问题的FMM算法进行CUDA加速的软件对于近程计算部分存在访存的瓶颈，并且对计算量未进行优化，因此计算性能不理想。

发明内容

为了解决上述现有技术存在的问题，本发明旨在提供一种N-body问题近程作用计算在GPU结构的优化映射方法，以高效实现CPU算法到GPU结构映射的计算任务划分和数据存取，从而为求解典型N-body应用问题提供多级并行求解方法。

本发明所述的一种N-body问题近程作用计算在GPU结构的优化映射方法，用于计算粒子受到本盒子中其他粒子以及邻居盒子中粒子的近程作用，所述方法包括以下步骤：

步骤S0，建立数组Neighbor List，依次存储编号大于当前盒子编号的邻居盒子信息；建立数组acc记录全局存储器中粒子所受其他粒子近程作用的计算结果；建立多个block与所有盒子的一一对应关系，使每个block处理与其对应的盒子中粒子近程作用的计算，每个block中的线程数不小于与其对应的盒子中的粒子数，每个盒子中的粒子数均为n；

步骤S1，对应当前盒子的block在共享存储器中申请大小均为n的数组A和数组B，其中，数组A用于保存当前盒子中的n个粒子的信息，数组B用于存放当前盒子中每个粒子所受当前盒子中其他n-1个粒子的反作用力数据；

步骤S2，建立对应当前盒子的block中的前n个线程与当前盒子中的n个粒子的一一对应关系，使每个线程将与其对应的当前盒子中的粒子信息装载到共享存储器中的数组A中；

步骤S3，对应当前盒子的block的前n个线程中的每个线程在寄存器中申请一个局部变量d_acc，每个局部变量d_acc用于保存当前线程对应的当前盒子中粒子所受的作用力，设当前线程对应的粒子为第threadIdx.x个粒子，threadIdx.x的取值范围为0至n-1，然后当前线程从共享存储器的数组A中顺序读取第threadIdx.x个粒子后面的n-threadIdx.x-1个粒子的信息，并分别将n-threadIdx.x-1个粒子与第threadIdx.x个粒子进行近程作用计算；每对粒子间的近程作用计算结果同时更新到局部变量d_acc和数组B；完成当前盒子中所有粒子间的近程作用计算后，对应当前盒子的block的前n个线程中的每个线程将数组B中存放的与该线程对应的当前盒子中粒子所受的反作用力数据叠加到数组acc中的元素acc[i]的位置，i表示当前线程所对应的当前盒子中的粒子；

步骤S4，根据数组Neighbor List，建立对应当前盒子的block中的前n个线程与当前盒子的下一个邻居盒子中的n个粒子的一一对应关系，使每个线程将与其对应的邻居盒子中粒子信息装载到共享存储器中的数组A中，以覆盖数组A中当前盒子中粒子的信息，同时，将数组B中存放的数据清零，使数组B用于存放当前盒子的下一个邻居盒子中每个粒子所受当前盒子中n个粒子的反作用力数据；

步骤S5，对应当前盒子的block的前n个线程中的每个线程从共享存储器的数组A中顺序读取当前盒子的下一个邻居盒子中的n个粒子的信息，并将该邻居盒子中的n个粒子分别与当前线程对应的当前盒子中的粒子进行近程作用计算；每对粒子间的计算结果需要同时更新到局部变量d_acc和数组B；完成当前盒子中n个粒子与当前盒子的下一个邻居盒子中n个粒子间的近程作用计算后，对应当前盒子的block的前n个线程中的每个线程将数组B中存放的与该线程对应的邻居盒子中粒子所受的反作用力数据叠加到数组acc中的元素acc[j]位置，j表示当前线程所对应的邻居盒子中的粒子；

步骤S6，根据数组Neighbor List，判断是否所有编号大于当前盒子编号的邻居盒子中的粒子信息都已经被装载到过共享存储器中用来与当前盒子中的粒子进行近程作用计算，若是，则当前盒子中粒子受到邻居盒子中粒子的近程作用计算完成，否则回到步骤S4继续装载再下一个邻居盒子中的粒子信息并依次执行步骤S5和步骤S6，直至完成当前盒子中粒子与其所有邻居盒子中粒子间的近程作用计算；

步骤S7，对应当前盒子的block的前n个线程中的每个线程将该线程申请的局部变量d_acc一一对应地写入全局存储器的数组acc中的元素acc[i]的位置，i表示当前线程所对应的当前盒子中的粒子。

在上述的N-body问题近程作用计算在GPU结构的优化映射方法中，所述每个block中的前n个线程处理其对应盒子中的n个粒子，每个block中的其余线程闲置。

在上述的N-body问题近程作用计算在GPU结构的优化映射方法中，所述每个block中的线程数最大为512。

由于采用了上述的技术解决方案，本发明是一种根据牛顿第三定律所构建的基于GPU硬件加速的N-body问题的近程计算并行求解方法，本发明采用基于GPU可编程线程实现N-body近程计算过程的加速，提出GPU数据共享存储方案，数据访问流程，近程力计算编码方案均基于GPU可编程实现。本发明针对GPU存储空间有限及单指令多数据的执行方式，提出复用共享存储数据处理方法，与已有方法不同，本发明的方法对每个盒子的粒子近程作用计算时，可连续使用共享数据，减少数据存储访存次数，从而提高并行计算过程实现速度。

附图说明

图1(a)是现有技术中计算粒子所受本盒子其它粒子近程作用计算的线程示意图；

图1(b)是本发明中计算粒子所受本盒子其它粒子近程作用计算的具体实施例的线程示意图；

图2是采用本发明计算粒子所受本盒子其它粒子近程作用计算的具体实施例的数据读取流图；

图3是采用发明计算粒子所受邻居盒子中粒子近程作用计算的具体实施例的数据读取流图。

具体实施方式

下面结合附图，给出本发明的较佳实施例，并予以详细描述。

以下先对本发明的思路进行介绍：本发明研究的是N-body问题在二维空间的粒子间作用力计算；由于物体间的互相作用遵守牛顿第三定律，因此本发明提出根据这个物理原理构建GPU映射方法，即提出N-body问题树形算法中近程作用计算的数据存储策略与任务划分粒度，建立CPU算法到GPU的映射的任务分解方式与数据结构，构建线程执行流程。

牛顿第三定律的内容主要有：两个物体之间的作用力和反作用力，总是同时在同一条直线上，大小相等，方向相反。即F1＝-F2(N＝N’)定律包括以下几方面：(1)力的作用是相互的，同时出现，同时消失；(2)相互作用力一定是相同性质的力；(3)作用力和反作用力作用在两个物体上，产生的作用不能相互抵消；(4)作用力也可以叫做反作用力，只是选择的参照物不同；(5)作用力和反作用力因为作用点不在同一个物体上，所以不能求合力。

通过对牛顿第三定律的描述，可以发现当计算粒子A受到来自粒子B的作用力为F时，粒子B受到来自粒子A的作用力为-F不用再重复计算，依照这个原理在近程作用计算的过程中总的计算量可以减少为原来的1/2，本发明提出近程作用计算的算法，分别从粒子受到本盒子粒子作用和粒子受到邻居盒子粒子的作用进行优化，实现计算量减半的结果。

本发明基于任务分解和数据划分的策略，高效实现了将CPU算法映射到GPU结构，提出了算法高效映射到GPU结构上的原则与实现机制，即N-body问题近程作用计算在GPU上的数据存储方案和算法映射过程的同步操作机制，建立全局存储与共享存储的数据结构，构建线程执行流程。

上述的任务分解策略为尽可能细化任务粒度从而获得更高的算法的并行度和计算密度，同时任务划分还需要保证各线程的负载均衡；数据存取原则从提高存储器利用率角度优化算法的数据存取，主要包括最小化低带宽数据传输和最大化设备共享存储器，其中：

最小化低带宽数据传输，即最小化主机和设备之间的数据传输，因此对于主机与GPU结构之间的数据存取，为了提高传输效率，应尽量通过数据结构的转换用传输整块大数据块来代替传输分组的小数据块；

最大化设备共享存储器的使用，可以尽可能地减少设备和全局存储器之间的数据传输，这也就意味着对于GPU内数据的存取，可以通过可在片上缓存的常量存储器存储数据量较小的只读数据来优化空间局部行为，并且寻找线程间公用的数据集合将其装载到高速共享存储器用数据复用的手段加速GPU内数据的存取。

根据最优的存储器访问模式，尽可能以最优方式组织存储器访问，这是由于根据各类存储器的访问模式的不同，有效带宽可能会出现数量级的变化。这种优化对于全局存储器访问尤为重要，因为全局存储器的带宽较低，其延迟可能达到数百个时钟周期。对于共享存储器来说通常只有在共享存储器访问中存在严重的存储体冲突时才应进行共享存储器访问优化。

下面对本发明，即一种N-body问题近程作用计算在GPU结构的优化映射方法进行说明，本发明用于计算粒子受到本盒子中其他粒子以及邻居盒子中粒子的近程作用，所述方法包括以下步骤：

步骤S0，建立数组Neighbor List，依次存储编号大于当前盒子编号的邻居盒子信息；建立数组acc记录全局存储器中粒子所受其他粒子近程作用的计算结果；建立多个block与所有盒子的一一对应关系，使每个block处理与其对应的盒子中粒子近程作用的计算，每个block中的线程数不小于与其对应的盒子中的粒子数，每个盒子中的粒子数均为n；每个block中的前n个线程处理其对应盒子中的n个粒子，每个block中的其余线程闲置(由于CUDA中定义每个block中的线程数最大为512，所以每个盒子中的粒子数不能超过512个)；

由于实现算法到GPU结构映射的过程中需要将系统中粒子信息传输到GPU，但如果直接将指针数组传输到GPU上，则需在全局存储器上申请许多不连续的小的存储空间，将指针数组以许多小存储块的形式传输到申请的存储空间，这种方式效率较低，因此在数据传输前需将数据结构转换为一维线性数据块，以盒子编号排序从盒子0中所有的粒子开始后面跟着盒子1中所有的粒子依此类推，并创建一个盒子索引用来记录每个盒子中粒子一维数组中的起始索引，将盒子索引和邻居列表(即数组Neighbor List)一起存储于全局存储器。

CUDA架构，即GPU结构的线程执行模型中，线程可以从GPU的存储器中取数据并对数据进行处理。一个线程映射到一个Stream Processor(SP)中执行。一组线程组成一个Block，一个Block映射到由一组SP组成的Streaming Multiprocessor(SM)，SM带有共享存储器，因此一个Block中的所有线程都能共享其中的数据。存储优化可将将邻居列表和盒子索引这两个线性存储块存放到常量寄存器，因为是只读的且容量较小，适合存储在容量有限的常量存储器，可在片上缓存提高访存效率，但受限于常量存储器的容量只能处理有限的盒子数，因此当盒子数较多时仍然存储在全局存储器中。

映射过程应用一个block处理一个盒子中粒子近程作用的计算，同步操作保证所有粒子受本盒子其它粒子近程作用计算部分完成；然后，各线程分别从共享存储器中顺序读取邻居盒子中的粒子进行近程作用计算，并做同步操作保证所有粒子受第一个邻居盒子中粒子的近程作用计算部分完成；若所有邻居盒子都已经被装载到过共享存储器用来计算，则计算完成；否则装载下一个邻居盒子的粒子到共享存储器来覆盖之前的盒子信息并做计算。按照数据装载到共享存储器的顺序，每次装载一个盒子中的粒子信息，因此盒子中粒子个数还受限于共享存储器的大小，例如，具体线程处理流程为线程0处理盒子0中的第一个粒子，即粒子0，首先装载粒子0的信息到共享存储器同步后，计算本盒子其它粒子的近程作用，假设盒子0有2个邻居那么依次装载邻居盒子到共享存储器进行计算。

在计算粒子所受所属盒子中其它粒子作用力的计算过程中，如图1所示T_i表示计算粒子i所受作用力的线程，P_i表示受力粒子i(i取为0-7)，阴影方块表示计算P_i所受对应P_j的作用力，P_j表示施力粒子j(j取为0-7，且j不等于i)。图1(a)是改进前的算法每个线程需要计算对应粒子所受本盒子所有其它粒子的作用力，T_i计算P_i受到P_j的作用力时仅叠加P_i所受的作用力，此时通过牛顿第三定律的改进，T_i需要同时更新P_i和P₁所受的作用力，那么T₁就不需要重复计算P_i对P_j的作用力了，改进后每个线程的计算内容如图1(b)所示，可见阴影方块少了一半，即计算量减半。

在计算粒子所受邻居盒子中粒子的作用力时，通过结合牛顿第三定律改进，遵守如下原则：当计算盒子A中粒子受到盒子B中粒子的作用时，同时更新盒子A和盒子B中粒子受到的来自于对方盒子中粒子的作用力，因此无需重复计算盒子B中粒子受到来自于盒子A中粒子的作用。根据划分后的粒子空间被Z-SFC曲线线性化，按照编码顺序规划每个盒子中粒子所受邻居盒子中粒子作用的计算过程，而且计算盒子A中某个粒子受其邻居盒子中粒子的近程作用时只计算编号大于A的邻居盒子即可，因此总体被计算的邻居盒子的个数为原来改进前算法的1/2，而改进前被计算的邻居盒子总个数N为：

N＝(2^level-1)(2^level+1-1)*4

其中，level为数的高度。

下面举例对上述步骤分别从粒子受到本盒子粒子作用和粒子受到邻居盒子粒子的作用两个部分进行详细解释。

如图2所示，设盒子A中的6个粒子，依次为i、i+1、。。。、i+5，这些粒子的信息存储在全局存储器中；对应盒子A的block在共享存储器中申请大小均为6的数组A和数组B，数组A用于存放盒子A中6个粒子的信息，数组B用于存放盒子A中每个粒子受到盒子A中其他5个粒子的反作用力数据。

对应盒子A的block中的前6个线程T₀-T₅分别从全局存储器读取盒子A中的6个粒子的信息，并一一对应地将信息装载到共享存储器中的数组A中，即存放在元素A[0]-A[5]的位置。

对应盒子A的block中对应粒子i的第0个线程T₀申请一个局部变量d_acc来保存粒子i所受的作用力，然后从数组A中顺序读取粒子i后面的5个粒子的信息，即粒子i+1、。。。、i+5的信息，并将这5个粒子分别与粒子i进行近程作用计算；每对粒子间的计算结果同时更新到局部变量d_acc和数组B，即第0个线程T₀申请的局部变量d_acc的最终值为粒子i受到粒子i+1、。。。、i+5的作用力的总和，数组B中的元素B[1]-B[5]的位置依次存放粒子i+1、。。。、i+5分别受到粒子i的反作用力数据。

对应盒子A的block中对应粒子i+1的第1个线程T₁申请一个局部变量d_acc来保存粒子i+1所受的作用力，然后从数组A中顺序读取粒子i+1后面的4个粒子的信息，即粒子i+2、。。。、i+5的信息，并将这4个粒子分别与粒子i+1进行近程作用计算；每对粒子间的计算结果同时更新到局部变量d_acc和数组B，即第1个线程T₁申请的局部变量d_acc的最终值为粒子i+1受到粒子i+2、。。。、i+5的作用力的总和，数组B中的元素B[2]-B[5]的位置依次叠加存放粒子i+2、。。。、i+5分别受到粒子i+1的反作用力数据。

以此类推，完成盒子A中的6粒子间的近程作用计算；然后，对应盒子A的block中的前6个线程T₀-T₅分别将数组B中6个元素位置所存放的数据依次一一对应地写入全局存储器的数组acc中元素acc[i]-acc[i+5]的位置；例如，第2个线程T₂将数组B中元素B[2]的位置所存放的粒子i+2所受的所有反作用力数据，即粒子i+2受到粒子i、i+1的反作用力的总和写入数组acc中元素acc[i+2]的位置。

预设数组Neighbor List中存储编号大于当前盒子编号的邻居盒子信息，本实施例中，设编号大于盒子A编号的邻居盒子只有一个，即为盒子B，则数组Neighbor List中仅存储盒子B的信息。

如图3所示，设盒子B中的6个粒子，依次为j、j+1、。。。、j+5，这些粒子的信息存储在全局存储器中。

对应盒子A的block中的前6个线程T₀-T₅分别从全局存储器读取盒子B中的6个粒子的信息，并一一对应地将信息装载到共享存储器中的数组A中，即存放在元素A[0]-A[5]的位置，从而覆盖原先位置上盒子A中6个粒子的信息；同时使数组B中的每个元素位置上存放的数据清零，并使数组B中存放盒子B中每个粒子受到盒子A中6个粒子的反作用力数据。

对应盒子A的block中对应粒子i的第0个线程T₀从数组A中顺序读取盒子B中的6个粒子的信息，即粒子j、j+1、。。。、j+5的信息，并将这6个粒子分别与粒子i进行近程作用计算；每对粒子间的计算结果同时更新到局部变量d_acc和数组B，即第0个线程T₀申请的局部变量d_acc的最终值为粒子i受到粒子i+1、。。。、i+5以及粒子j、j+1、。。。、j+5的作用力的总和，数组B中的元素B[0]-B[5]的位置依次存放粒子j、j+1、。。。、j+5分别受到粒子i的反作用力数据。

对应盒子A的block中对应粒子i+1的第1个线程T₁从数组A中顺序读取盒子B中的6个粒子的信息，即粒子j、j+1、。。。、j+5的信息，并将这6个粒子分别与粒子i+1进行近程作用计算；每对粒子的计算结果同时更新到局部变量d_acc和数组B，即第1个线程T₁申请的局部变量d_acc的最终值为粒子i+1受到粒子i+2、。。。、i+5以及粒子j、j+1、。。。、j+5的作用力的总和，数组B中的元素B[0]-B[5]的位置依次叠加存放粒子j、j+1、。。。、j+5分别受到粒子i+1的反作用力数据。

以此类推，完成盒子A中的6粒子与盒子B中的6个粒子间的近程作用计算；然后，对应盒子A的block中的前6个线程T₀-T₅分别将数组B中6个元素位置所存放的数据依次一一对应地写入全局存储器的数组acc中元素acc[j]-acc[j+5]的位置；例如，第2个线程T₂将数组B中B[2]的位置所存放的粒子j+2所受的所有反作用力数据，即粒子j+2受到的粒子i、i+1、。。。、i+5的反作用力的总和写入数组acc中元素acc[j+2]的位置。

根据数组Neighbor List中存储的信息可知，盒子A的所有邻居盒子中粒子已经都被装载到过共享存储器中用来与盒子A中的粒子进行近程作用计算，因此，盒子A中粒子受到邻居盒子中粒子的近程作用计算完成。

至此执行最后一步，即对应盒子A的block中的前6个线程T₀-T₅分别将它们各自申请的局部变量d_acc一一对应地写入全局存储器的数组acc中元素acc[i]-acc[i+5]的位置，此时元素acc[i]-acc[i+5]的位置存放的即为盒子A中粒子i、i+1、。。。、i+5分别受到的盒子A中其他粒子的作用力和反作用力以及盒子B中所有粒子的作用力数据。

以盒子B作为当前盒子，对应盒子B的block中的前6个线程在处理盒子B中6个粒子j、j+1、。。。、j+5间的近程作用计算时，与上述过程相同，完成盒子B中的6粒子间的近程作用计算后，全局存储器的数组acc中元素acc[j]-acc[j+5]的位置存放的即为盒子B中粒子j、j+1、。。。、j+5分别受到的盒子B中其他粒子的反作用力以及盒子A中所有粒子的反作用力数据；根据数组Neighbor List中存储的信息可知，盒子B没有编号大于其编号的邻居盒子，因此，直接执行最后一步，即对应盒子B的block中的前6个线程分别将它们各自申请的局部变量d_acc一一对应地写入全局存储器的数组acc中元素acc[j]-acc[j+5]的位置，此时元素acc[j]-acc[j+5]的位置存放的即为盒子B中粒子j、j+1、。。。、j+5分别受到的盒子B中其他粒子的作用力和反作用力以及盒子A中所有粒子的反作用力数据。

若盒子A具有两个编号大于其编号的邻居盒子，并依次为盒子B和盒子C(盒子C的编号大于盒子B的编号)，则数组Neighbor List中依次存储盒子B和盒子C的信息。

那么在完成盒子A中的6粒子与盒子B中的6个粒子间的近程作用计算后，根据数组Neighbor List中存储的信息可知，对应盒子A的block中的前6个线程T₀-T₅还需要进行盒子A中的6粒子与盒子C中的6个粒子间的近程作用计算，此时，重复上述步骤，将盒子C中的6个粒子的信息装载共享存储器中的数组A中，即存放在元素A[0]-A[5]的位置，从而覆盖原先位置上盒子B中6个粒子的信息；同时置数组B中的每个元素位置上存放的数据清零，并使数组B中存放盒子C中每个粒子受到盒子A中6个粒子的反作用力数据，然后做相应的近程作用计算即可。当执行最后一步后，数组acc中元素acc[i]-acc[i+5]的位置存放的即为盒子A中粒子i、i+1、。。。、i+5分别受到的盒子A中其他粒子的作用力和反作用力以及盒子B、盒子C中所有粒子的作用力数据。

同理可得，处理盒子B中粒子的近程作用计算时，只需计算盒子B中粒子受到的盒子B中其他粒子的作用力和反作用力以及盒子C的作用力即可；而在处理盒子C中粒子的近程作用计算时，只需计算盒子C中粒子受到的盒子C中其他粒子的作用力和反作用力即可。

由此可见，以上执行过程中，计算量仅为原来一半，寄存器的读写操作数也减为原来一半，但由于寄存器的读写延迟非常低，因此在性能评估中可忽略不计。但对于共享存储器和全局存储器各自总的读写操作数是增加还是减少和树的层数以及每个盒子包含的粒子数有关，从而使得算法从计算量和访存量都得到了优化。

综上所述，本发明通过数据结构转换将分组的小数据块转换为整块大数据块来优化主机和GPU之间的数据存取，并提出了将邻居列表和盒子所以存储到常量存储器的优化方案，然后基于CUDA的编程模型，利用共享存储器的数据复用优化GPU内数据的存取。本发明提出的通过高速存取的片上共享存储来复用数据，解决了算法访存的瓶颈，利用牛顿第三定律减少了算法的计算量，总体时间得到了进一步优化，实现近程计算部分在GPU结构上的高效执行，提高了计算效率。

以上所述的，仅为本发明的较佳实施例，并非用以限定本发明的范围，本发明的上述实施例还可以做出各种变化。即凡是依据本发明申请的权利要求书及说明书内容所作的简单、等效变化与修饰，皆落入本发明专利的权利要求保护范围。本发明未详尽描述的均为常规技术内容。

Claims

1.一种N-body问题近程作用计算在GPU结构的优化映射方法，用于计算粒子受到本盒子中其他粒子以及邻居盒子中粒子的近程作用，其特征在于，所述方法包括以下步骤：

步骤S3，对应当前盒子的block的前n个线程中的每个线程分别在寄存器中申请一个局部变量d_acc，每个局部变量d_acc用于保存其所属当前线程对应的当前盒子中一个粒子所受的近程作用力，设当前线程对应的粒子为第threadIdx.x个粒子，threadIdx.x的取值范围为0至n-1，然后当前线程从共享存储器的数组A中顺序读取第threadIdx.x个粒子后面的n-threadIdx.x-1个粒子的信息，并分别将n-threadIdx.x-1个粒子与第threadIdx.x个粒子进行近程作用计算；每对粒子间的近程作用计算结果同时叠加到相应线程的局部变量d_acc和数组B中相应位置；完成当前盒子中所有粒子间的近程作用计算后，对应当前盒子的block的前n个线程中的每个线程分别将数组B中存放的与该线程对应的当前盒子中粒子所受的反作用力数据叠加到全局存储器的数组acc中的元素acc[i]的位置，i表示当前线程所对应的当前盒子中的粒子编号；

步骤S5，对应当前盒子的block的前n个线程中的每个线程从共享存储器的数组A中顺序读取当前盒子的下一个邻居盒子中的n个粒子的信息，并将该邻居盒子中的n个粒子分别与当前线程对应的当前盒子中的粒子进行近程作用计算；每对粒子间的计算结果需要同时叠加到局部变量d_acc和数组B；完成当前盒子中n个粒子与当前盒子的下一个邻居盒子中n个粒子间的近程作用计算后，对应当前盒子的block的前n个线程中的每个线程将数组B中存放的与该线程对应的邻居盒子中粒子所受的反作用力数据叠加到数组acc中的元素acc[j]位置，j表示当前线程所对应的邻居盒子中的粒子；

步骤S7，对应当前盒子的block的前n个线程中的每个线程将该线程申请的局部变量d_acc一一对应地叠加到全局存储器的数组acc中的元素acc[i]的位置，i表示当前线程所对应的当前盒子中的粒子编号。

2.根据权利要求1所述的N-body问题近程作用计算在GPU结构的优化映射方法，其特征在于，所述每个block中的前n个线程处理其对应盒子中的n个粒子，每个block中的其余线程闲置。

3.根据权利要求1或2所述的N-body问题近程作用计算在GPU结构的优化映射方法，其特征在于，所述每个block中的线程数最大为512。