C++ 虚拟数组实现，由用户系统组合显存支持
仅头文件的 C++ 工具，支持基本的数组式使用模式，并使用系统中多张显卡作为存储，带有 LRU 缓存。
引言
所有流行的现代操作系统都有一些分页技术，允许我们使用比 RAM 更大的数据。这对于继续测试/学习更大规模数据的算法至关重要，而无需购买额外的 RAM，并且可以在同一台机器上运行数千个应用程序而不会溢出。
但是，它们（尚未）提供使用系统中通过足够高吞吐量 PCIe 桥接连接的闲置显存的选项。这个仅头文件的 C++ 工具（`VirtualMultiArray`）支持基本的数组式使用模式，并由一个分页系统支持，该系统通过 OpenCL 使用多张显卡。
Github 页面：https://github.com/tugrul512bit/VirtualMultiArray/wiki/How-it-is-used
背景
这个工具的基本思想是减少因大缓冲区过度分配和过度使用而导致的交换死锁中浪费的 CPU 周期。浪费一些显存有时比在接下来的 10 分钟内浪费所有 CPU 周期试图 **Ctrl-C** 应用程序（并在暴风雨天气中保护硬盘免受永久损坏，SSD 免受磨损等）更好。
使用这个简单的数组工具，数组所有索引的访问都是无缝的，分页仅在需要时由访问数组的线程处理。它是线程安全的，即使在相同元素上也是如此。
性能取决于访问模式，但对于低内存廉价系统上的简单顺序多线程访问，它可以将
变成这样
即使系统从近乎交换死锁的状态中恢复，也需要几分钟才能从硬盘回滚这些页面。因此，实际时间损失远不止 10 分 33 秒，但当闲置的显存得到一些工作时，它既快速又完全响应。
Using the Code
一些环境准备
- 由于它使用 OpenCL 并在 Ubuntu 18.04 和 Nvidia 显卡下编写，其 OpenCL 库是“`lib64`”。对于其他操作系统和显卡供应商（如 Amd/Intel），这可能会有所不同。
- 仅用于多线程测试，使用了一些 OpenMP（可以使用任何其他并行任务库）。
- 使用 g++-7 (c++1y dialect) 和 g++-10 测试
- 对于一个简短的 *main.cpp* 源文件，编译和链接如下所示
g++ -std=c++1y -O3 -g3 -Wall -c -fmessage-length=0 -march=native -mavx -fopenmp -pthread
-fPIC -MMD -MP -MF"src/foo.d" -MT"src/foo.d" -o "src/foo.o" "../src/foo.cpp"
g++ -L/usr/local/cuda/lib64 -o "your_app_name" ./src/foo.o -lOpenCL -lpthread -lgomp
用法简单
GraphicsCardSupplyDepot depot;
const size_t n = 1000000000ull; // number of elements of array, uses only VRAM
const size_t pageSize=1000; // RAM usage
const int activePagesPerVirtualGpu = 5; // RAM usage for caching (per GPU)
VirtualMultiArray<Bar> data(n,depot.requestGpus(),pageSize, activePagesPerVirtualGpu);
data.set(100,Bar()); // or: data[100]=Bar();
Bar value = data.get(100); // or: Bar value = data[100];
- 创建一个 `GraphicsCardSupplyDepot`。
- 从中请求 `gpus`。
- 将它们与数组大小、页面大小和每个 `gpu` 的活动页面一起提供给 `VirtualArray
` 模板化虚拟数组。 - 使用 `get`/`set` 方法（或带有代理类的运算符 `[]` 重载，`T var=data[i]; data[j]=x;`）。
- 数组大小必须是页面大小的整数倍。
- 不要使用太多活动页面，它们会消耗 RAM。
#include "GraphicsCardSupplyDepot.h"
#include "VirtualMultiArray.h"
// for testing
#include <iostream>
#include <numeric>
#include <random>
// a big enough object with plain data
template<typename T>
class ArrayData
{
public:
ArrayData(){
std::random_device rd;
std::mt19937 rng(rd());
std::uniform_real_distribution<float> rnd(-1,1000);
for(int i=0;i<1000;i++)
{
data[i]=rnd(rng);
}
}
T calc(){ return std::accumulate(data,data+1000,0); }
private:
T data[1000];
};
int main(int argC, char ** argV)
{
// gets all gpus. use depot(true) for some info about graphics cards in system
GraphicsCardSupplyDepot depot;
// n needs to be integer multiple of pageSize !!!!
const size_t n = 800000;
const size_t pageSize=10;
const int maxActivePagesPerGpu = 5;
VirtualMultiArray<ArrayData<int>> data(n,depot.requestGpus(),pageSize,
maxActivePagesPerGpu,{2,4,4});
//std::vector<ArrayData<int>> data;
//data.resize(n);
#pragma omp parallel for
for(size_t i=0;i<n;i++)
{
data.set(i,ArrayData<int>());
//data[i]=ArrayData<int>();
}
#pragma omp parallel for
for(size_t i=0;i<n;i++)
{
if(data.get(i).calc()<0){ std::cout<<"error: doctor"<<std::endl;};
//if(data[i].calc()<0){ std::cout<<"error: doctor"<<std::endl;};
}
return 0;
}
在具有以下空闲统计信息的系统上运行
（使用 vector 10 分钟 vs 使用虚拟数组 11 秒）
{2,4,4} 参数用于选择每张显卡显存使用比例（数量和最大 CPU 线程）。测试系统有 3 张卡，第一张空闲内存较少，因此值为 2。其他卡完全用于程序，因此它们获得 4。此值也是每张物理显卡的虚拟显卡数量。2 表示有 2 个通道在 2 个虚拟 GPU 上工作，可供 2 个 CPU 线程同时访问数组元素。4 表示有 4 个线程正在访问元素，仅由该卡处理。{2,4,4,} 允许 10 个 CPU 线程同时访问数组。如果需要从其中移除一张卡，只需将值设为零：{0,4,4} 禁用第一张卡。如果第二张卡有 24GB 内存而其他卡只有 4GB，那么 {4,24,4} 应该可以工作。未在高端系统上测试。如果未给出参数，其默认值将所有卡分配 4 个虚拟卡，这相当于 {4,4,...,4}，可能需要大量 CPU 核心才能充分利用。
{2,4,4} 也意味着，每个通道有“`maxActivePagesPerGpu=5`”数量的活跃计算/保存在 RAM 中的页面。第一张卡有 2 个克隆，因此它有 10 个页面。其他卡各有 20 个。总共 50 个页面。在相同的示例代码中，每个页面有 10 个元素。这意味着 RAM 中总共有 500 个元素，用于 10 路线程。由于每个元素需要 4000 字节，RAM 只容纳 2MB 数据。更大的页面提供更好的顺序访问性能，在下面的测试机器上高达 3.3GB/s
在配备 FX8150(2.1GHz)、GT1030(2GB, pcie v2.0 4 通道)、2xK420(2GB, pcie v2.0 4 通道, 8 通道) 和 4GB 单通道 1333MHz DDR RAM 的 PC 上，性能测量如下
// cpu active
// testing method object size throughput page size threads n objects pages/gpu
// random 44 bytes 3.1 MB/s 128 objects 8 100k 4
// random 4kB 238.3 MB/s 1 object 8 100k 4
// serial per thread 4kB 496.4 MB/s 1 object 8 100k 4
// serial per thread 4kB 2467.0MB/s 32 objects 8 100k 4
// serial per thread 44 bytes 142.9 MB/s 32 objects 8 100k 4
// serial per thread 44 bytes 162.3 MB/s 32 objects 8 1M 4
// serial per thread 44 bytes 287.0 MB/s 1k objects 8 1M 4
// serial per thread 44 bytes 140.8 MB/s 10k objects 8 10M 4
// serial per thread 44 bytes 427.1 MB/s 10k objects 8 10M 100
// serial per thread 44 bytes 299.9 MB/s 10k objects 8 100M 100
// serial per thread 44 bytes 280.5 MB/s 10k objects 8 100M 50
// serial per thread 44 bytes 249.1 MB/s 10k objects 8 100M 25
// serial per thread 44 bytes 70.8 MB/s 100kobjects 8 100M 8
// serial per thread 44 bytes 251.1 MB/s 1k objects 8 100M 1000
// interleaved thread44 bytes 236.1 MB/s 1k objects 8 100M 1000
// interleaved thread44 bytes 139.5 MB/s 32 objects 8 100M 1000
// interleaved thread44 bytes 153.6 MB/s 32 objects 8 100M 100
// interleaved thread4kB 2474.0MB/s 32 objects 8 1M 5
它处理 4.5GB 数据，仅消耗 100MB-1GB RAM（取决于总活动页面）。
另一个实际执行某些操作的示例（Mandelbrot 集生成代码，来自 此处）
#include "GraphicsCardSupplyDepot.h"
#include "VirtualMultiArray.h"
// testing
#include <iostream>
#include <fstream>
#include <complex>
class Color
{
public:
Color(){}
Color(int v){val=v;}
int val;
};
const int width = 1024;
const int height = 1024;
int value ( int x, int y) {std::complex<float> point((float)x/width-1.5,
(float)y/height-0.5);std::complex<float> z(0, 0);
unsigned int nb_iter = 0;
while (abs (z) < 2 && nb_iter <= 34) {
z = z * z + point;
nb_iter++;
}
if (nb_iter < 34) return (255*nb_iter)/33;
else return 0;
}
int main(int argC, char ** argV)
{
GraphicsCardSupplyDepot depot;
// n needs to be integer multiple of pageSize !!!!
const size_t n = width*height;
const size_t pageSize=width;
const int maxActivePagesPerGpu = 10;
VirtualMultiArray<Color> img(n,depot.requestGpus(),pageSize,maxActivePagesPerGpu);
std::cout<<"Starting cpu-compute on gpu-array"<<std::endl;
#pragma omp parallel for
for (int i = 0; i < width; i++) {
for (int j = 0; j < height; j++) {
img.set(j+i*width,Color(value(i, j)));
}
}
std::cout<<"mandelbrot compute finished"<<std::endl;
std::ofstream my_Image ("mandelbrot.ppm");
if (my_Image.is_open ()) {
my_Image << "P3\n" << width << " " << height << " 255\n";
for (int i = 0; i < width; i++) {
for (int j = 0; j < height; j++) {
int val = img.get(j+i*width).val;
my_Image << val<< ' ' << 0 << ' ' << 0 << "\n";
}
}
my_Image.close();
}
else std::cout << "Could not open the file";
std::cout<<"File saved"<<std::endl;
return 0;
}
验证
在 setter 方法中使用 openmp 并行化的 `for` 循环生成，通过 `get` 方法进行单线程串行读取。
重叠 I/O 的好处
CPU 缓存只有几个周期的开销，这只有纳秒级。RAM 有数十纳秒的延迟。即使只从显卡请求 1 字节数据，PCIe 延迟也介于 1 微秒到 20 微秒之间。因此，在等待数据时不做任何事情是严重的 CPU 周期损失。这与使用此工具的好处（不因等待 HDD 交换死锁而损失周期）相悖。幸运的是，CPU 可以运行的线程数超过其逻辑核心数。为此，一些线程在进入未知持续时间的等待状态（例如等待微秒级的 PCIe 数据到来，对于 CPU 来说微秒是永恒的）时，需要 `yield()` 当前工作。
默认的 OpenCL 等待方式（未定义任何空闲/忙等待策略）
clFinish(commandQueue);
这很简单，但也可能对 CPU 非常不友好，特别是当 GPU 供应商选择忙等待策略作为 OpenCL 实现的默认策略时。此时，CPU 线程无法切换到其他地方，必须等待 `clFinish()` 返回。大量 CPU 周期被浪费，而没有使用任何浮点流水线、整数流水线或 L1 缓存。
一种解决方案，使其在 8 核/8 线程 CPU 上可扩展到 64 个线程（为简单起见，删除了所有错误处理）
// this is the OpenCL data transmission command to be idle-waited
cl_event evt;
clEnqueueReadBuffer(q->getQueue(),...,.., &evt);
clFlush(q->getQueue());
const cl_event_info evtInf = CL_EVENT_COMMAND_EXECUTION_STATUS;
cl_int evtStatus0 = 0;
clGetEventInfo(evt, evtInf,sizeof(cl_int), &evtStatus0, nullptr);
while (evtStatus0 != CL_COMPLETE)
{
clGetEventInfo(evt, evtInf,sizeof(cl_int), &evtStatus0, nullptr);
std::this_thread::yield();
}
`clFlush()` 部分只是向 OpenCL 驱动程序发出一个开始在命令队列中异步发出命令的信号。因此，此命令会立即返回，驱动程序开始向显卡发送命令。
事件对象数据由 `clEnqueueReadBuffer()`（它将数据从显存读取到内存）的最后一个参数填充，事件跟踪开始。
`clGetEventInfo()` 是查询命令，用于了解事件是否已达到所需状态。当事件状态变为 `CL_COMPLETE` 时，命令完成，数据已准备好在目标缓冲区中。
在事件变为 `CL_COMPLETE` 之前，`while 循环` 会不断将工作让给其他线程，以执行一些有用的工作。这可能是一个 nbody 算法的计算循环，或者 CPU L1 缓存内部的一些数据复制，这些都独立于 CPU 外部进行的 DMA 工作，甚至是另一个并行启动 DMA 操作的 I/O 线程，或者隐藏另一个数组访问的 LRU 缓存的延迟（LRU 实现信息如下）。
这是 FX8150（8 线程/8 核）上，对 120KB 到 4.6GB 范围内数据集的 64 线程随机访问基准测试
当随机访问的数据集大小大于 LRU 容量时，I/O 会与其他 I/O 重叠（包括多张显卡之间和同一张卡内部）。
当在 L3 大小的数据集中进行随机访问时，任何 I/O 都与 `get()` 方法到 `Object` 变量的对象复制过程中的 L1 数据传输重叠。测试 CPU 具有 8 MB L3 容量，因此在图中，大约 8MB 数据集大小处性能提升很快。当数据集小于 LRU 并且 LRU 几乎缓存了所有内容时，I/O 的可能性非常非常低。但在其填满之前，它会继续从显卡的 VRAM 中获取数据。由于重叠 I/O，它比 8 线程版本更快，并且比单线程版本快得多。
此图的一个问题是性能在左侧再次下降。这是由页面锁定争用引起的，因为数组中存在的对象越少，线程争夺同一锁的竞争就越激烈。该基准测试中的每个对象长 60kB，因此整个 4.6GB 数据数组仅由大约 75000 个对象组成。
最近最少使用 (LRU) 缓存实现
根据用户所需的缓存行数，遵循以下三种不同的数据路径之一
**1 个缓存行** (`numActivePage = 1`) 每虚拟 GPU：按需直接淘汰。每当所需的虚拟索引未缓存在活动页面（缓存行）中时，它都会从 VRAM 中获取（如果旧数据已编辑，则在发送回之前）。
fImplementation=[&](const size_t & ind)
{
if(directCache->getTargetGpuPage()!=ind) // if virtual index not found
{
updatePage(directCache, ind); // update the cache line
directCache->reset(); // uncheck the "edited" boolean
}
return directCache;
};
然后，`get`/`set` 方法简单地使用返回的页面指针对其进行任何读/写操作（如果进行编辑，则设置编辑状态）。
**2-12 个缓存行** (`numActivePage = 2..12`) 每虚拟 GPU：这是“优先级队列”到代码的直接且快速映射，以缓存友好的方式移动数据。
Page<T> * const accessFast(const size_t & index)
{
Page<T> * result=nullptr;
auto it = std::find_if(usage.begin(),usage.end(),
[index](const CacheNode<T>& n){ return n.index == index; });
if(it == usage.end())
{
if(usage[0].page->getTargetGpuPage()!=index)
{
updatePage(usage[0].page, index);
usage[0].page->reset();
}
usage[0].index=index;
usage[0].used=ctr++;
result = usage[0].page;
}
else
{
it->used=ctr++;
result = it->page;
}
insertionSort(usage.data(),usage.size());
return result;
}
首先，它检查所需的元素（虚拟索引）是否在 `std::vector` 中。如果未找到，则简单地用所需的虚拟索引替换第一个元素，以便将来访问时将其缓存。在这里，使用一个计数器直接赋予它最高的访问频率，然后再淘汰页面内容。这部分可以针对其他类型的淘汰策略进行调整。（例如，如果需要 LFU（最不常用）淘汰，则可以删除计数器，只需使用 `usage[0].index = usage.[usage.size()/2].index + 1;` 将其发送到向量中间。当它到达向量中间时，左侧元素继续生存，以避免下次成为受害者，右侧元素变得特权并更容易留在缓存中。更像是 LFU 的“分段”版本。）
如果已在缓存中找到，则将其计数器值更新为向量中的最大值。
最后，插入排序执行必要的排序（此版本中右侧元素的移位）。
**13 个或更多缓存行** (`numActivePage = 13+`) 每虚拟 GPU：此版本使用 `std::unordered_map` 和 `std::list` 来实现对更多缓存行（例如 500 个）的必要扩展，而不会减慢速度。
Page<T> * const accessScalable(const size_t & index)
{
Page<T> * result=nullptr;
typename std::unordered_map<size_t,
typename std::list<Page<T>*>::iterator>::iterator it =
scalableMapping.find(index);
if(it == scalableMapping.end())
{
// not found in cache
Page<T> * old = scalableCounts.back();
size_t oldIndex = old->getTargetGpuPage();
if(old->getTargetGpuPage()!=index)
{
updatePage(old, index);
old->reset();
}
scalableCounts.pop_back();
scalableMapping.erase(oldIndex);
// add a new
scalableCounts.push_front(old);
typename std::list<Page<T>*>::iterator iter = scalableCounts.begin();
scalableMapping[index]=iter;
result = old;
}
else
{
// found in cache
// remove
Page<T> * old = *(it->second);
scalableCounts.erase(it->second);
// add a new
scalableCounts.push_front(old);
auto iter = scalableCounts.begin();
scalableMapping[index]=iter;
result = old;
}
return result;
}
基本上相同的算法，只是对于更多的缓存行具有更高的效率。无序映射（unordered map）保存从虚拟索引（VRAM 缓冲区）到活动页面（物理页面/RAM 缓冲区）的必要映射，但缓冲区指针保存在列表中。因为在列表中，从尾部删除节点并将其添加到头部很快（O(1)）。无序映射也很快（O(1)）地进行映射，这得益于其高效（哈希表？）的实现。直接查找元素比线性搜索向量更好。一旦未找到，就淘汰“最近最少使用”的页面（缓存行）。如果已找到，则不淘汰，只需访问 RAM（或 L1/L2/L3）中的缓冲区。
将列表的尾部节点取回并移动到头部，完成了 LRU 的循环，这样下次当有人在缓存中找不到东西时，下一个受害者就在新的尾部节点处准备就绪，而最新的头部节点在很长一段时间内都是安全的，不会被淘汰（或者如果用户的访问模式中发生缓存抖动，则时间很短）。
一旦选择了实现并将其分配给 `fImplementation` 函数成员，它就会在请求页面时被调用
// same access for both reading and writing
Page<T> * const access(const size_t & index)
{
return fImplementation(index);
}
根据 *VirtualArray.h* 计算冻结（VRAM）页面索引后
T get(const size_t & index)
{
const size_t selectedPage = index/szp;
Page<T> * sel = pageCache->access(selectedPage);
return sel->get(index - selectedPage * szp);
}
用户从 *VirtualMultiArray.h* 调用 `get()` 方法后
// get data at index
// index: minimum value=0, maximum value=size-1
// but not checked for overflowing/underflowing
T get(const size_t & index) const{
const size_t selectedPage = index/pageSize;
const size_t numInterleave = selectedPage/numDevice;
const size_t selectedVirtualArray = selectedPage%numDevice;
const size_t selectedElement = numInterleave*pageSize + (index%pageSize);
std::unique_lock<std::mutex> lock(pageLock.get()[selectedVirtualArray].m);
return va.get()[selectedVirtualArray].get(selectedElement);
}
它对 LRU 进行交错访问。每 K 个页面属于另一个虚拟显卡，每个虚拟显卡都缓存。因此，如果访问页面 1,5,9,13,17,21, ...，如果虚拟显卡设置是这样的：`std::vector
对于一个简单的图像处理基准测试（在 Mandelbrot 图像上计算高斯模糊），在添加 LRU（结合瓦片处理）后，吞吐量增加了 300%。类似地，重叠的顺序读写受益于 LRU，甚至在一定程度上，随机访问模式也获得了性能提升。
特点
批量读写
这些方法大大减少了每个元素的平均页面锁定次数，并实现了比基本 `get`/`set` 方法高得多的带宽（有时高达 50 倍）。
写入数组
std::vector<Obj> vec;
vec.push_back(Obj(1));
vec.push_back(Obj(1));
// write 2 elements to vArray starting from index i
vArray.writeOnlySetN(i,vec);
从数组中读取
// read 2 elements from vArray
std::vector<Obj> result = vArray.readOnlyGetN(i,2)
映射
此方法接受用户函数，并为其提供原始缓冲区指针，作为涉及 SIMD/向量指令、更快的数据复制以及任何需要对齐缓冲区的批量操作的更好替代方案。当页面足够大时，它比批量读/写方法略有优势。
简而言之，它选择性地将整个区域读入缓冲区，运行函数，然后选择性地将数据写回虚拟数组。
// T = int
// map region [303,803]
arr.mappedReadWriteAccess(303,501,[](int * buf){
// use same index for access
for(int i=303;i<303+501;i++)
buf[i]=i;
},pinned=false,read=true,write=true,userPtr=nullptr);
如果 `pinned` 参数为 `true`，则使用 Linux 的 `mlock`/`munlock` 命令来阻止操作系统将缓冲区分页出去。
如果 `read` 参数为 `true`，则在运行用户函数之前，缓冲区将填充来自虚拟数组的最新数据位。
如果 `write` 参数为 `true`，则在运行用户函数后，缓冲区内容将写回虚拟数组。
如果 `userPtr` 参数不是 `nullptr`，则跳过内部分配，并将 `userPtr` 用于用户函数+读写副本。用户函数采用 `userPtr` 的负偏移版本，因此函数内部的算法使用与虚拟数组相同的索引基数。如果 `index=100` 且 `range=100`，则在函数内部，用户可以访问 `buf` 为 `buf[100]...buf[199]`。
由于映射内的临时缓冲区比 `std::vector` 对齐得更好，并且一次性分配（不像批量读取的 K 次分配，其中 K=锁定的页面数），它倾向于更快完成。如果用户提供专用指针，则完全避免了分配，从而在运行用户函数之前减少了延迟。
未缓存（流式）读写操作
对于某些算法，任何页面大小（缓存行大小）或任何活动页面数（缓存行数）都不会带来性能提升。未缓存版本的 `get`/`set` 方法（目前仅适用于标量版本）可降低访问延迟。它与 `get`/`set` 相同，但额外地用 `streamStart`/`Stop` 命令装饰任何未缓存的访问块。
arr.streamStart(); // flushes any edited data to vram
arr.setUncached(..); // just like set but uncached, 2x lower latency
arr.getUncached(..); // just like get but uncached, ~10% lower latency
arr.streamStop(); // brings updated data from vram into LRU cache
如果在作用域结束之前没有缓存访问，则不需要这些额外的 `start`/`stop` 命令。由于缓存写入需要逐出并带回新数据，其成本是缓存读取操作的两倍。但未缓存写入只有单向数据移动，因此在延迟方面速度是两倍。
批量读写和映射的线程安全性
批量和映射操作在每个页面内都是线程安全的（页面始终是线程安全的），但当其区域跨越多个页面时，用户需要通过显式同步来确保数据一致性。多个页面不会同时锁定。一次只锁定和处理 1 个页面。因此，用户需要停止任何重叠/触及同一区域的并发访问。一个大型批量/映射操作可能存在多个读/写访问，这些访问在完成之前可能与其他线程的读/写发生冲突。
即使在重叠区域，并发使用只读版本（映射/批量读取/获取）也是安全的。数据始终保持完整，除非在同一区域使用写入方法。这可能会或可能不会在从另一个线程读取数据之前使上传到显卡的数据失效。
在非重叠区域进行并发批量写入/写入/映射写入是安全的。
在非重叠区域进行任何（批量或非批量）写入/任何（批量或非批量）读取操作是安全的。
为什么选择批量读/写？
一些算法通过一次加载更多数据（例如“平铺”）或显式缓存一个区域来优化速度。然后，数据在更快的内存上进行更有效的计算。例如，在 nbody 算法中，RAM 中的数据可以被平铺成更小的“缓存”大小的块，以比始终从 RAM 流式传输数据更快地完成力计算组。有时，块可以小到几个 CPU 寄存器（寄存器平铺）。同样的规则也适用于虚拟内存。“分页”这个虚拟数组只是一种隐式缓存。然后，用户可以通过使用自定义大小的“瓦片”并从“页面”加载到“瓦片”来进一步改进它，这些瓦片既具有 C++ 编译器 SIMD/矢量化的适当对齐，又具有“每页只锁定一次而不是每个元素锁定一次”的较低延迟。
为什么要映射？
映射提供了一个 `raw pointer`，它没有用于数组订阅 `operator []` 重载的 `SetterGetter proxy class` 的可怕副作用。
映射给出的原始指针也对齐到 `4096`，这有助于提高设备（如显卡）和 RAM 之间的数据复制性能，并运行 SIMD 指令的对齐版本。
为了进一步优化，可以选择允许固定缓冲区和用户定义的指针，以避免在每次映射时都进行分配。
GPU 加速的 find()
为了在常数时间内找到虚拟数组中的一个或多个元素，
// obj: object instance
// member: member of object to be used for searching value
// indexListMaxSize: maximum number of indices found
template<typename S>
std::vector<size_t> find(T & obj, S & member, const int indexListMaxSize=1);
被使用。第一个参数是一个对象实例，其中包含一个搜索值作为成员。第二个参数是要搜索的成员。此成员必须是“`obj`”对象的直接成员。这样，此方法可以计算其在对象中的字节偏移量，并在所有 GPU 的“find”内核中使用它来搜索所有对象中的成员值。最后一个参数是从每个 OpenCL 数据通道返回的最大索引数。如果给出 1000，并且总共有 10 个通道，则此方法最多可以返回 10000 个索引。这些索引处的所有元素都具有相同的成员值。用法很简单
纯数据
char searchObj = 'X';
std::vector<size_t> found = charArray.find(searchObj,searchObj);
这将搜索每个 OpenCL 通道中最多 1 个包含“`X`”值的元素。
对象
MyObject obj;
obj.id=450;
std::vector<size_t> found = charArray.find(obj,obj.id,100);
这将搜索所有包含 `id=450` 的对象，每个 OpenCL 通道最多返回 100 个结果。
在多次调用 `find()` 之间，向量中结果（索引）的顺序不保证相同。
由于它使用 GPU，搜索性能依赖于显存+GPU。对于配备单个 GTX 1070 的系统，搜索 1GB 数据中约 200 个元素最多需要 45 毫秒，而 GT1030 + 2x K420 系统需要 250 毫秒才能找到包含所需成员值的相同元素。
如果待查找的元素成员值在数组末尾，则此操作完成得**快得多**，比 `std::find` 快得多。如果元素成员在数组开头找到，则 `std::find` 速度更快。但是，在搜索过程中，不使用 RAM 带宽，只需要 VRAM 带宽。只有结果索引数组从 VRAM 复制到 RAM，与 `**std::find**` 的带宽要求相比可以忽略不计。如果它是一个可以装入 L1 的小数组，那么无论如何都无需使用虚拟数组。这个虚拟数组适用于不适合 RAM 的大数据。
由于搜索时间是准恒定的，并且它返回一个元素索引数组，所以它有点像使用带有读写多对一关系增强功能的 `std::map`。
目前，这是唯一一个 GPU 加速的方法。也许将来可以添加类似 `std::transform` 的方法。但如果用户的项目已经有 OpenCL/CUDA 加速的计算库，可能根本不需要。
在调用方法运行“`find`”内核之前，RAM 上的所有活动页面都会刷新到显卡的 VRAM 中，以便 `find()` 实际返回更新的结果。不建议在 `find()` 调用期间对虚拟数组执行并发 `write` 操作。并发读取是可以的。
一个矢量化和瓦片化 N-body 计算基准测试，使用虚拟数组
N体算法用于计算质量在由质量产生的引力场下的运动。这个例子重点关注其最瓶颈的部分，即N个质量的力计算。
由于 CPU 每个核心拥有不止一个浮点流水线，即 SIMD，因此在像 N 体算法这种计算密集型算法中不使用所有这些流水线将是一种浪费。为了方便实现这一点，使用了 mipp 库（这是一个仅包含头文件的库，可以使用 SSE/AVX/AVX512）。开发人员需要做的就是使用它的向量类来以更大的块进行计算，例如一次 4 个，一次 8 个，16 个，等等。
矢量化只是优化的一部分。另一部分是分块。分块利用对内存较小区域的重复访问来更好地利用缓存。
- L1 缓存分块（假设只有 4000 个 x,y,z 坐标可以放入 CPU 的 L1 中）
- RAM 分块（一次性将虚拟数组的大块数据获取到 RAM 中，高效地处理它们）
缓存分块很简单。只需反复使用一小组变量（某些粒子的坐标）就足以利用缓存，并且许多 CPU 至少有数十 KB 的 L1 缓存，足以容纳数千个质量坐标数据（x=4 字节，y=4 字节，z=4 字节：每个粒子 12 字节）。为此，准备一个粒子块，并对粒子数组中的每个粒子进行重新扫描。实际上这里没有粒子，只有数组。用于 3D 位置数据的 X 数组、Y 数组、Z 数组，以及用于 3D 速度数据的 VX 数组、VY 数组、VZ 数组。
RAM 分块只是为了优化虚拟数组的 x,y,z 读取效率。一次读取 4000 个粒子意味着页面锁定开销和 PCIe 延迟平均每个粒子分摊 4000 份。对虚拟数组的映射访问执行此操作，并提供原始指针来处理数据。由于映射访问还将临时缓冲区对齐到 4096，因此可以在其上使用 AVX/SSE 指令的任何对齐版本。对齐的 SSE/AVX 数据加载通常比未对齐的数据加载更快。
基准测试代码
#include "GraphicsCardSupplyDepot.h"
#include "VirtualMultiArray.h"
#include "PcieBandwidthBenchmarker.h"
#include "CpuBenchmarker.h"
// testing
#include <iostream>
#include "omp.h"
// a simd tool from github
#define MIPP_ALIGNED_LOADS
#include "mipp/mipp.h"
int main()
{
const int simd = mipp::N<float>();
std::cout<<"simd width: "<<simd<<std::endl;
const int n = simd * 40000;
std::cout<<"particles: "<< n <<std::endl;
GraphicsCardSupplyDepot depot;
VirtualMultiArray<float> xVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
VirtualMultiArray<float> yVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
VirtualMultiArray<float> zVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
VirtualMultiArray<float> vxVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
VirtualMultiArray<float> vyVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
VirtualMultiArray<float> vzVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
// if you don't initialize data,
// floating point NaN values slow down the speed too much
{
CpuBenchmarker init(0,"init");
#pragma omp parallel for
for(int i=0;i<n;i+=1000)
{
xVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
yVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
zVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
vxVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
vyVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
vzVir.mappedReadWriteAccess(i,1000,[i](float * ptr)
{ for(int j=i;j<i+1000;j++) ptr[j]=j; },false,false,true);
}
}
mipp::Reg<float> smoothing = 0.0001f;
// mapped array access
{
CpuBenchmarker bench(((size_t)n) * n * sizeof(float) * 3,
"mapped access for force-calc",((size_t)n)*n);
const int tileSize = 4000;
const int regTile = 500;
#pragma omp parallel for num_threads(32) schedule(dynamic)
for (int i = 0; i < n; i += regTile)
{
std::vector<float> x0 = xVir.readOnlyGetN(i,regTile);
std::vector<float> y0 = yVir.readOnlyGetN(i,regTile);
std::vector<float> z0 = zVir.readOnlyGetN(i,regTile);
mipp::Reg<float> fma1[regTile];
mipp::Reg<float> fma2[regTile];
mipp::Reg<float> fma3[regTile];
for(int j=0;j<regTile;j++)
{
fma1[j]=0.0f;
fma2[j]=0.0f;
fma3[j]=0.0f;
}
for (int ii = 0; ii < n; ii += tileSize)
{
xVir.mappedReadWriteAccess(ii, tileSize, [&,ii](float* ptrX1) {
const float* __restrict__ const ptrX = ptrX1 + ii;
yVir.mappedReadWriteAccess(ii, tileSize, [&,ii](float* ptrY1) {
const float* __restrict__ const ptrY = ptrY1 + ii;
zVir.mappedReadWriteAccess(ii, tileSize, [&, ii](float* ptrZ1) {
const float* __restrict__ const ptrZ = ptrZ1 + ii;
for (int ld = 0; ld < tileSize; ld += simd)
{
mipp::Reg<float> x = mipp::load(ptrX + ld);
mipp::Reg<float> y = mipp::load(ptrY + ld);
mipp::Reg<float> z = mipp::load(ptrZ + ld);
for(int reg0 = 0; reg0 < regTile; reg0++)
{
const int reg = reg0 ;
const mipp::Reg<float> x0r = x0[reg];
const mipp::Reg<float> y0r = y0[reg];
const mipp::Reg<float> z0r = z0[reg];
const mipp::Reg<float> dx = mipp::sub(x,x0r);
const mipp::Reg<float> dy = mipp::sub(y,y0r);
const mipp::Reg<float> dz = mipp::sub(z,z0r);
const mipp::Reg<float> dx2 = mipp::mul(dx,dx);
const mipp::Reg<float> dy2 = mipp::mul(dy,dy);
const mipp::Reg<float> dz2 = mipp::mul(dz,dz);
const mipp::Reg<float> dxy2 = mipp::add(dx2,dy2);
const mipp::Reg<float> dz2s =
mipp::add(dz2,smoothing);
const mipp::Reg<float> smoothed =
mipp::add(dxy2,dz2s);
const mipp::Reg<float> r = mipp::rsqrt(smoothed);
const mipp::Reg<float> r3 =
mipp::mul(mipp::mul(r, r), r);
fma1[reg] = mipp::fmadd(dx, r3, fma1[reg]);
fma2[reg] = mipp::fmadd(dy, r3, fma2[reg]);
fma3[reg] = mipp::fmadd(dz, r3, fma3[reg]);
}
}
}, false, true, false);
}, false, true, false);
}, false, true, false);
}
for(int j=0;j<regTile;j++)
{
vxVir.set(i+j, vxVir.get(i+j) + mipp::hadd(fma1[j]));
vyVir.set(i+j, vyVir.get(i+j) + mipp::hadd(fma2[j]));
vzVir.set(i+j, vzVir.get(i+j) + mipp::hadd(fma3[j]));
}
}
}
return 0;
}
首先，分配虚拟数组（这对于慢速计算机可能需要几秒钟）。
VirtualMultiArray<float> xVir(n, depot.requestGpus(), 4000, 1,{5,15,10});
这里，4000（元素）是页面大小，5、15、10 是并行 OpenCL 数据通道，1 是每个通道的页面缓存。这意味着 RAM 中总共有 (5+15+10)*1*4000 = 120000 个元素被缓存。这略高于所有粒子的 1/3。因此，如果所有粒子都同时使用，最多 1/3 的粒子具有 RAM 速度。对于所有 PCIe 桥接器具有相同数量的 PCIe 通道的系统，可以使用 {5,5,5} 或 {10,10,10} 或类似的任意乘数集来最大化 PCIe 使用效率。对于开发计算机，第二张卡具有更高的 PCIe 带宽，而第一张卡由于操作系统功能而存在一些延迟问题。为了自动化找到最佳带宽，可以使用 `PcieBandwidthBenchmarker().bestBandwidth(5 /* 或 10 */ )` 代替 {...}。
然后，将数组初始化为任意值，以抵消任何非规范化浮点运算减速效应（不初始化时，在开发机器上会慢 10 倍）。最后，以暴力方式计算力（O(N^2) 力计算）。SIMD 长度（在 fx8150 和 -mavx -march=native g++ 编译器选项下）为 8。这意味着它使用 CPU 的 AVX 指令一次计算 8 个数据。当 SIMD 长度为 8 时，需要模拟 320000 个粒子。在 SSE CPU 上，它将创建 160000 个粒子。
由于 I/O 操作的可扩展性，OpenMP pragma 指令中使用了 32 个线程。这适用于 PCIe v2.0 上的 3 张显卡、一个 8 线程 CPU 和 64 位构建。在其他系统上可能会有不同的性能增益。但是对虚拟数组的每次“映射访问”都意味着 VRAM-RAM 数据移动，这需要一些时间，并且由于 DMA 引擎而与 CPU 异步。因此，当一个线程在等待 VRAM 数据时被阻塞时，另一个线程可以切换到其位置并继续执行其他任务。这样，部分（或全部？）I/O 延迟被隐藏在新任务后面，整体性能提高。通常，NBody 算法在 8 核 CPU 上以 8 个线程运行得很好，因为它纯粹是数学运算。但是虚拟数组访问增加了 I/O。
相同的多线程可扩展性可以发生在任何其他 I/O 瓶颈算法上，而不仅仅是 N-body。特别是当系统拥有比 CPU 核心数更多的异步显卡 DMA 引擎时。例如，一些 Nvidia Tesla 卡每张卡有 5 个引擎。因此，即使是双核 CPU 也将在双 Tesla 系统上使用 10-20 个线程。
映射部分
xVir.mappedReadWriteAccess(ii, tileSize, [&,ii](float* ptrX1) {
const float* __restrict__ const ptrX = ptrX1 + ii;
...
},false,true,false);
本身不足以充分发挥 SIMD 指令的性能。编译器需要知道指针不是别名。这通过为指针添加 `__restrict__` 关键字来实现。然后编译器可以执行必要的优化，这可以导致数学代码的矢量化。
如果此 N-body 算法通过 Barnes-Hut 优化，则可以轻松处理数百万个粒子。如果是快速傅里叶变换方法，则可以处理数十亿个粒子。在这些更大规模的模拟中，将从虚拟数组中读取更大的数据块，这可能导致操作系统在计算过程中将其中一些分页进出。为了阻止这种情况，可以将第一个“`false`”参数切换为“`true`”，以阻止操作系统执行此操作。由于这个简单的示例不太可能感受到操作系统分页，因此不需要它，甚至可能导致性能下降，因为固定数组也需要一些时间。
矢量化部分很简单，感谢 `mipp` 库，它也适用于任何其他矢量化库
// for all elements in virtual array's RAM tile
for (int ld = 0; ld < tileSize; ld += simd)
{
// aligned data load from memory
// unaligned version: mipp::loadu(ptr) is slower
mipp::Reg<float> x = mipp::load(ptrX + ld);
mipp::Reg<float> y = mipp::load(ptrY + ld);
mipp::Reg<float> z = mipp::load(ptrZ + ld);
// for all elements in cache tile
for(int reg0 = 0; reg0 < regTile; reg0++)
{
// select a particle id in cache tile
const int reg = reg0 ;
// broadcast particle data to all lanes of AVX vector (8 element copy)
const mipp::Reg<float> x0r = x0[reg];
const mipp::Reg<float> y0r = y0[reg];
const mipp::Reg<float> z0r = z0[reg];
// find 3D relative position between cache tile particle and RAM tile particle
// in 8-element chunks
const mipp::Reg<float> dx = mipp::sub(x,x0r);
const mipp::Reg<float> dy = mipp::sub(y,y0r);
const mipp::Reg<float> dz = mipp::sub(z,z0r);
// compute squares of relative positions
// in 8-element chunks
const mipp::Reg<float> dx2 = mipp::mul(dx,dx);
const mipp::Reg<float> dy2 = mipp::mul(dy,dy);
const mipp::Reg<float> dz2 = mipp::mul(dz,dz);
// find sum of squared distances in all dimensions + add a smoothing
// in 8-element chunks
const mipp::Reg<float> dxy2 = mipp::add(dx2,dy2);
const mipp::Reg<float> dz2s = mipp::add(dz2,smoothing);
const mipp::Reg<float> smoothed = mipp::add(dxy2,dz2s);
// find r (distance) between 2 selected particles
// in 8-element chunks
const mipp::Reg<float> r = mipp::rsqrt(smoothed);
// we need 3rd power of it because (dx+dy+dz)/r^3 equals (unit vector) *(1/r^2)
// in 8-element chunks
const mipp::Reg<float> r3 = mipp::mul(mipp::mul(r, r), r);
// multiply with dx,dy,dz and add to force components fx,fy,fz
// no mass value used, all particles with same mass assumed
// in 8-element chunks
fma1[reg] = mipp::fmadd(dx, r3, fma1[reg]);
fma2[reg] = mipp::fmadd(dy, r3, fma2[reg]);
fma3[reg] = mipp::fmadd(dz, r3, fma3[reg]);
}
}
优化力存储（`fmadd` 部分）性能的诀窍是，不使用“`store`”命令。而是在 `fmadd`（最后的 3 条指令）操作中直接使用一个 CPU 寄存器向量。这使得编译器只在必要时进行加载和存储，并在可能的情况下直接使用寄存器。给予编译器一些自由有助于获得一些性能。对于 3.6GHz（无睿频）的 FX8150，这实现了 **90GFLOPS** 的计算性能。在 3.6GHz 时，CPU 的绝对峰值为 **230 GFLOPS**，但 Nbody 算法的“+”和“*”数学运算数量并不相等。因此，绝对可达到的**最大性能实际上是峰值的 70%**，即 161 GFLOPS。要获得 **161 GLFLOPS**，必须采用寄存器平铺方案，更频繁地使用寄存器而不是内存访问。为简单起见，这里仅使用缓存平铺，并且仍然实现了超过 50% 的可达峰值计算性能。在 Ryzen R9 3950x CPU 上，它应该接近 1TFLOPS。
当粒子的力寄存器累积了所有来自其他粒子的力分量后，水平加法操作将其中 8 个元素相加，并将其作为一个元素写入速度（虚拟）数组。
基本错误处理
目前，`VirtualMultiArray` 的构造函数会检查 `页面大小` - `数组长度` 以及 `GPU 数量` - `页面数量` 的不一致性，并抛出相关的错误消息。
try
{
// throws error when number of pages (n/pageSize) < 20+20+20
// throws error when n/pageSize division is not integer result
VirtualMultiArray<char> img(n,gpus,pageSize,maxActivePagesPerGpu,{20,20,20});
}
catch(std::exception & e)
{
std::cout<<e.what()<<std::endl;
}
此外，活动页面和显卡之间的数据复制（在 `get`/`set`/`map`/`read`/`write` 操作期间）也有“`throw`”，提示哪个部分出现故障。
关注点
**编辑**：LRU 缓存开销通过让所有 I/O 在空闲循环中等待而被优化掉了。现在，一个 8 逻辑核心 CPU 可以运行 64 个线程，其中 48 个在空闲等待 I/O，而 8 个在进行数学计算。
欢迎添加任何想法，例如将其打造成**内存数据库**、**GPU 加速科学工具**、**错误修复**等。
优化
由于用户可以使用“`memMult`”参数（`vector
如果物理卡不在同类型的 PCIe 桥接器上（例如，一张在 16x 插槽，一张在 8x 插槽等），`VirtualMultiArray` 的 `memMult` (`std::vector
如果页面没有足够的容量来弥补 PCIe 延迟，那么使用小对象效率不高。一个慢速的 PCIe 可能每秒只能进行 100k 次读/写操作（或约 10 微秒延迟），因此对象数据中的字节越多，效率越高。图像处理的一个潜在优化可能是将像素转换为 16x16 瓦片（256 像素，每个 3 个浮点数，每个瓦片对象 3KB 数据），以实现空间局部性和 PCIe 效率。
目前，所有活动页面都是固定缓冲区，无法被操作系统分页系统移动。在未来版本中，将提供一个选项来禁用此功能，因为固定内存是相当稀缺的资源，可能不适用于过大的页面。如果在 `std::cout` 上观察到“error: buffer read”或类似输出，则可能是 `pageSize` 过高，或页面过多，或对象太大无法被操作系统/OpenCL 固定（或者只是剩余内存不足）。
编辑：一些新功能
为了自动调整数据共享比例以获得最大带宽，有一个可选的基准测试类
#include "PcieBandwidthBenchmarker.h"
// by default, uses 128MB vram of each card to test their pcie-copy performance
PcieBandwidthBenchmarker bench;
// slowest pcie card receives 2 channels,
// others get higher "integer" number of channels
// that depends on their pcie performance
std::vector<int> ratios = bench.bestBandwidth(2);
// which can be adjusted for better benchmark results
PcieBandwidthBenchmarker bench(250);
// slowest card gets 10 channels,
// better chance of overlapping multiple data-copies in multithreaded element access
std::vector<int> ratios = bench.bestBandwidth(10);
// in development computer, this had best results (3500MB/s) for sizeof(Obj)=1MB
VirtualMultiArray<Obj> data1(..,..,..,..,bench.bestBandwidth(2));
为了自动调整数据共享比例以获得最大大小，`VirtualMultiArray` 的构造函数有一个新参数 (`mem`)
// array data shared between cards in tune with their vram size
// 2 GB card serves 1GB, if 24 GB card serves 12 GB
auto maximizeMem = VirtualMultiArray<Type>::MemMult::UseVramRatios;
// this is default value, just keeps what user gives in "memMult" parameter {1,2,3}
auto customMem = VirtualMultiArray<Type>::MemMult::UseDefault;
VirtualMultiArray<Obj> data1(..,..,..,..,bench.bestBandwidth(2),customMem);
VirtualMultiArray<Obj> data1(..,..,..,..,overridenParameter,maximizeMem);
如果 `mem` 参数被赋值为“`UseVramRatios`”，那么它会覆盖 `memMult` 参数中的比例，除了零值元素
// array data shared between cards in tune with their vram size
// 2 GB card serves 1GB, if 24 GB card serves 12 GB
auto maximizeMem = VirtualMultiArray<Type>::MemMult::UseVramRatios;
VirtualMultiArray<Obj> data1(..,..,..,..,{0,overridden,overridden},maximizeMem);
VirtualMultiArray<Obj> data2(..,..,..,..,{0,0,overridden}, maximizeMem);
这样，组合内存大小仍然最大化，但只使用非零比例的显卡。
基准测试方法同时使用设备到主机和主机到设备的复制来测量组合带宽，而不是仅仅测量单向性能。也许以后可以添加读优化或写优化的基准测试选项。目前，基准测试同时测量读写性能。
优化 `memMult` 参数以获得最大容量可能会使组合带宽低于等数据分布情况（当用户不使用 `VirtualMultiArray` 构造函数的最后两个参数时的默认情况）。例如，一张具有 24GB 显存的显卡将有效地通过“`MemMult::UseVramRatios`”限制所有其他 2GB 显卡的吞吐量。
基准测试类也不考虑“组合大小”，如果在非常不对称的 GPU 系统中可能会导致容量问题。
Virtual array=5GB
PCIE lanes bandwidth maximize bandwidth maximize capacity default
Gtx 1080ti: 1x 1GB/s 1GB/s 1GB/s 1.0GB/s
Gtx 1070 : 4x 2GB/s 2GB/s 0.7GB/s 1.0GB/s
Gt1030 : 4x 2GB/s overflow error 0.2GB/s 1.0GB/s
PCIE lanes vram maximize bandwidth maximize capacity default
Gtx 1080ti: 1x 11GB 550MB 2.6GB 1.7GB
Gtx 1070 : 4x 8 GB 2.2GB 1.9GB 1.7GB
Gt1030 : 4x 2 GB overflow error 0.5GB 1.7GB
目前，算法并未考虑这些点。用户可能需要手动更改基准测试结果中的比例，以避免缓冲区错误。特别是系统的主卡（向屏幕发送帧的卡）的操作系统会持续占用内存，可能高达数百 MB。该卡（可能在比例数组中是第一张卡）的乘数可以手动降低到 `{1,3,3}`，而不是直接使用基准测试输出的 `{2,3,3}`。
历史
- 2021年2月5日：目前，这是一个非常基础的版本，文档很少，但至少在随机访问性能上比慢速硬盘快至少 15 倍，在顺序读写方面击败了旧式 SSD。此外，当与足够大的对象结合时，它甚至能击败 NVMe。
- 2021年2月8日：`VirtualMultiArray` 类构造函数中添加了自动调整参数选项。它调整每张卡的内存消耗以最大化允许的数组大小（但不检查其他应用程序是否使用相同的显卡，因此主卡在接近最大值时可能会溢出）。还添加了一个基准测试类，用于微调数据共享比例以最大化带宽（而不是大小）。
- 2021年2月13日：增加了批量读写和映射操作，以大幅提高吞吐量
- 2021年3月6日：增加了 GPU 加速的 `find()` 方法和 nbody 算法示例
- 2021年3月15日：增加了 LRU 缓存，其开销如何通过重叠 I/O 进行平衡，以及基准测试结果