oneAPI GPU 优化指南 - 子组（sub-group）和 SIMD 向量化

已于 2023-11-17 11:53:45 修改
阅读量164 收藏 1
点赞数
分类专栏： oneAPI GPU 优化指南 文章标签： oneapi 1024程序员节
于 2023-11-17 11:22:15 首次发布
原文链接：https://www.intel.com/content/www/us/en/docs/oneapi/optimization-guide-gpu/2023-1/sub-groups-and-simd-vectorization.html
版权

本章节翻译by chenshusmail@163.com 原文：Sub-Groups and SIMD Vectorization (intel.com)

目录

最大化内存带宽利用率

内存块加载和存储

数据共享

Sub-Group 大小 vs. 最大 Sub-Group 大小

ND-Range kernel 的索引空间被划分为 work-group，sub-group 和 work-item 。 work-item 是最基本的单位。一组 work-item 组成一个 sub-group， 一组 sub-group 组成一个 work-group。 work-item 和 work-group 到硬件向量引擎 (VE) 的映射是依赖于实现的。 所有的 work-group 都同时运行，但根据资源的可用性，它们可能被安排在不同的时间运行。 work-group 的执行可能会或者不会被抢占，这取决于底层硬件的能力。 同一 work-group 中的 work-item 保证同时运行。 同一 sub-group 中的 work-item 可能有额外的调度保证，并且可以访问额外的功能。

sub-group 是全局索引空间中连续的 work-item 的集合，它们在同一个 VE 线程中执行。 当设备编译器编译 kernel 时，多个 work-item 通过向量化被打包到一个 sub-group 中， 以便生成的 SIMD 指令流可以同时执行多个 work-item 的任务。 正确地将 work-item 划分为 sub-group 可以产生很大的性能差异。

让我们从一个简单的例子开始，来说明 sub-group：

  q.submit([&](auto &h) {
    sycl::stream out(65536, 256, h);
    h.parallel_for(sycl::nd_range(sycl::range{32}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int groupId = it.get_group(0);
                     int globalId = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     int sgGroupId = sg.get_group_id()[0];
                     int sgId = sg.get_local_id()[0];

                     out << "globalId = " << sycl::setw(2) << globalId
                         << " groupId = " << groupId
                         << " sgGroupId = " << sgGroupId << " sgId = " << sgId
                         << " sgSize = " << sycl::setw(2) << sgSize
                         << sycl::endl;
                   });
  });

这个例子的输出可能是这样的：

