1

I have been exploring the field of parallel programming and have written basic kernels in Cuda and SYCL. I have encountered a situation where I had to print inside the kernel and I noticed that std::cout inside the kernel does not work whereas printf works. For example, consider the following SYCL Codes - This works -

void print(float*A, size_t N){
    buffer<float, 1> Buffer{A, {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]]);
       });
    });
}

whereas if I replace the printf with std::cout<<accessor[idx[0]] it raises a compile time error saying - Accessing non-const global variable is not allowed within SYCL device code. A similar thing happens with CUDA kernels. This got me thinking that what may be the difference between printf and std::coout which causes such behavior.

Also suppose If I wanted to implement a custom print function to be called from the GPU, how should I do it?
TIA

4

3 回答 3

3

在 SYCL 中，由于与CUDA 代码std::cout的答案中列出的类似原因，您不能将其用于未在主机上运行的代码的输出。

这意味着如果您在“设备”（例如 GPU）上运行内核代码，那么您需要使用stream该类。在名为 Logging 的 SYCL 开发人员指南部分中有更多信息。

于 2021-02-02T09:01:56.993 回答
2

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

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

std::cout没有得到这种编译器辅助的替换/劫持——它的代码在 GPU 上根本不相关。

但是 - 我已经实现了一种std::cout用于 GPU 内核的类似机制；有关更多信息和链接，请在此处查看我的这个答案。

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

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

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

于 2021-02-01T22:53:34.513 回答
0

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

于 2021-02-01T17:26:52.153 回答