时间戳

首页 > 解决方案 > 如何在 GPU 上计划这种最小二乘计算？

opencl - 如何在 GPU 上计划这种最小二乘计算？

问题描述

我正在编写一个名为 PerfectTIN ( https://github.com/phma/perfecttin ) 的程序，它对 TIN 进行大量最小二乘调整以适应点云。每次调整都需要一些连续的三角形组，并调整最多 8 个点的高程，这些点是三角形的角，以适应三角形中的点。我让它在 SMP 上工作。在开始处理时，它一次只进行一次调整，因此它将调整拆分为任务，每个任务需要一些点，所有这些点都在同一个三角形中。每个线程从队列中获取一个任务并计算一个小方阵和一个小列向量。当它们都准备好后，调整例程将矩阵和向量相加并完成最小二乘计算。

我想在 GPU 和 CPU 上处理任务。任务所需的数据是

点数是 1024 的倍数，除了一些我可以在 CPU 中处理的任务（三角形中的点总数可以是任何非负整数）。对于 5600 万个点的相当大的点云，一些任务大于 131072 个点。

这是输出的一部分clinfo（如果您需要其他部分，请告诉我）：

  Platform Name                                   Clover
Number of devices                                 1
  Device Name                                     Radeon RX 590 Series (POLARIS10, DRM 3.33.0, 5.3.0-7625-generic, LLVM 9.0.0)
  Device Vendor                                   AMD
  Device Vendor ID                                0x1002
  Device Version                                  OpenCL 1.1 Mesa 19.2.8
  Driver Version                                  19.2.8
  Device OpenCL C Version                         OpenCL C 1.1 
  Device Type                                     GPU
  Device Profile                                  FULL_PROFILE
  Device Available                                Yes
  Compiler Available                              Yes
  Max compute units                               36
  Max clock frequency                             1545MHz
  Max work item dimensions                        3
  Max work item sizes                             256x256x256
  Max work group size                             256
  Preferred work group size multiple              64
  Preferred / native vector sizes                 
    char                                                16 / 16      
    short                                                8 / 8       
    int                                                  4 / 4       
    long                                                 2 / 2       
    half                                                 8 / 8        (cl_khr_fp16)
    float                                                4 / 4       
    double                                               2 / 2        (cl_khr_fp64)
  Double-precision Floating-point support         (cl_khr_fp64)
    Denormals                                     Yes
    Infinity and NANs                             Yes
    Round to nearest                              Yes
    Round to zero                                 Yes
    Round to infinity                             Yes
    IEEE754-2008 fused multiply-add               Yes
    Support is emulated in software               No

如果我没看错的话，如果我在 GPU 的每个核心上放一个点，那么我一次可以处理的点总数是 36×256=9×1024=9216。我可以在每个核心中放置四个点，因为一个工作组将有 1024 个点吗？在这种情况下，我可以一次处理 36864 个点。每个核心应该处理多少个点？如果一个任务比 GPU 可以容纳的大怎么办？如果多个任务（可能来自不同的三角形）适合 GPU 怎么办？

当然，我希望这段代码可以在我以外的其他 GPU 上运行。我将使用 OpenCL 来实现可移植性。我会遇到哪些不同的 GPU（描述而不是名称）？

标签： openclgpgpu

解决方案

如果我没看错的话，如果我在 GPU 的每个核心上放一个点，那么我一次可以处理的点总数是 36×256=9×1024=9216。

不完全的。点的总数不受最大工作组大小的限制。工作组大小是一个组中同步工作的 GPU 线程数。在工作组内，您可以通过本地内存共享数据，这对于加速某些计算（例如矩阵乘法）很有用（例如缓存平铺）。

GPU 并行化的想法是将问题分解成尽可能多的独立部分。在 C++ 中类似

void example(float* data, const int N) {
    for(int n=0; n<N; n++) {
        data[n] += 1.0f;
    }
}

在 OpenCL 中变成这个

kernel void example(global float* data ) {
    const int n = get_global_id(0) ;
    data[n] += 1.0f;
}

其中内核的全局范围设置为N.

每个 GPU 线程应该只处理一个点。线程数（全局范围）可以而且应该远大于可用的 GPU 内核数。如果您不明确需要本地内存（所有线程都可以独立工作），您可以将本地工作组大小设置为 32、64、128 或 256 - 没关系 - 但这些之间可能存在一些性能差异价值观。但是，全局范围（线程总数/点）必须是工作组大小的倍数。

如果一个任务比 GPU 可以容纳的大怎么办？

我假设您的意思是当您的数据集不能一次全部放入视频内存时。在这种情况下，您可以分批进行计算，因此使用 PCIe 传输交换 GPU 缓冲区。但这会带来很大的性能损失。

当然，我希望这段代码可以在我以外的其他 GPU 上运行。我将使用 OpenCL 来实现可移植性。我会遇到哪些不同的 GPU（描述而不是名称）？

OpenCL 非常适合跨设备和操作系统的可移植性。除了 AMD GPU 和 CPU，您还会遇到仅支持 OpenCL 1.2 的 Nvidia GPU 和 Intel GPU 和 CPU。如果安装了图形驱动程序，您的代码将在所有这些驱动程序上运行而不会出现问题。请注意，视频内存的数量可能会有很大差异。

推荐阅读