Device: Intel(R) Gen12HP
globalId = 0 groupId = 0 sgGroupId = 0 sgId = 0 sgSize = 16
globalId = 1 groupId = 0 sgGroupId = 0 sgId = 1 sgSize = 16
globalId = 2 groupId = 0 sgGroupId = 0 sgId = 2 sgSize = 16
globalId = 3 groupId = 0 sgGroupId = 0 sgId = 3 sgSize = 16
globalId = 4 groupId = 0 sgGroupId = 0 sgId = 4 sgSize = 16
globalId = 5 groupId = 0 sgGroupId = 0 sgId = 5 sgSize = 16
globalId = 6 groupId = 0 sgGroupId = 0 sgId = 6 sgSize = 16
globalId = 7 groupId = 0 sgGroupId = 0 sgId = 7 sgSize = 16
globalId = 16 groupId = 0 sgGroupId = 1 sgId = 0 sgSize = 16
globalId = 17 groupId = 0 sgGroupId = 1 sgId = 1 sgSize = 16
globalId = 18 groupId = 0 sgGroupId = 1 sgId = 2 sgSize = 16
globalId = 19 groupId = 0 sgGroupId = 1 sgId = 3 sgSize = 16
globalId = 20 groupId = 0 sgGroupId = 1 sgId = 4 sgSize = 16
globalId = 21 groupId = 0 sgGroupId = 1 sgId = 5 sgSize = 16
globalId = 22 groupId = 0 sgGroupId = 1 sgId = 6 sgSize = 16
globalId = 23 groupId = 0 sgGroupId = 1 sgId = 7 sgSize = 16
globalId = 8 groupId = 0 sgGroupId = 0 sgId = 8 sgSize = 16
globalId = 9 groupId = 0 sgGroupId = 0 sgId = 9 sgSize = 16
globalId = 10 groupId = 0 sgGroupId = 0 sgId = 10 sgSize = 16
globalId = 11 groupId = 0 sgGroupId = 0 sgId = 11 sgSize = 16
globalId = 12 groupId = 0 sgGroupId = 0 sgId = 12 sgSize = 16
globalId = 13 groupId = 0 sgGroupId = 0 sgId = 13 sgSize = 16
globalId = 14 groupId = 0 sgGroupId = 0 sgId = 14 sgSize = 16
globalId = 15 groupId = 0 sgGroupId = 0 sgId = 15 sgSize = 16
globalId = 24 groupId = 0 sgGroupId = 1 sgId = 8 sgSize = 16
globalId = 25 groupId = 0 sgGroupId = 1 sgId = 9 sgSize = 16
globalId = 26 groupId = 0 sgGroupId = 1 sgId = 10 sgSize = 16
globalId = 27 groupId = 0 sgGroupId = 1 sgId = 11 sgSize = 16
globalId = 28 groupId = 0 sgGroupId = 1 sgId = 12 sgSize = 16
globalId = 29 groupId = 0 sgGroupId = 1 sgId = 13 sgSize = 16
globalId = 30 groupId = 0 sgGroupId = 1 sgId = 14 sgSize = 16
globalId = 31 groupId = 0 sgGroupId = 1 sgId = 15 sgSize = 16

在这个例子中，每个 sub-group 有16个 work-item，也就是 sub-group 的大小为 16。 这意味着每个线程同时执行 16 个 work-item，32 个 work-item 由两个 VE 线程执行。

默认情况下，编译器使用设备特定信息和一些启发式方法来选择 sub-group 大小。 用户可以使用 kernel 属性 intel::reqd_sub_group_size 来指定最大的 sub-group 大小，从而覆盖编译器的选择。 有时候，显式地指定 sub-group 大小可能有助于提高性能，但并不总是如此。

  q.submit([&](auto &h) {
    sycl::stream out(65536, 256, h);
    h.parallel_for(sycl::nd_range(sycl::range{32}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(32)]] {
                     int groupId = it.get_group(0);
                     int globalId = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     int sgGroupId = sg.get_group_id()[0];
                     int sgId = sg.get_local_id()[0];

                     out << "globalId = " << sycl::setw(2) << globalId
                         << " groupId = " << groupId
                         << " sgGroupId = " << sgGroupId << " sgId = " << sgId
                         << " sgSize = " << sycl::setw(2) << sgSize
                         << sycl::endl;
                   });
  });

输出将是：

