问题标签 [dynamic-parallelism]

关于Kepler的资料很少dynamic parallelism，从这项新技术的描述来看，是不是意味着同一个warp中线程控制流发散的问题已经解决了？

它允许recursion从设备代码中获取内核，这是否意味着不同线程中的控制路径可以同时执行？

我正在尝试使用动态并行性来改进我在 CUDA 中的算法。在我最初的 CUDA 解决方案中，每个线程都会计算每个块通用的数字。我想要做的是首先启动一个粗略（或低分辨率）内核，其中线程只计算一次公共值（就像每个线程代表一个块一样）。然后每个线程创建一个由 1 个块组成的小网格（16x16 线程），并为它启动一个子内核以传递公共值。理论上它应该更快，因为它可以节省许多冗余操作。但在实践中，该解决方案的工作速度非常慢，我不知道为什么。

这是代码，非常简化，只是想法。

child_kernels 的数量很多，每个线程一个，并且必须有大约 400x400 个线程。据我了解，GPU 应该并行处理所有这些内核，对吧？

或者子内核以某种方式按顺序处理？

我的结果表明，性能比我原来的解决方案慢 10 倍以上。

我有一个已经并行化的 CUDA 内核，它执行一些需要频繁插值的任务。

所以有一个内核

它调用一次或多次此插值设备函数：

插值算法在三个维度上连续进行 WENO 插值。这是一个高度可并行化的任务，我迫切希望并行化！

很明显，内核complexStuff()可以很容易地通过使用<<<...>>>语法从主机代码调用它来实现并行化。complexStuff()已经并行化也很重要。

但是我不清楚如何从 CUDA 设备函数内部并行化某些东西/创建新线程......这甚至可能吗？有人知道吗？

我有以下代码片段，并且正在试验新的开普勒架构的功能。内核在具有固定 NUM_ITERATIONS 的循环中被调用多次。您是否认为将循环转移到父内核会有所帮助，即，与 CPU 相比，从 GPU 调用时内核开销是否更小？

是否可以使用动态并行来提高以下算法的性能？如果是这样，您能否建议一个类似的动态并行用例来帮助我在自己的程序中实现它？

I have 2 cpp files setup and functions, 6 .cu files main, flood, timestep, discharge, continuity and copy. I'm trying to compile this to the main call the cpp files and so the flood kernel global and then flood call timestep, discharge, continuity and copy kernels all device.

Something like this:

Main
~functions
~setup
~flood
~~timestep
~~discharge
~~continuity
~~copy

I'm using a GK110 board with CUDA 5.5, but i don't know how to compile in separate sources (I got an error, that device functions requires separation compile mode). I'm also dont know how to use -dc or -rtc={true} and flags for dynamic parallelism.

I tried to do this, but doesn't work:

When i reach de fifth line I got the error message that the device function call cannot be configured.

Someone can help me?

我的电脑有一个 GTX 580（计算能力 2.0）。

我想编译一个使用动态并行性的 CUDA 源代码，这是计算能力 3.5 中引入的一项功能。

我知道我将无法在我的 GPU 上运行该程序，但是，应该可以在我的机器上编译此代码。我假设这是因为我可以毫无问题地编译使用 3.5 功能的 CUDA 示例。这些示例带有“手动生成”的 Visual Studio 项目（我猜）。

我相信我的问题在于 CMake。我正在使用 CMake 生成 Visual Studio 2012 项目。

我的第一个 CMakeLists.txt 看起来像这样：

然后，在使用生成的 Visual Studio 2012 项目进行编译时，我收到一个警告，然后是一个错误：

warning : The 'compute_10' and 'sm_10' architectures are deprecated, and may be removed in a future release.

error : calling a __global__ function from a __global__ function is only allowed on the compute_35 architecture or above

预期什么。然后我加了

list(APPEND CUDA_NVCC_FLAGS -gencode arch=compute_35,code=sm_35)

到 CMakeLists。警告消失了，但我得到了：

error : kernel launch from __device__ or __global__ functions requires separate compilation mode

好的。所以我添加到 CMakeLists：

set(CUDA_SEPARABLE_COMPILATION ON)

...并收到了这个：

fatal error : nvcc supports '--relocatable-device-code=true (-rdc=true)', '--device-c (-dc)', and '--device-link (-dlink)' only when targeting sm_20 or higher

奇怪的是，我以为我的目标是 sm_35（高于 sm_20）。

后来我发现我可以直接在 CUDA_ADD_EXECUTABLE 命令中设置一些选项。所以我删除了将值附加到 CUDA_NVCC_FLAGS 的行并将 CUDA_ADD_EXECUTABLE 命令更改为：

我得到的是：

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\bin\crt\link.stub : fatal error C1083: Cannot open compiler generated file: 'C:/Users/sms/Desktop/sample-cuda-tests/CMakeFiles/sample-cuda-tests.dir/Debug/sample-cuda-tests_intermediate_link.obj': No such file or directory

现在不知道该去哪里。感谢任何帮助。

我在 Windows 7 上使用 CUDA SDK 6.0。

这个部署套件给我留下了深刻的印象。而不是购买可能需要新主板等的新 CUDA 卡，这张卡似乎提供了所有功能。

在它的规格上，它说它具有 CUDA 计算能力 3.2。AFAIK 动态并行和更多来自 cm_35，cuda 计算能力 3.5。此卡是否支持 Kepler 架构的 Dynamic Parallelism 和 HyperQ 功能？

当内核中的内核调用（甚至递归调用）使用纹理内存来获取值时，我似乎遇到了麻烦。

如果子内核，比如说另一个内核，不使用纹理内存，一切都很好。如果我不在内核中调用内核，则结果是预期的。 只要我使用由于空间局部性和快速过滤而在我的情况下非常有用的纹理内存，cuda-memcheck 就会返回"Invalid __global__ write of size 4"。

我已经看到，在编程​​指南中的动态并行中，使用纹理内存时必须小心，这可能会导致数据不一致，但这里子内核甚至没有启动。

我试过 __syncthreads() 和 cudaDeviceSynchronize 在调用纹理内存之前或之后放置，但没有。

是否有一些已经报告的案例，是我做错了什么还是只是你不能那样使用纹理内存？

系统：gtx泰坦黑（sm_3.5），CUDA6.0。

编辑：一些示例代码来说明。

显然，之前已声明并填充了EField 。HANDLE_ERROR来自 book.h 包括来自CUDA 的示例

这是一个可编译的代码：

在 CUDA 编程指南部分 C.4.3.1.2 下。《嵌套与同步深度》中提到：

“如果系统检测到在父内核从不调用 cudaDeviceSynchronize() 的情况下不需要为父状态保留空间，则允许进行优化。在这种情况下，由于从不发生显式父/子同步，因此所需的内存占用程序将远小于保守最大值。这样的程序可以指定较浅的最大同步深度以避免后备存储的过度分配"

这是否意味着编译器支持动态并行下的尾递归？例如，如果我有一个递归调用自身的内核：

在主机上启动：

如果编译器可以优化尾递归，它是否仍然受到最大递归级别 24 的限制，因为“父/子同步永远不会发生”？如果不限制，编译器怎么开启优化呢？

谢谢！

我写了一个简单的代码来理解动态并行。从打印的值中，我看到子内核已正确执行，但是当我回到父内核时，我看到使用了错误的值来代替在子内核中正确更新的临时数组。当我尝试更新“d_cin 数组”时，它给了我错误的值。这些是正在使用的编译标志：

有人能帮我吗 ？这是代码。