Статус поддержки Apple GPU в AdaptiveCPP

0

2

В прошлый раз я писал когда поддержку смержили в мастер: В AdaptiveCpp смержили Metal backend для Apple GPU

С тех пор много изменилось. Сегодня был заапрувлен финальный PR (https://github.com/AdaptiveCpp/AdaptiveCpp/pull/2072), добавляющий более менее полную поддержку USM (unified shared memory). Теперь можно работать с косвенными структурами данных (списки, деревья) и даже запускать студ мапу внутри Apple GPU в подобном коде:


  using usm_map = std::map<int, int, std::less<int>,
                           usm_allocator<std::pair<const int, int>>>;

  usm_map* m = static_cast<usm_map*>(sycl::malloc_shared(sizeof(usm_map), q));
  new (m) usm_map(usm_allocator<std::pair<const int, int>>(q));

  for (int i = 0; i < N; ++i)
      m->insert({i * 2, i * 2 * 10});

  q.submit([&](sycl::handler& h) {
      h.parallel_for(sycl::range<1>(N), [=](sycl::id<1> idx) {
          auto it = m->find(keys[idx]);
          if (it != m->end())
              results[idx] = it->second;
      });
  }).wait();

  for (int i = 0; i < N; ++i)
      assert(i % 2 == 0 ? results[i] == i * 10 : results[i] == -1);

Как это работает. Сначала выделяем терабайт виртуальных адресов с помощью mmap, а потом используем вариант вызова newBuffer, который первым аргументом принимает cpu адрес. В этом случае мы можем контролировать отображение cpu адреса на gpu адрес и даже сделать так что адреса будут отличаться на константу (зависящую от конкретного запуска программы).

Алгоритм достаточно простой, можно посмотреть в metal_allocator.cpp.

Далее делаются LLVM пасы, которые заменяют каждый релевантный load адреса на +delta и каждый store адреса на -delta. Используется соглашение что в буферах всегда лежат cpu адреса. Сама delta просто проталкивается в ядро в рантайме.

Работает все замечательно. Я добавил десяток тестов с древовидными структурами данных. Один из них кстати выявил деградацию на AMD

Похожие темы