OpenCL - 改善内存大小使用答案

【问题标题】：OpenCL - improve memory size usageOpenCL - 改善内存大小使用
【发布时间】：2016-10-08 20:50:16
【问题描述】：

我现在正在处理一个使用 GPU 的项目，该项目的结果比 CPU 慢。原因是我将一个太小的数据数组作为输入排队（长度 = 1024）

我想将更多数据排入队列，但由于内存使用情况而卡住了。我在我的内核中计算 283 个函数，每个函数都在 481 个周期上进行评估。

因此，为了取回我的结果，我必须创建一个大小为 N(此处为 1024) x 481 x 283 的双精度数组（因为 283 个函数返回双精度值）

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

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

周期 = 481
数据 = get_global_id(0) 处的输入
OUTPUT(get_global_id(0), t, x) 是存储结果的存取器（三维数组）

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

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

例如：在 2D 问题中（数据，PERIODS）——函数 2 需要函数 1 在 T=4 时的结果。但工作项并非全部同步。所以也许价值函数2需要存在，也许不存在。解决方案是确保通过在所有调用的函数周围放置 for 循环并从 2D 问题转到 1D （看起来真的很糟糕，2D 组织本来可以很棒..但我没有找到任何方法通过全局同步所有线程记忆）

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

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

【问题讨论】：

对于您所描述的问题，您的问题似乎不太适合 SIMD。你有一堆函数，里面有多个 if（否则你不需要不同的数据）。执行将非常缓慢。您必须将您的问题转换为可以在所有线程中同时运行而无需分支的问题。
感谢您的回复。分支是我读过的关于 GPU 使用的最糟糕的事情。我真的不能阻止他们..你可能是对的但我在写作的那一刻有一个想法。如果我将所有 283 个函数转换为 283 个不同的内核会怎样？从主机我可以做：入队一个内核，读回数据，入队下一个内核，读回数据为了构造我的三维输出，但在主机端有了这个想法，我可以再注入 283 个输入，因为输出重 283少。
283 个不同的内核，一个接一个地运行大小为 1 并具有 CPU 依赖性的内核，情况会更糟。您需要更改算法。
尺寸不是 1。但在我介绍的示例中，尺寸 N=1024。现在，如果我运行一个调用 283 个函数的内核，我需要一个大小为 N*481*283 的输出数组。所以我们接近 1GB 的 GPU 内存使用量。但是如果我做 283 个内核：输出数组的大小为 N*481。所以我们可以使用更多的内存，从而更多的输入数据，比如 N = 289792
这可能行得通，但尽量不要有任何 CPU 依赖性，因此您无需停止 GPU 执行等待 CPU 做出决定。一次将所有元素排队（283 个内核）。无论如何，只有 1024 大小，要做的工作量非常少......

标签： 2d opencl dimensions large-data

【解决方案1】：

我现在正在处理一个使用 GPU 的项目，结果速度较慢比CPU。原因是我将一个太小的数据数组排队为输入（长度=1024）

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

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

所以为了取回我的结果，我必须创建一个大小为 N(here 1024) x 481 x 283 of doubles（因为 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 参数。

【讨论】：

感谢您的回复。我正在编写具有大量函数调用的 opencl 中大小为 8000 行的代码。我的内核大约是 1.1 MB，有时我会因为这个太大而出现错误。这就是为什么我必须使用尽可能大的 OUTPUT 数组。现在function2需要function1的结果时，不需要调用function1，因为kernel size会太大，可以直接查看function2调用前计算的OUTPUT数组中的结果。
我使用 d_output = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR , Sizeof.cl_double * MODELPOINTS_LENGTH * (T_MAX+1) * NBR_FX_MAX, Pointer.to(h_output), null);但是当我将内核排入队列时，内存使用量将达到以下大小：N * T_MAX * NBR_FX * sizeof(double) 字节。为什么设备不使用指针但需要复制所有数据？就像你说的我希望设备使用指针而不是复制所有数据
也许驱动程序被窃听了？
在检查内存使用量增长的位置后：当我将设备指针设置为内核参数然后使用主机代码中的 clEnqueueNDRangeKernel 函数时会发生这种情况。如果我没有将此指针设置为内核参数，则内存不会增长，我将检查它是否是驱动程序问题。但要澄清：当使用 CL_MEM_USE_HOST_PTR 时，它不会在设备上分配内存？也许我对这个概念有误
您创建了一个缓冲区，但它不应该创建一个精确的副本，这将不利于它的使用。 “OpenCL 实现允许将 host_ptr 指向的缓冲区内容缓存在设备内存中。当内核在设备上执行时，可以使用此缓存副本。” khronos 说。