☰

为什么 Nvidia Pascal 同时拥有 FP32 和 FP64 核心？为什么我不能同时使用它们？

tag-icon • tag-icon 64-bit nvidia-graphics-card gpu 32-bit gpgpu
为什么 Nvidia Pascal 同时拥有 FP32 和 FP64 核心？为什么我不能同时使用它们？

我正在尝试了解 Nvidia 的 GPU 架构，但我对一些看似相当简单的事情有点困惑。Pascal 中的每个流式多处理器都由 64xFP32 和 32xFP64 核心组成。以下是我的两个问题：

  • Nvidia 为什么要在芯片中同时放置 FP32 和 FP64 单元？为什么不只放置能够每条指令执行 2xFP32 操作的 FP64 单元（就像 CPU 中的 SIMD 指令集一样）。
  • 为什么我不能同时使用所有 FP32 和 FP64 单元？

我猜两者都是硬件设计决定，但我想了解有关此主题的更多详细信息。任何有关此方面的信息都非常欢迎！

编辑1：

  • 如果可以同时执行 FP32 和 FP64，这是否意味着具有 8TFLOPS SP 和 4TFLOPS DP 的 GPU 可以为您提供（理论上）12 TFLOPS 混合 TFLOPS？
    • 如果是 CUDA，如何实现这一点？我是否只在内核中同时使用双精度和浮点？还是我需要将某种标志传递给 NVCC？

答案1

为什么 Nvidia 要在芯片中同时放置 FP32 和 FP64 单元？

我认为这与市场渗透有关，尽可能多地销售。如果没有 FP64，科研人员甚至无法尝试使用 FP64 的科学上重要的 gpgpu 软件的演示（甚至游戏有时也可能使用一些双精度）。如果没有 FP32，游戏物理和模拟将非常缓慢，或者 GPU 将需要核反应堆。没有 FP16，就没有快速神经网络。如果只有 FP32，神经网络模拟将以半速运行，或者某些 FP64 求和将不起作用。

谁知道呢，也许未来会有 FP_raytrace 专用核心，可以超快地进行光线追踪，这样就不再需要 DX12 DX11 DX9 痛苦的升级和更好的图形。

最终，我不会拒绝基于 FPGA 的 GPU，它可以将某些内核从 FP64 转换为 FP32 或某些特殊功能内核，用于某个应用程序，然后为另一个应用程序将所有内核转换为 FP64，甚至将所有内容转换为执行顺序工作（例如编译着色器）的单个胖内核。这将使在计算机上做许多不同事情的人受益。例如，我可能需要比加法更多的乘法，而 FPGA 可以在这里提供帮助。但现在，金钱至上，它说“目前是固定功能”，最好的收入是通过 FP64 和 FP32（以及最近的 FP16）的混合来实现的。

为什么不直接放置能够每条指令执行 2xFP32 操作的 FP64 单元（就像 CPU 中的 SIMD 指令集一样）。

SIMD 总是期望对多个数据执行相同的操作，而标量 GPGPU 内核则没那么有趣。此外，用 FP64 制作 2xFP32 需要比纯 FP64 更多的晶体管，可能产生更多热量，延迟也更大。

晶体管数量越多 = 生产失败概率越大，因此 1024 FP32 GPU 的生产可能性比 512 FP64_flexible GPU 更大。

为什么我不能同时使用所有 FP32 和 FP64 单元？

混合精度计算可以在 cuda 和 opencl 中完成，因此您可以使用所有核心获得更快的速度，但仅适用于非内存瓶颈情况，这种情况很少见且难以编码。

编辑 1 的答案：

这是一个详细的来源http://www.nvidia.com/content/PDF/sc_2010/CUDA_Tutorial/SC10_Accelerating_GPU_Computation_Through_Mixed-Precision_Methods.pdf

长话短说，它们不添加，存在“收益递减”，由于不同精度计算之间需要“额外周期”，因此无法在所有核心上进行 %100 缩放。当它们不混合时，它们需要块之间的“额外迭代”，这也不允许 %100 缩放。它似乎更有用，因为可以加快“FP64”而不是降低“FP32”（但拥有许多 FP64 核心应该是有益的（对于提高 FP32），您可以使用类似 nbody 内核（它没有内存瓶颈）来测试它们）。FP64 非常耗费内存（和缓存行（和本地内存）），这就是为什么我建议使用 nbody 算法，它将一些数据重复使用 N（例如>64k）次。我的 GPU 有 1/24 FP64 功率，所以我不信任我的电脑。你有泰坦吗？你应该尝试一下，也许它的功率比其广告中的 GFLOP 值高出 50%。（但广告中的 TDP 值可能会以这种方式限制其频率，并崩溃）

此来源：http://www.nvidia.com/content/nvision2008/tech_presentations/NVIDIA_Research_Summit/NVISION08-Mixed_Precision_Methods_on_GPUs.pdf

说是“出色的性能和准确性”，但我找不到使用 FP32 + FP32（截断的 FP64）的游戏物理解算器，也许金钱再次发挥作用，如果有人做到这一点，那将是游戏上的“出色的性能和崩溃”。（可能比 furmark 爆炸的 gpu 更糟糕）

人们甚至在浮点数上使用整数（整数点积）：https://devblogs.nvidia.com/parallelforall/mixed-precision-programming-cuda-8/

如果是 CUDA，如何实现这一点？我是否只在内核中同时使用双精度和浮点？还是我需要将某种标志传递给 NVCC？

在同一函数中使用 fp64+fp32 进行迭代细化的示例：

https://www.sciencesmaths-paris.fr/upload/Contenu/HM2012/07-dongarra_part2.pdf

第26-28页。

对于 opencl 部分，这里是 amd evergreen（hd5000 系列），能够每周期发出 1dp fma + 1 sp（或 1 sf）。

http://www.microway.com/download/whitepaper/gpgpu_architecture_and_performance_comparison_2010.pdf

明天我将在我的 R7-240 上测试类似 nbody 的东西，它是 fp32 的 1/24 或 1/26 次方，即 fp64。

编辑：它正在工作。

__kernel void sumGPU(__global float * a,__global float * b)
{
  int idx = get_global_id(0);
  float a0=a[idx];
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  //a0=convert_float(convert_double(a0)+2.0);
  //a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  a0+=2.0f;
  b[idx] = a0;

}

当仅禁用其中一条注释时，它会在 13.02 毫秒和 12.85 毫秒之间切换。

注意：内核本身不是 fp32。没有内核。有调度程序将硬件资源（fp32、fp64、special_function、registers）绑定到线程的内核指令。线程也不是真正的线程。因此，当您使用 fp32、fp64、fp32、fp64_square_root 时，它会在需要时保留必要的资源。当不需要时，它们是其他工作项目的选项。（但我怀疑单个工作项目不能使用超过 1-2 个 fp32 ALU（我不知道，这是我编造的））

编辑（2018/03）：FP_raytrace（上述答案的第二段）正在成为现实吗？

(英伟达) https://www.geforce.com/whats-new/articles/nvidia-rtx-real-time-game-ray-tracing

(AMD) https://www.gamingonlinux.com/articles/amd-has-announced-radeon-rays-an-open-source-ray-tracing-sdk-using-vulkan.11461

或者这只是另一个营销噱头？如果它有硬件方面，那么光线追踪器人员可以更快地工作，但对 moba 游戏玩家或无光线追踪器的物理模拟器没有帮助。如果我要编辑一些视频，为什么我要为这些光线追踪器支付更多费用？也许这些也可以与其他产品一样细分，但可能需要更多钱。

相关内容