FPGA 上的 OpenCL 单个工作项 VS NDRange 内核答案

【问题标题】：OpenCL single work-item VS NDRange kernel on FPGAFPGA 上的 OpenCL 单个工作项 VS NDRange 内核
【发布时间】：2017-08-09 22:13:00
【问题描述】：

我是 OpenCL 新手，正在使用 FPGA 上的 OpenCL 进行分组密码加密。我读了一些论文，知道 Opencl 中有两种内核（单个工作项和 NDRange）。 NDRange 内核的函数不会被编译器自动流水线化，而单个工作项内核的函数会。

是否建议实现单个工作项内核而不是 NDRange 内核在 FPGA 上？为什么？

如果我想让内核循环运行直到读取所有数据，那么内核（一次从主机获取一些数据——在 FPGA 上运行——写回）。管道如何实现？

【问题讨论】：

我建议您查看 FPGA 供应商的指导。他们最了解他们的编译器。我也在他们的网站上看过一些不错的视频。因此，您不仅会了解这个问题的答案，而且可能会在此过程中学到更多。

标签： opencl gpu fpga pipeline

【解决方案1】：

单个工作项内核允许您将计算循环移动到内核中，您可以生成自定义管道，对累积进行巧妙优化，并通过“pragma”控制访问模式。 NDRange 内核依赖于您在工作项之间划分数据，编译器生成 SIMD 类型的硬件，每个单元都由您的内核描述。如果您的问题具有常规的数据并行性，从而使分区变得容易，那就太好了。 OpenCL 的 NDRange 内核专为 GPU 等 SIMD 计算单元而设计。您还可以利用“通道”在流应用程序中的单个工作项内核之间移动数据，从而减轻 DRAM 带宽。对于 NDRange 内核，您必须使用全局内存作为内核之间数据共享的媒介。

【讨论】：

【解决方案2】：

Shreedutt 的回答在 2016 年之前通常是可以接受的。英特尔的管道和通道实现远远超出了这一点：

如果您没有 work-item-variant 或对工作项 ID 的依赖，用户可以让多个内核或多个工作项访问管道。
可以控制和保证从多个工作项（例如 NDRange 内核）访问管道的确定顺序：“当管道存在于具有多个工作项的循环体中时，如下所示，每个循环迭代在后续迭代之前执行。这意味着工作组中每个工作项的循环迭代 0 在工作组中每个工作项的迭代 1 之前执行，依此类推。”
```
__kernel void ordering (__global int * data, write_only pipe int __attribute__((blocking)) req)
{
write_pipe (req, &data[get_global_id(0)]);
}
```

Channels 扩展是一种可行的替代方案，可以在具有多个工作项的循环中执行：

__kernel void ordering (__global int * data, int X) {
  int n = 0;
  while (n < X)
  {    
    write_channel_intel (req, data[get_global_id(0)]);
    n++;
  }
}

在 UG-OCL002 | 第 5 节中可以找到一些限制和警告。 2018.05.23 用于渠道和管道。我建议阅读它并观看最新的培训块https://www.youtube.com/watch?v=_0RtAKeRl00。另一个巨大的警告是，大公司决定为 OpenCL 提供单独的代码语法，需要不同的 pragma，一个多一个少。
然而，我应该从这个 IWOCL 演示开始：https://www.iwocl.org/wp-content/uploads/iwocl2017-kapre-patel-opencl-pipes.pdf。原因在于这些是新的计算模型，通过正确构建并行应用程序可以获得巨大的性能提升。更重要的是学习如何移动而不是移动数据。查看有关移除转置的最新技巧 GPU：https://devblogs.nvidia.com/tensor-core-ai-performance-milestones/ 我们可以在 FPGA 中做更多这样的技巧，可以吗？

我把它留给感兴趣的读者和贡献者来衡量 XILINX OpenCL 管道。

恕我直言，这是自切片面包以来软件定义的 FPGA 最重要的话题，特别是如果我们要在 ML/AI GPU 与 FPGA 的比赛中获胜。我支持 FPGA 团队。

【讨论】：