问题标签 [gpu-cooperative-groups]

CUDA 运行时 API 允许我们使用 variable-number-of-arguments 三人字形语法启动内核：

但关于“协作”内核，CUDA Programming Guide 说（第 C.3 节）：

要启用网格同步，在启动内核时需要使用CUDA 运行时启动 API ，而不是<<<...>>>执行配置语法：cuLaunchCooperativeKernel

（或等效的 CUDA 驱动程序）。

我宁愿不必编写自己的包装器代码来构建指针数组……运行时 API 中真的没有设施可以避免这种情况吗？

在阅读了 CUDA 9 中的合作组之后，我一直在尝试在网格级别进行同步。

我正在使用 Visual Studio 2017、GTX 1060 和 CUDA 9.1。

我改变了我的代码如下：

而且，如文档中所述，我这样称呼我的内核：

不幸的是，我在内核中添加“网格”部分后立即出现链接错误。

我已经设置了 -rdc=true 和 sm_61 但找不到它为什么不起作用...有什么想法吗？

非常感谢 ！

遵循CUDA 示例中的相同步骤来启动内核并跨网格同步cooperative_groups::this_grid().sync()会导致任何 CUDA API 调用失败。虽然使用 cooperative_groups::this_thread_block().sync()工作正常并给出正确的结果。

我使用以下代码和CMakeLists.txt（cmake 版本 3.11.1）在带有 Ubuntu 16.04.5 LTS 的 TITAN V GPU（驱动程序版本 410.73）上使用 CUDA 10 对其进行测试。该代码也可以在github上找到，以便于重现错误。

该代码读取一个数组，然后将其反转（从[0 1 2 ... 9]到[9 8 7 ... 0]）。为了做到这一点，每个线程从数组中读取单个元素，同步，然后将其元素写入正确的目的地。可以轻松修改代码以确保this_thread_block().sync()正常工作。只需更改arr_size为更少的 1024 并cg::thread_block barrier = cg::this_thread_block();改用。

test_cg.cu

CMakeLists.txt

运行此代码给出：

unknown error in /home/ahdhn/test_cg/test_cg.cu at line 67

这是使用cudaMalloc. 我通过从设备查询确保代码编译为正确的架构，结果为 700。如果您发现我在代码或文件__CUDA_ARCH__中做错了什么，请告诉我。CMakeLists.txt

我的 GPU 是 GeForce MX150，pascal 架构，CC。6.1，CUDA 9.1，windows 10。虽然我的GPU是pascal但合作组不起作用。我想将它用于块间同步。
我发现我的 tcc 模式没有激活。
我还发现它在 Windows 中的 wddm 中不起作用。
如何使用合作组？
如何在 windows 中激活 tcc 模式？
感谢您的回复。

假设我有两个任务要在 GPU 上运行，第二个任务基本上依赖第一个任务的所有工作。传统上，我基本上必须将这些任务编写为两个单独的内核，并安排第二个在第一个之后的某个时间运行。但是 - 使用 CUDA 9，我现在可以在整个网格上同步，完成第一个任务的工作 -使用协作组功能，然后继续让网格完成第二个任务的工作。

我的问题是：

• 我们能否就编写两个内核的性能以及何时使用全网格同步提供一个经验法则？
• 如果是这样，那会是什么？
• 如果不是 - 为什么很难确定在哪些情况下哪个更可取？

我正在为我的仓库构建示例cuda-api-wrappers。前段时间，我遇到一个问题：每当我尝试.cu使用“合作组”机制编译一个编码的文件时，我都会收到一条错误消息：

这很奇怪，因为我只是在编译。好吧，我使用这个 nVIDIA DevTalk 论坛帖子上的建议解决了这个问题：

我建议...指定-rdc=true

我做到了，这似乎奏效了。然而，现在（在 Devuan 3 ~= Debian 10 上使用 CUDA 10.1），我在链接过程中遇到错误：

为什么我仍然收到这些错误，我该怎么办？

考虑以下 CUDA 程序，在一个名为 的文件中foo.cu：

这个程序做的不多——但它是一个有效的程序（如果你的计算能力足够高，可以支持整个网格作为一个合作组）。它应该编译链接并运行。但是，我明白了：

奇怪！如果我添加一些特定-l的-L标志并没有帮助，例如：

为什么会这样？我应该如何修改 nvcc 命令行以使其找到未解析的符号？

笔记：

• 我正在使用 Devuan GNU/Linux 3.0。
• CUDA 10.1 作为分发包安装，因此其库位于/usr/lib/x86_64-linux-gnu.
• 带有 GeForce 1050 Ti 卡的 x86_64 机器。

当我遇到内核根本不会做任何事情时，我试图利用协作组功能实现一些 cuda 内核。我把它归结为一个小例子，如下所示。一旦我调用该sync()函数，cuda 内核就不再工作了。我已经确保我的 GPU（GTX 1070 Ti）计算能力足够高，并且我还在 nvcc 中添加了适当的开关（-rdc、-gencode）。

编译命令：

没有同步（）的内核

输出（正如预期的那样）

带有 sync() 的内核失败示例

输出（全部为 0）

完整代码：

问题

为什么会这样？

如何解决？

或有关故障排除的任何提示？cudaGetLastError()即使内核没有做任何事情，内核也会通过它返回成功，这真的很奇怪。

PS nvidia-smi 输出

所以我使用 GTX 1050，计算能力为 6.1 和 CUDA 11.0。我需要在我的程序中使用网格同步，所以cudaLaunchCooperativeKernel()需要。我检查了我的设备查询，所以 GPU 确实支持合作组。我无法执行以下功能

打电话后，

收到错误“不允许操作”（代码为 800）。现在，当设备不支持协作组时返回（在这种情况下不支持）。那么，什么可能导致这个问题呢？

我们正在尝试使用 CUDA 中的一些协作组功能来编写一个小型应用程序。我们正在使用带有 CUDA 11.0 的 Tesla V100 卡。但是在 thread_block 中使用 is_valid() 方法时，观察到以下错误：

在CUDA提供的示例simpleCooperativeGroups.cu中，在内核cgkernel()中使用is_valid方法时，可以看到这个错误。修改后的示例中的代码片段如下所示。

任何解决此问题的建议都会有很大帮助。