Device: Intel(R) Gen12HP
globalId = 0 groupId = 0 sgGroupId = 0 sgId = 0 sgSize = 32
globalId = 1 groupId = 0 sgGroupId = 0 sgId = 1 sgSize = 32
globalId = 2 groupId = 0 sgGroupId = 0 sgId = 2 sgSize = 32
globalId = 3 groupId = 0 sgGroupId = 0 sgId = 3 sgSize = 32
globalId = 4 groupId = 0 sgGroupId = 0 sgId = 4 sgSize = 32
globalId = 5 groupId = 0 sgGroupId = 0 sgId = 5 sgSize = 32
globalId = 6 groupId = 0 sgGroupId = 0 sgId = 6 sgSize = 32
globalId = 7 groupId = 0 sgGroupId = 0 sgId = 7 sgSize = 32
globalId = 8 groupId = 0 sgGroupId = 0 sgId = 8 sgSize = 32
globalId = 9 groupId = 0 sgGroupId = 0 sgId = 9 sgSize = 32
globalId = 10 groupId = 0 sgGroupId = 0 sgId = 10 sgSize = 32
globalId = 11 groupId = 0 sgGroupId = 0 sgId = 11 sgSize = 32
globalId = 12 groupId = 0 sgGroupId = 0 sgId = 12 sgSize = 32
globalId = 13 groupId = 0 sgGroupId = 0 sgId = 13 sgSize = 32
globalId = 14 groupId = 0 sgGroupId = 0 sgId = 14 sgSize = 32
globalId = 15 groupId = 0 sgGroupId = 0 sgId = 15 sgSize = 32
globalId = 16 groupId = 0 sgGroupId = 0 sgId = 16 sgSize = 32
globalId = 17 groupId = 0 sgGroupId = 0 sgId = 17 sgSize = 32
globalId = 18 groupId = 0 sgGroupId = 0 sgId = 18 sgSize = 32
globalId = 19 groupId = 0 sgGroupId = 0 sgId = 19 sgSize = 32
globalId = 20 groupId = 0 sgGroupId = 0 sgId = 20 sgSize = 32
globalId = 21 groupId = 0 sgGroupId = 0 sgId = 21 sgSize = 32
globalId = 22 groupId = 0 sgGroupId = 0 sgId = 22 sgSize = 32
globalId = 23 groupId = 0 sgGroupId = 0 sgId = 23 sgSize = 32
globalId = 24 groupId = 0 sgGroupId = 0 sgId = 24 sgSize = 32
globalId = 25 groupId = 0 sgGroupId = 0 sgId = 25 sgSize = 32
globalId = 26 groupId = 0 sgGroupId = 0 sgId = 26 sgSize = 32
globalId = 27 groupId = 0 sgGroupId = 0 sgId = 27 sgSize = 32
globalId = 28 groupId = 0 sgGroupId = 0 sgId = 28 sgSize = 32
globalId = 29 groupId = 0 sgGroupId = 0 sgId = 29 sgSize = 32
globalId = 30 groupId = 0 sgGroupId = 0 sgId = 30 sgSize = 32
globalId = 31 groupId = 0 sgGroupId = 0 sgId = 31 sgSize = 32

有效的 sub-group 大小与设备相关。您可以查询设备获取相关信息：

  std::cout << "Sub-group Sizes: ";
  for (const auto &s :
       q.get_device().get_info<sycl::info::device::sub_group_sizes>()) {
    std::cout << s << " ";
  }
  std::cout << std::endl;

支持的有效 sub-group 大小可能包括：

Device: Intel(R) Gen12HP
Subgroup Sizes: 8 16 32

接下来，我们将展示如何使用 sub-group 来提高性能。

向量化和内存访问

Intel® 的图形设备具有多个 VE。每个 VE 都是一个多线程的 SIMD 处理器。 编译器生成 SIMD 指令，将一个 sub-group 中的多个 work-item 打包，以便在 VE 线程中同时执行。 编译器选择的 SIMD 宽度（即 sub-group 大小）基于设备特征和启发式方法， 或由 kernel 显式指定，可以是 8、16 或 32。

在给定 SIMD 宽度的情况下，最大化 SIMD 通道利用率可以获得最佳的指令性能。 如果一个或多个通道（或 kernel 实例或 work-item）发生分歧，线程在路径合并之前执行两个分支路径， 增加动态指令计数。 SIMD 分歧对性能产生负面影响。 编译器努力最小化分歧，但如果可能的话，最好在源代码中避免分歧。

work-item 中内存的访问方式会影响 sub-group 中内存的访问方式或 SIMD 通道的利用方式。 在 work-item 中访问连续内存通常不是最佳选择。例如：

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int i = it.get_global_linear_id();
                     i = i * 16;
                     for (int j = i; j < (i + 16); j++) {
                       data[j] = -1;
                     }
                   });
  });
  q.wait();

