CUDA thrust库中的函数是否隐式同步？

Question

当我使用thrust库中的函数时，遇到了一些问题。我不确定在它之前是否需要手动添加cudaDeviceSynchronize。例如：

double dt = 0;
kernel_1<<<blocks, threads>>>(it);
dt = *(thrust::max_element(it, it + 10));
printf("%f\n", dt);

由于kernel_1是非阻塞的，主机将执行下一条语句。问题在于我不确定thrust::max_element是否是阻塞的。如果它是阻塞的，那么它就能正常工作；否则，主机会跳过它并执行“printf”语句吗？

谢谢

- fois

Answer 1

你的代码至少有两处错误。

1. it is presumably a device pointer:

   kernel_1<<<blocks, threads>>>(it);
                                 ^^
   

   it is not allowed to use a raw device pointer as an argument to a thrust algorithm:

   dt = *(thrust::max_element(it, it + 10));
                              ^^
   

   unless you wrap that pointer in a thrust::device_ptr or else use the thrust::device execution policy explicitly as an argument to the algorithm. Without any of these clues, thrust will dispatch the host code path (which will probably seg fault) as discussed in the thrust quick start guide.

2. If you fixed the above item using either thrust::device_ptr or thrust::device, then thrust::max_element will return an iterator of a type consistent with the iterators passed to it. If you pass a thrust::device_ptr it will return a thrust::device_ptr. If you use thrust::device with your raw pointer, it will return a raw pointer. In either case, it is illegal to dereference such in host code:

   dt = *(thrust::max_element(it, it + 10));
        ^
   

   again, I would expect such usage to seg fault.

关于异步性，可以安全地假设返回存储在堆栈变量中的标量数量的所有推力算法都是同步的。这意味着CPU线程在推力调用之后不会继续进行，直到堆栈变量被填充了正确的值。

关于GPU活动，除非使用流，否则所有GPU活动都将发出到同一个（默认）流中。这意味着所有CUDA活动都将按顺序执行，并且给定的CUDA操作不会在前面的CUDA活动完成之前开始。因此，即使您的内核启动是异步的，并且CPU线程将继续执行thrust :: max_element调用，从该调用产生的任何CUDA活动也不会开始执行，直到前一个内核启动完成。因此，kernel_1对由it引用的数据所做的任何更改都应该在thrust :: max_element中的任何CUDA处理开始之前完成并完全有效。正如我们所看到的，thrust :: max_element本身将插入同步。

因此，一旦您修复代码中的缺陷，就不应该在任何地方插入额外的同步要求。

- Robert Crovella

---

感谢您的详细解释。 - fois

Answer 2

这个函数似乎不是异步的。

这两个页面都解释了max_element()的行为，但它们没有将其明确说明为异步，因此我认为它是阻塞的：

• thrust : Extrema
• max_element

由于它使用迭代器处理所有元素并找到值的最大值，我无法想象它是异步的。

您仍然可以使用cudaDeviceSynchronize来尝试实际运行它，但不要忘记在设备上设置相应的标志。

- Taro