时间戳

首页 > 解决方案 > 将单个浮点值更新到 GPU 以在 CUDA 内核中访问它的最快方法是什么？

cuda - 将单个浮点值更新到 GPU 以在 CUDA 内核中访问它的最快方法是什么？

问题描述

我有一个 opengl 粒子模拟，其中每个粒子的位置都是在 CUDA 内核中计算的。大多数内存驻留在 GPU 内存中，但只有一个浮点值，我必须从 CPU 更新每一帧。

目前我cudaMemcpyAsync()用来将浮点值复制到 GPU，但是（至少据我所知），这会大大降低性能。我尝试使用 nvproof 来查看哪些调用耗时最长，结果如下：

Calls    Avg       Min       Max   Name
477  2.9740us  2.8160us  4.5440us  simulation(float3*, float*, float3*, float*)
477  89.033us  18.600us  283.00us  cudaLaunchKernel
477  47.819us  10.200us  120.70us  cudaMemcpyAsync

我认为我不能对内核启动本身做太多事情，但从调用来看，每帧发生cudaMemcpyAsync()的时间似乎最长。

我也尝试过使用固定内存，并且cudaHostGetDevicePointer() 如此处所述，但是由于某种原因，这会进一步增加内核启动时间，从而为不需要 memcopy 功能节省的时间增加了更多。

我想必须有更好/更快的方法将我的单个浮点变量更新到 GPU？

标签： cuda

解决方案

最简单的方法是，您可以向模拟内核函数添加一个额外的参数作为简单浮点值，而不是作为浮点指针，以便数据直接通过内核启动参数结构，当您启动时 CUDA 发送到 GPU核心。然后，您完全避开了该数据复制命令。（我假设 CUDA 将内核的整个函数参数描述符数据打包到一个复制命令中，因为内核参数描述符空间受到几 kB 或更少的限制）。

simulation(fooPointer, 
           barPointer, 
           fooBarPointer, 
           floatVariable 
);

或者，尝试在数据更新和渲染之间或数据更新和计算之间进行双缓冲，以便模拟图像在模拟计算之后落后 1-2 帧（每帧时间变得更糟），但“每秒帧数”会增加。

如果它不是交互式模拟，那么通过双缓冲或三缓冲隐藏计算/渲染/数据延迟应该可以工作。

如果您在最小化每帧时间（更快地响应用户输入到仿真中？），那么您应该将浮点变量嵌入到您已经在仿真中发送/使用的数组的末尾或您正在使用的任何结构的末尾。如果您已经有一个 1MB+ 浮点缓冲区要发送到 GPU，那么将 4B(float) 附加到它的末尾应该没有太大区别，那么您可以从那里访问它。1 次复制操作应该比具有相同总大小的 2 次复制操作快。

如果您实际上在每帧仅向 GPU 发送 4B（使用一个简单的函数来生成该数据），那么（正如 3Dave 在评论中所说）您可以尝试添加一个额外的内核函数来更新 GPU 中的值，并且只需内核启动命令的开销，而不是复制命令开销和数据复制开销。从积极的方面来说，如果每个帧自动运行一个内核“图形”，而无需一次又一次地将所有内核排入队列，那么额外的内核开销可能会被隐藏。

这里，

https://devblogs.nvidia.com/cuda-graphs/

那个部分

我们将创建一个模仿这种模式的简单代码。然后，我们将使用它来演示与标准启动机制相关的开销，并展示如何引入包含多个内核的 CUDA Graph，这些内核可以在单个操作中从应用程序启动。

cudaGraphLaunch(instance, stream);

他们说，当算法中有许多（20）个内核时，这个“图形”功能中的每个内核启动开销只有 3-4 微秒。

由于图形也支持其他命令，您可以尝试在图形内的并行 cuda 流中复制和计算部分，并使用双缓冲切换它们的输入，以便所有 CUDA 事物在将输出发送到渲染之前都可以保留在 CUDA 的上下文中。

（也许）你甚至根本不需要改变数据机制。只需尝试将浮点数据作为二进制表示形式发送到指针值中，并且只从内核读取指针值（不是数据值）并将其转换回浮点数。如果您不尝试在内核中到达浮点数据表示的（错误）指针地址，我不知道 CUDA 是否会为此返回错误。

simulation(fooPointer, 
           barPointer, 
           fooBarPointer, 
           toPtr(floatData) // <----- float to 64/32 bit pointer value
);

在内核中

float val = fromPtrToFloat(parameter4); // converts pointer itself, not the data

但是，如果您可以简单地使用“值”类型参数，这可能不是首选做法。

推荐阅读