这个简单的 kernel 初始化了一个 1024 x 1024 的整数数组。每个 work-item 初始化 16 个连续的整数。 假设编译器选择的 sub-group 大小为 16，则每个 sub-group 或线程中初始化 256 个整数。 然而，16 个 SIMD 通道中的存储是分散的。

与其在一个 work-item 中初始化 16 个连续的整数，不如在一个 SIMD 指令中初始化 16 个连续的整数更有效率。

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int i = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     i = (i / sgSize) * sgSize * 16 + (i % sgSize);
                     for (int j = 0; j < sgSize * 16; j += sgSize) {
                       data[i + j] = -1;
                     }
                   });
  });

我们在示例中使用了内存写入，但同样的技术也适用于内存读取。

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);
  int *data2 = sycl::malloc_shared<int>(N, q);
  memset(data2, 0xFF, sizeof(int) * N);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int i = it.get_global_linear_id();
                     i = i * 16;
                     for (int j = i; j < (i + 16); j++) {
                       data[j] = data2[j];
                     }
                   });
  });

这个 kernel 将一个 1024 x 1024 的整数数组复制到另一个相同大小的整数数组中。 每个 work-item 复制 16 个连续的整数。然而，从 data2 中读取的数据是聚集的， 而存储到 data 中的数据是分散的。更改代码， 使每个 sub-group 中而不是每个 work-item 读取和存储连续的整数会更有效率。

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);
  int *data2 = sycl::malloc_shared<int>(N, q);
  memset(data2, 0xFF, sizeof(int) * N);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int i = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     i = (i / sgSize) * sgSize * 16 + (i % sgSize);
                     for (int j = 0; j < sgSize * 16; j += sgSize) {
                       data[i + j] = data2[i + j];
                     }
                   });
  });

最大化内存带宽利用率

在上面的示例中，每个 work-item 在每次循环迭代中加载和存储 1 个整数或 4 个字节 (data[i + j] = data2[i + j];)， 或者每个向量化内存操作加载/存储 64 个字节（假设 sub-group 大小为 16），使内存带宽不饱和。

增加每个 work-item 在一次内存操作中加载或存储的有效负载或数据大小将导致更好的带宽利用率。

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);
  int *data2 = sycl::malloc_shared<int>(N, q);
  memset(data2, 0xFF, sizeof(int) * N);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
                   [=](sycl::nd_item<1> it) {
                     int i = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     i = (i / sgSize) * sgSize * 16 + (i % sgSize) * 4;
                     for (int j = 0; j < 4; j++) {
                       sycl::vec<int, 4> x;
                       sycl::vec<int, 4> *q =
                           (sycl::vec<int, 4> *)(&(data2[i + j * sgSize * 4]));
                       x = *q;
                       sycl::vec<int, 4> *r =
                           (sycl::vec<int, 4> *)(&(data[i + j * sgSize * 4]));
                       *r = x;
                     }
                   });
  });

每个 work-item 在每次循环迭代中加载/存储一个 sycl::vec<int,4> 而不是一个整数。 读取/写入 256 个连续字节的内存（假设 sub-group 大小为 16）， 每个向量化内存操作的有效负载增加了四倍。

最大带宽因硬件而异。可以使用 Intel® VTune Profiler 来测量带宽并找到最佳大小。

然而，使用向量类型可能会增加寄存器压力。建议仅在寄存器不溢出时使用长向量。 请参阅“寄存器化和避免寄存器溢出”一章，了解避免寄存器溢出的技巧。

内存块加载和存储

