Статус поддержки Apple GPU в AdaptiveCPP
В прошлый раз я писал когда поддержку смержили в мастер: В 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