为啥 printf() 可以在内核中工作，但使用 std::cout 不能？

Posted 2023-03-10

【中文标题】为啥 printf() 可以在内核中工作，但使用 std::cout 不能？

【英文标题】：Why does printf() work within a kernel, but using std::cout doesn't?为什么 printf() 可以在内核中工作，但使用 std::cout 不能？

【发布时间】：2021-02-01 17:20:08

【问题描述】：

我一直在探索并行编程领域，并用 Cuda 和 SYCL 编写了基本内核。我遇到了必须在内核内部打印的情况，我注意到内核内部的std::cout 不起作用，而printf 起作用。例如，考虑以下 SYCL 代码 - 这行得通 -

void print(float*A, size_t N)
    buffer<float, 1> BufferA, N;
    queue Queue((intel_selector()));
    Queue.submit([&Buffer, N](handler& Handler)
       auto accessor = Buffer.get_access<access::mode::read>(Handler);
       Handler.parallel_for<dummyClass>(range<1>N, [accessor](id<1>idx)
           printf("%f", accessor[idx[0]]);
       );
    );

而如果我将 printf 替换为 std::cout<<accessor[idx[0]] 则会引发编译时错误 - Accessing non-const global variable is not allowed within SYCL device code. CUDA 内核也会发生类似的事情。 这让我想到 printf 和 std::coout 之间可能有什么区别导致这种行为。

另外假设如果我想实现一个从 GPU 调用的自定义打印函数，我应该怎么做？ TIA

【参考方案1】：

在 SYCL 中，您不能使用 std::cout 输出未在主机上运行的代码，原因与 answer for CUDA code 中列出的原因类似。

这意味着如果您在“设备”（例如 GPU）上运行内核代码，那么您需要使用 stream 类。在SYCL developer guide section called Logging 中有更多相关信息。

【讨论】：

您能否详细说明流类与std::ostream 的相似或不同之处？

用法非常相似，即调用 mystream

【参考方案2】：

这让我想到 printf 和 std::cout 之间可能有什么区别导致这种行为。

是的，有区别。在您的内核中运行的printf() 不是标准C 库printf()。对设备上的函数（其代码已关闭，如果它在 CUDA C 中存在的话）进行不同的调用。该函数使用 NVIDIA GPU 上的硬件机制——内核线程打印的缓冲区，该缓冲区被发送回主机端，然后 CUDA 驱动程序将其转发到启动内核的进程的标准输出文件描述符。

std::cout 没有得到这种编译器辅助的替换/劫持 - 它的代码与 GPU 完全无关。

但是 - 我已经实现了一个类似std::cout 的机制，用于 GPU 内核；有关更多信息和链接，请参见我的this answer。

这意味着我必须自己回答你的第二个问题：

如果我想实现一个从 GPU 调用的自定义打印函数，我应该怎么做？

除非您有权访问未公开的 NVIDIA 内部 - 唯一的方法是使用 printf() 调用，而不是在主机端使用 C 标准库或系统调用。您本质上需要通过低级原始 I/O 设施对整个流进行模块化。这远非微不足道。

【讨论】：

非常感谢@einpoklum 的回答，这就解决了

@AtharvaDubey：如果回答了问题，请接受...

为了说明 printf() 在 SYCL 示例代码中起作用的原因，我想补充一点，这不是标准的 SYCL，但某些实现恰好支持内核中的 printf() 作为扩展。对于类似 cout 的机制，SYCL 标准提供了 stream 类，正如 Rod 在他的回答中指出的那样。

【参考方案3】：

没有__device__ 版本的std::cout，所以只能在设备代码中使用printf。

【讨论】：

感谢您的回复@Oblivion。你能详细说明一下吗？

@AtharvaDubey 从内核运行一个函数，它必须被定义为 global 或 device。您可以调用 printf，因为定义了 device printf。

哦，好的，非常感谢。有没有办法定义我的自定义函数？

@Oblivion 我猜你的意思是“有 no __device__ 版本”

this 可能感兴趣