Intel® 图形具有针对内存块加载/存储优化的指令。 因此，如果 sub-group 中的 work-item 访问连续的内存块， 您可以使用 sub-group 块访问函数来利用块加载/存储指令。

  constexpr int N = 1024 * 1024;
  int *data = sycl::malloc_shared<int>(N, q);
  int *data2 = sycl::malloc_shared<int>(N, q);
  memset(data2, 0xFF, sizeof(int) * N);

  auto e = q.submit([&](auto &h) {
    h.parallel_for(
        sycl::nd_range(sycl::range{N / 16}, sycl::range{32}),
        [=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
          auto sg = it.get_sub_group();
          sycl::vec<int, 8> x;

          using global_ptr =
              sycl::multi_ptr<int, sycl::access::address_space::global_space>;
          int base = (it.get_group(0) * 32 +
                      sg.get_group_id()[0] * sg.get_local_range()[0]) *
                     16;
          x = sg.load<8>(global_ptr(&(data2[base + 0])));
          sg.store<8>(global_ptr(&(data[base + 0])), x);
          x = sg.load<8>(global_ptr(&(data2[base + 128])));
          sg.store<8>(global_ptr(&(data[base + 128])), x);
        });
  });

这个示例也使用 sycl::vec 来提高性能，但是 x 中的整数在内存中并不连续， 而是以 sub-group 大小为步长！

您可能已经注意到，显式指定了 sub-group 大小 16。 当您使用 sub-group 函数时，最好始终覆盖编译器的选择， 以确保 sub-group 大小始终符合您的预期。

数据共享

由于 sub-group 中的 work-item 在同一线程中执行，因此在 work-item 之间共享数据更有效率， 即使数据对每个 work-item 是私有的。 在 sub-group 中共享数据比使用共享本地内存或 SLM 在 work-group 中共享数据更有效率。 在 sub-group 中的 work-item 之间共享数据的一种方法是使用 shuffle 函数。

    constexpr size_t BLOCK_SIZE = 16;
    sycl::buffer<uint, 2> m(matrix.data(), sycl::range<2>(N, N));

    auto e = q.submit([&](auto &h) {
      sycl::accessor marr(m, h);
      sycl::local_accessor<uint, 2> barr1(
          sycl::range<2>(BLOCK_SIZE, BLOCK_SIZE), h);
      sycl::local_accessor<uint, 2> barr2(
          sycl::range<2>(BLOCK_SIZE, BLOCK_SIZE), h);

      h.parallel_for(
          sycl::nd_range<2>(sycl::range<2>(N / BLOCK_SIZE, N),
                            sycl::range<2>(1, BLOCK_SIZE)),
          [=](sycl::nd_item<2> it) [[intel::reqd_sub_group_size(16)]] {
            int gi = it.get_group(0);
            int gj = it.get_group(1);

            auto sg = it.get_sub_group();
            uint sgId = sg.get_local_id()[0];

            uint bcol[BLOCK_SIZE];
            int ai = BLOCK_SIZE * gi;
            int aj = BLOCK_SIZE * gj;

            for (uint k = 0; k < BLOCK_SIZE; k++) {
              bcol[k] = sg.load(marr.get_pointer() + (ai + k) * N + aj);
            }

            uint tcol[BLOCK_SIZE];
            for (uint n = 0; n < BLOCK_SIZE; n++) {
              if (sgId == n) {
                for (uint k = 0; k < BLOCK_SIZE; k++) {
                  tcol[k] = sg.shuffle(bcol[n], k);
                }
              }
            }

            for (uint k = 0; k < BLOCK_SIZE; k++) {
              sg.store(marr.get_pointer() + (ai + k) * N + aj, tcol[k]);
            }
          });
    });

这个 kernel 转置一个 16 x 16 的矩阵。它看起来比前面的例子更复杂，但思想很简单： 一个 sub-group 加载一个 16 x 16 的子矩阵，然后使用 sub-group shuffle 函数对子矩阵进行转置。 只有一个子矩阵，子矩阵就是矩阵，所以只需要一个 sub-group。更大的矩阵，比如 4096 x 4096， 可以使用相同的技术进行转置：每个 sub-group 加载一个子矩阵，然后使用 sub-group shuffle 函数对子矩阵进行转置。 这留给读者作为练习。

