OpenCL-提高内存大小使用率



我现在正在用GPU做一个项目,这个项目的结果比CPU慢。原因是我将一个太小的数据数组作为输入(长度=1024)排队

我想将更多的数据排入队列,但由于内存使用问题,我被卡住了。我在内核中计算283个函数,每个函数的求值周期为481。

因此,为了返回我的结果,我必须创建一个数组,其大小为N(此处为1024)x481 x 283倍(因为283个函数返回双倍值)

这个长度太大了。由于我想放更多的数据,输出将大481 x 283,我将达到GPU的内存限制。我不知道如何使用更少的内存。

这是我的内核函数的一个例子:

PERIOD=481
data=input at get_global_id(0)
OUTPUT(get_global-id(0),t,x)是存储结果(三维数组)的访问器

for(int t=0;t<PERIODS;t++)
OUTPUT(get_global_id(0),t,1)=function1(data,t)
for(int t=0;t<PERIODS;t++)
OUTPUT(get_global_id(0),t,2)=function2(data,t)
for(int t=0;t<PERIODS;t++)
OUTPUT(get_global_id(0),t,3)=function3(data,t)

当然它看起来很糟糕,但问题是我的"被调用"函数有时需要T=12或T=24的值。因此,我必须计算每个函数的所有句点,以确保它们所需的值存在于OUTPUT访问器中。

例如:在二维问题(数据,周期)中,函数2需要函数1在T=4时的结果。但是工作项并不是全部同步的。所以,函数2所需要的值可能存在,也可能不存在。解决方案是确保通过在所有调用的函数周围设置循环,并从2D问题转到1D(这看起来真的很糟糕,2D组织本可以很好……但我没有找到任何方法通过全局内存同步所有线程)

我必须使用更少内存的第一个想法是调用481次内核函数,参数为T=T。因此,输出数组的重量将比现在少481,我可以多放481个数据。但要使用这个解决方案,我必须分解for循环,我想这是不可能的。(就像我说的:因为函数2可能需要函数1在T=4时的结果)

如果你有任何想法或解决方案,我很乐意听到。

我现在正在用GPU做一个项目,这个项目的结果较慢而不是CPU。原因是我将一个太小的数据数组作为输入(长度=1024)

假设您的gpu具有异步dma模块,为什么不将其实现为流水线?上传第一次迭代,在上传第二次的同时计算第一次,在计算第二次时下载第一次,同时上传第三次,下载2+计算3+上传4,下载3+计算4+上传5。这至少应该隐藏对gpu的读取或写入。

我想将更多数据排入队列,但由于内存问题,我被卡住了用法我在内核中计算283个函数,每个函数都在481个周期。

因此,为了得到我的结果,我必须创建一个大小为N(此处为1024)x 481 x 283的双(因为283函数返回双倍值)

假设"函数"的递归性不超过GB,则可以将类型的每个请求排入队列

OUTPUT[thread_num], period_num, f_num) = function_f_num(data, period_num);

当你只需要它(而不是计算所有东西)时,把它放进一个列表。然后,当队列达到大约1M个元素时,将其上传到gpu。添加伪递归性(使用带有后期修复的多重复制内核名称或支持内核的一些半堆栈结构)来查找任何需要递归性的元素。代替递归性,可以将不可计算的元素添加到新列表中,在主机端具有递归性来完成所有子列表。这也应该为软件提供明确的多gpu使用优势,因为您可以在不同的gpu中计算多个队列。或者,您可以简单地通过gpu使用主机指针来使用所有主机端缓冲区。查看CL_MEM_USE_HOST_PTR参数。

最新更新