oneAPI GPU 优化指南 - 在 kernel 中进行 I/O

本章节翻译by chenshusmail@163.com 原文：Doing I/O in the Kernel (intel.com)

打印语句是查看程序结果所需的最基本功能。在加速器上，打印语句出奇地困难， 而且在性能开销方面相当高昂。

SYCL* 提供了一些功能，以帮助使这项任务类似于标准的 I/O C/C++ 程序， 但是由于加速器的工作方式，您需要理解一些特性。从 SYCL* kernel 中无法进行文件 I/O。

SYCL* 提供了 stream 类，让您可以从 kernel 内部向控制台打印信息， 也提供了一种简单的方式来调试简单问题，而无需求助于调试器。 stream 类提供的功能非常类似于 C++ STL 的 ostream 类，其使用方法与 STL 类似。 下面我们描述如何使用 SYCL stream 类从 queue 的 kernel 中输出信息。

要使用该类，我们必须首先实例化它。stream 构造函数的签名如下：

stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH);

构造函数接受三个参数：

• BufferSize：在整个 kernel 范围内可能打印的字符总数

• MaxStatementSize：对 stream 类的任何一次调用中的最大字符数

• CGH：在 sycl::queue::submit 调用中对 sycl::handler 参数的引用

使用方法与 C++ STL 的 ostream std::cout 类非常相似。 需要打印的消息或数据通过适当的 operator<< 方法发送到 SYCL stream 实例。 SYCL 为所有内置数据类型（如 int、char 和 float）以及一些常见类 （如 sycl::nd_range 和 sycl::group）提供了实现。

以下是一个 SYCL stream 实例的使用示例：

void out1() {
  constexpr int N = 16;
  sycl::queue q;
  q.submit([&](auto &cgh) {
     sycl::stream str(8192, 1024, cgh);
     cgh.parallel_for(N, [=](sycl::item<1> it) {
       int id = it[0];
       /* Send the identifier to a stream to be printed on the console */
       str << "ID=" << id << sycl::endl;
     });
   }).wait();
} // end out1

使用 sycl::endl 类似于使用 C++ STL 的 std::endl ostream 引用 ——它用于插入新行以及刷新 stream。

编译并执行上述 kernel 将给出以下输出：

ID=0
ID=1
ID=2
ID=3
ID=4
ID=5
ID=6
ID=7
ID=8
ID=9
ID=10
ID=11
ID=12
ID=13
ID=14
ID=15

在选择适当的 BufferSize 和 MaxStatementSize 参数时必须小心。如果大小不足， 可能会导致语句不被打印，或者打印的信息比预期的少。请参考以下 kernel：

void out2() {
  sycl::queue q;
  q.submit([&](auto &cgh) {
     sycl::stream str(8192, 4, cgh);
     cgh.parallel_for(1, [=](sycl::item<1>) {
       str << "ABC" << sycl::endl;     // Print statement 1
       str << "ABCDEFG" << sycl::endl; // Print statement 2
     });
   }).wait();
} // end out2

编译并运行此 kernel 将给出以下输出：

ABC

第一条语句成功打印出来，因为要打印的字符数是 4（包括由 sycl::endl 引入的换行符）， 并且最大语句大小（由 sycl::stream 构造函数的 MaxStatementSize 参数指定）也是 4。 然而，只有第二个语句的换行符被打印出来。

以下 kernel 显示了增加允许的最大字符大小的影响：

void out3() {
  sycl::queue q;
  q.submit([&](auto &cgh) {
     sycl::stream str(8192, 10, cgh);
     cgh.parallel_for(1, [=](sycl::item<1>) {
       str << "ABC" << sycl::endl;     // Print statement 1
       str << "ABCDEFG" << sycl::endl; // Print statement 2
     });
   }).wait();
} // end out3

编译并运行上述 kernel 将给出预期的输出：

ABC
ABCDEFG

上述示例使用了具有单个 work-item 的简单 kernel。更真实的 kernel 通常会包含多个 work-item。 在这些情况下，不能保证打印到控制台的语句的特定顺序，您应该预料到来自不同 work-item 的语句被交错输出。 请参考以下 kernel：

void out4() {
  sycl::queue q;
  q.submit([&](auto &cgh) {
     sycl::stream str(8192, 1024, cgh);
     cgh.parallel_for(sycl::nd_range<1>(32, 4), [=](sycl::nd_item<1> it) {
       int id = it.get_global_id();
       str << "ID=" << id << sycl::endl;
     });
   }).wait();
} // end out4

一次运行可以产生以下输出。

ID=0
ID=1
ID=2
ID=3
ID=4
ID=5
[snip]
ID=26
ID=27
ID=28
ID=29
ID=30
ID=31

当再次运行此程序时，我们可能会以完全不同的顺序得到输出，这取决于线程的执行顺序。

ID=4
ID=5
ID=6
ID=7
ID=0
ID=1
[snip]
ID=14
ID=15
ID=28
ID=29
ID=30
ID=31

sycl::stream 的输出是在 kernel 完成执行后打印的。在大多数情况下，这无关紧要。 然而，如果 kernel 出现故障或抛出异常，将不会打印任何语句。为了说明这一点， 请考虑以下 kernel，它引发了一个异常：

void out5() {
  int *m = NULL;
  sycl::queue q;
  q.submit([&](auto &cgh) {
     sycl::stream str(8192, 1024, cgh);
     cgh.parallel_for(sycl::nd_range<1>(32, 4), [=](sycl::nd_item<1> it) {
       int id = it.get_global_id();
       str << "ID=" << id << sycl::endl;
       if (id == 31)
         *m = id;
     });
   }).wait();
} // end out5

编译并执行上述代码会由于写入空指针而产生段错误。

Segmentation fault (core dumped)

没有任何打印语句实际上被打印到控制台。相反，您会看到一个关于段错误的错误消息。这与传统的 C/C++ stream 不一样。

上一章                                         主目录    上级目录                                                               下一章