SYCL 提供了多种 sub-group shuffle 函数的变体。每种变体都针对特定设备上的特定用途进行了优化。 使用这些优化过的函数（如果它们符合您的需求）会是更好的选择，而不是创建您自己的函数。

Sub-Group 大小 vs. 最大 Sub-Group 大小

到目前为止，在我们的示例中，work-group 大小可以被 sub-group 大小整除， work-group 大小和 sub-group 大小（由用户指定或由编译器自动选择）都是 2 的幂。 如果 work-group 大小可以被最大 sub-group 大小整除且两个大小都是 2 的幂， 则 sub-group 大小和最大 sub-group 大小相同。但是，如果 work-group 大小不能被 sub-group 大小整除会发生什么呢？ 请看以下示例：

  auto e = q.submit([&](auto &h) {
    sycl::stream out(65536, 128, h);
    h.parallel_for(sycl::nd_range<1>(7, 7),
                   [=](sycl::nd_item<1> it) [[intel::reqd_sub_group_size(16)]] {
                     int i = it.get_global_linear_id();
                     auto sg = it.get_sub_group();
                     int sgSize = sg.get_local_range()[0];
                     int sgMaxSize = sg.get_max_local_range()[0];
                     int sId = sg.get_local_id()[0];
                     int j = data[i];
                     int k = data[i + sgSize];
                     out << "globalId = " << i << " sgMaxSize = " << sgMaxSize
                         << " sgSize = " << sgSize << " sId = " << sId
                         << " j = " << j << " k = " << k << sycl::endl;
                   });
  });
  q.wait();

这个示例的输出是这样的：

globalId = 0 sgMaxSize = 16 sgSize = 7 sId = 0 j = 0 k = 7
globalId = 1 sgMaxSize = 16 sgSize = 7 sId = 1 j = 1 k = 8
globalId = 2 sgMaxSize = 16 sgSize = 7 sId = 2 j = 2 k = 9
globalId = 3 sgMaxSize = 16 sgSize = 7 sId = 3 j = 3 k = 10
globalId = 4 sgMaxSize = 16 sgSize = 7 sId = 4 j = 4 k = 11
globalId = 5 sgMaxSize = 16 sgSize = 7 sId = 5 j = 5 k = 12
globalId = 6 sgMaxSize = 16 sgSize = 7 sId = 6 j = 6 k = 13

sub-group 大小为七，尽管最大 sub-group 大小仍然为 16！最大 sub-group 大小实际上是 SIMD 宽度， 所以它不会改变，但是 sub-group 中的 work-item 少于八个，所以 sub-group 大小为七。 因此，当您的 work-group 大小不能被最大 sub-group 大小整除时要小心。 最后一个 work-item 较少的 sub-group 可能需要特殊处理。

上一章                                         主目录    上级目录                                                              下一章

