Opencl 启动并发内核答案

【问题标题】：Opencl launch concurrent kernelsOpencl 启动并发内核
【发布时间】：2021-09-24 18:16:57
【问题描述】：

据我了解，要执行并发内核（在我的情况下是相同的内核但不同的 I/O 数据），必须通过启动具有明显自己的工作组的独特计算单元（流式多处理器 -SM）来完成。

例如 gtx960m 有 5 个 SM（Opencl 中的计算单元）。以自己的 16x16 (2d) 工作组异步无序启动clEnqueueNDRangeKernel 5 次，会启动所有 5 个计算单元以同时执行它们吗？报告的本地内存为 64kb。这是针对所有计算单元，还是每个单元都有自己的 64kb？

【问题讨论】：

标签： concurrency gpu opencl nvidia

【解决方案1】：

每个 CU 有 4 个（麦克斯韦/帕斯卡）或 2 个（图灵/安培，AMD）扭曲。 Warp 是一组 32 个 CUDA 内核/硬件流处理器。

在一个 Warp 中运行的所有线程都必须执行完全相同的指令。在 Warp 中，甚至无法进行分支。一个 CU 中的两个 Warp 可以处理不同的分支，但不能同时处理不同的内核。

如果您在具有 5 个 CU 的 960m 上并行执行不同队列中的两个内核，例如，内核 1 可以有 3 个 CU，内核 2 可以有剩余的 2 个。但是不能拆分 CU 以同时运行多个内核时间。

在 OpenCL 中，您可以将工作组大小设置为 Warp 大小 (32) 的某个倍数。可以在一个 Maxwell CU 上同时执行 4 个（工作组大小为 32）、2 个（工作组大小为 64）或 1 个（工作组大小为 128 或更大）的 OpenCL 工作组。

本地内存量（在您的情况下为 64KB）是每个 CU。因此，如果您有一个大型工作组，例如 256 个线程，则每个线程的可用本地内存都比工作组大小为 64 的情况下要少，因为工作组中的所有线程在它们运行的一个 CU 上共享相同的本地内存。

【讨论】：

所有内核（同一个程序）都在同一个队列中，并以我们的顺序模式异步启动。每个工作组有 256 (16x16 - 2d) 个工作项或线程。所以我再次问，因为我不明白你的答案，他们会同时执行吗？关于本地内存，您的意思是更大的工作组（更多线程）需要更多内存？这是为什么？假设本地缓冲存储器是一个 12kb 的数组。如果工作组有 32 或 256 个线程，为什么这种情况会发生变化？
如果内核在同一个队列中，它们不能同时执行。为此，您需要 2 个队列。队列的异步操作只是意味着在主机端运行enqueue命令时clEnqueueNDRangeKernel命令没有阻塞，内核没有执行。见这里stackoverflow.com/q/58946050/9178992
本地内存是每个 CU 的，并且在工作组中的所有线程之间共享。如果你有一个更大的工作组，那么每个线程你有更少的本地内存和更少的可用寄存器空间（私有内存/临时变量）。 GPU 上的 __local 内存空间是 L2 缓存的一部分。如果您想在工作组内的线程之间传递数据，您只需要本地内存。在某些情况下，例如矩阵乘法中的缓存平铺，这可以显着提高速度。
我明白了。所以要同时运行 5 个内核，我需要 5 个单独的队列，对吗？就像 cuda 世界中的“流”？我在内核函数中的本地内存声明是__local unsigned int* tempArray，它是使用语句clSetKernelArg(kernelCalc, 1, 1000 * sizeof(integer), NULL) 从主机调用的。那么，如果我有 1 个线程或 100 个线程（工作项），那么本地内存的 1000 * 4 = 4000 字节的大小到底是如何变化的？我认为它保持不变，即在这种情况下约为 4KB
我想让你知道，Args 问题的根源是一些没有改变的 args 超出了在并发内核执行中调用内核的 for 循环。感谢您在并发级别提供的宝贵帮助。