TechVideoGPU
关注
  • 0
    点赞
  • 踩
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
cute-nucleotides:Rust中SIMD向量化核苷酸的二进制编码和解码的可爱技巧
03-13
Rust中SIMD向量化核苷酸的二进制编码和解码的可爱技巧。 跑步 要运行测试，请使用带有特殊标志的货物，告知其以您的CPU为目标，以实现最高效率。 RUSTFLAGS="-C target-cpu=native" cargo test 您还可以运行基准...
native-sql-engine:适用于Spark SQL的本机SQL Engine插件，具有矢量化SIMD优化
04-19
Spark本机SQL引擎用于Spark SQL的本机引擎，具有矢量化SIMD优化在线文件您可以在找到所有Native SQL Engine文档。介绍 Spark SQL与基于行的结构化数据配合得很好。 它使用WholeStageCodeGen通过Java JIT代码来提高...
oneAPI GPU 优化指南 - 线程映射和 GPU 占用
TechVideoGPU的博客
11-09 165
SYCL 执行模型公开了 GPU 执行的抽象视图。SYCL 线程层次结构由 1、2 或 3 维 work-item 网格组成。这些工作项被分组为大小相等的线程组， 称为 work-group。work-group 中的线程进一步分为大小相等的向量组 称为 sub-group （参见下图）。Work-itemwork-item 代表 kernel 函数的并行执行的集合之一。Sub-group。
oneAPI GPU 优化指南 - 寄存器化和避免寄存器溢出
TechVideoGPU的博客
11-24 77
然而，寄存器空间有限，比内存空间小得多。如果没有足够的寄存器来保存所有变量，则寄存器可能会溢出， 或者当前在寄存器中的某些变量可以移动到内存中以腾出空间给其他变量。如果同一 sub-group 中的 work-item 共享私有 bin， 那么整个 sub-group 只需要 256 个私有 bin，或者每个 work-item 只需要 16 个私有 bin。如果 bin 的数量很大，例如大于 1024，即使私有 bin 在同一 sub-group 中共享， 也不会有足够的寄存器空间用于存储私有 bin。
oneAPI GPU 优化指南 - 移除条件检查
TechVideoGPU的博客
11-17 45
在中, 我们了解到 SIMD 分歧会对性能产生负面影响。如果 sub-group 中的所有 work-item 执行相同的指令，则 SIMD 通道得到最大程度的利用。如果一个或多个 work-item 走了不同的路径，则两条路径都必须执行才能合并。分歧是由条件检查引起的，尽管并非所有条件检查都会导致分歧。一些条件检查，即使它们不会导致 SIMD 分歧，仍然可能是性能隐患。总的来说，移除条件检查有助于提高性能。
oneAPI GPU 优化指南 - 内核（kernel）
TechVideoGPU的博客
11-17 76
你可以在一个迭代空间上提交一个 kernel，这样就可以对指定的数据对象进行计算。在 oneAPI 卸载模型中，kernel 是执行计算的单元。本节将介绍与 kernel 编程、提交和执行相关的内容。
性能优化-OpenCL kernel 开发
发狂的小花的博客
01-22 1284
OpenCL kernel 部分的介绍。 OpenCL kernel 是运行在设备端的，采用OpenCL C 语言进行开发，本文接下来首先给出一个简单的OpenCL kernel 样例，然后对OpenCL C 语言的各个部分做详细的说明，最后会给出一个完整的OpenCL程序实例，相信通过本文的学习之后大家应该可以在实际工作中使用OpenCL来优化程序的性能。
oneAPI GPU 优化指南
TechVideoGPU的博客
09-16 213
oneAPI GPU Optimization Guide 的中文版
[引擎开发] 现代图形API - metal篇
(ง •̀_•́)ง
02-02 1903
在我们对Metal进行介绍前，先来了解一下Apple GPU。metal在工作提交、显存管理上和dx12,vulkan这样的API设计比较接近，而在资源管理和绑定上的设计差异较大，这使得我们在考虑跨平台的应用时，需要充分理解API层面的设计。此外，Apple的开发者文档较为详细且完善，并且每年也有开发者大会对整个API的设计，以及一些新推出的特性进行介绍，这也是值得我们持续关注和学习的。t=MBR7。
6.CUDA编程手册中文版---附录A&B
热门推荐
专注于人工智能领域的小何尚
04-16 1万+
附录A 支持GPU设备列表 https://developer.nvidia.com/cuda-gpus 列出了所有支持 CUDA 的设备及其计算能力。 可以使用运行时查询计算能力、多处理器数量、时钟频率、设备内存总量和其他属性（参见参考手册）。 附录B 对C++扩展的详细描述 B.1 函数执行空间说明符 函数执行空间说明符表示函数是在主机上执行还是在设备上执行，以及它是可从主机调用还是从设备调用。 B.1.1 __global__ __global__ 执行空间说明符将函数声明为内核。 它的功能是： 在
体系结构复习-Part 1-指令+流水
不定期更新系统使用、编程小技巧，哔哩哔哩账号：https://space.bilibili.com/399455629，欢迎关注
06-16 462
体系结构复习-Part 1-指令+流水 1. 绪论 基本概念&性能评价 指令系统设计 流水线技术基础 存储技术基础 高级流水线 层次存储系统优化 向量处理（GPU） 互连网络 分布/并行处理 ISA与CA的差距 Old Days： ISA == CA Nowadays： CA >> ISA Architecture与Micro Architecture Architecture 对上层可见，例如Opcodes、Addressing Modes等； Micro Ar
音视频-编解码-面向非多媒体程序的SIMD向量化方法及优化技术研究.pdf
04-18
音视频-编解码-面向非多媒体程序的SIMD向量化方法及优化技术研究.pdf
generic-simd:Rust的通用SIMD抽象
04-25
通用模拟 generic-simd提供了安全且惯用的零成本抽象，用于编写显式的跨平台SIMD操作。 执照 generic-simd是根据MIT许可证和Apache许可证（2.0版）的条款发行的。 有关详细信息，请参见和 。
EasySSE:SIMD 向量数据类型和内在函数的 C++ 包装器-开源
06-29
EasySSE 是一个 C++ 类库，由 SIMD 向量数据类型和内部函数的包装器组成。 使用运算符重载、成员函数和静态函数将内部函数分组为不同的向量数据类型。 这些类提供了对 SIMD 功能的轻松自然的访问，而无需通过复杂的 ...
计算机专业毕业设计范例845篇jsp2118基于Web停车场管理系统的设计与实现_Servlet_MySql演示录像.rar
05-24
博主给大家详细整理了计算机毕业设计最新项目,对项目有任何疑问(部署跟文档),都可以问博主哦~ 一、JavaWeb管理系统毕设项目【计算机毕设选题】计算机毕业设计选题，500个热门选题推荐,更多作品展示 计算机毕业设计|PHP毕业设计|JSP毕业程序设计|Android毕业设计|Python设计论文|微信小程序设计
Windows 10 平台 FFmpeg 开发环境搭建 博客资源
05-24
【FFmpeg】Windows 10 平台 FFmpeg 开发环境搭建 ④ ( FFmpeg 开发库内容说明 | 创建并配置 FFmpeg 项目 | 拷贝 DLL 动态库到 SysWOW64 目录 ) https://hanshuliang.blog.csdn.net/article/details/139172564 博客资源 一、FFmpeg 开发库 1、FFmpeg 开发库编译 2、FFmpeg 开发库内容说明 二、创建并配置 FFmpeg 项目 1、拷贝 dll 动态库到 C:\Windows\SysWOW64 目录 - 必须操作 特别关注 2、创建 Qt 项目 - C 语言程序 3、配置 FFmpeg 开发库 - C 语言项目 4、创建并配置 FFmpeg 开发库 - C++ 项目
基于 Spring Cloud 、Spring Boot、 OAuth2 的 RBAC 企业快速开发平台
最新发布
05-24
基于 Spring Cloud 、Spring Boot、 OAuth2 的 RBAC 企业快速开发平台， 同时支持微服务架构和单体架构。提供对 Spring Authorization Server 生产级实践，支持多种安全授权模式。提供对常见容器化方案支持 Kubernetes、Rancher2 、Kubesphere、EDAS、SAE 支持
安装Pillow-SIMD直接pip install Pillow-SIMD 就可以了吗
06-12
不是的，因为Pillow-SIMD需要依赖于libjpeg-turbo库，所以在安装Pillow-SIMD之前需要先安装libjpeg-turbo库。 在Ubuntu系统上，可以通过以下命令来安装： ``` sudo apt-get install libjpeg-turbo8-dev ``` 在CentOS系统上，可以通过以下命令来安装： ``` sudo yum install libjpeg-turbo-devel ``` 安装完libjpeg-turbo库之后，再使用pip命令来安装Pillow-SIMD： ``` pip install Pillow-SIMD ``` 安装完毕后，您可以在Python中导入Pillow-SIMD库来使用它的功能。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值