Возникли проблемы с созданием буферов данных для настраиваемых объектов oneAPI

Question

Я новичок в oneAPI и подобных фреймворках, поэтому у меня возникают проблемы с управлением данными с использованием буферов данных SYCL.

Моя задача - найти подстроки в заданной строке с помощью алгоритма Ахо-Корасика.

Моя идея заключалась в том, чтобы создать дерево-дерево и после этого отправить ядро, которое параллельно будет находить в нем подстроки. Поэтому для этого я создал очередь SYCL, создал буферы для строки (для поиска подстроки), для вектора (для хранения результата поиска) и для моего объекта Aho-Corasick, который содержит корень ранее созданного дерева. Однако насчет последнего я не уверен, поскольку я создаю буфер для объекта в памяти хоста, который содержит указатели на другие объекты (например, узлы, которые содержат указатели на другие узлы).

Структура объекта Node:

class Node {

        typedef Node *node_ptr;

    private:

        std::set<std::pair<int, std::string>> retVals;
        std::unordered_map<char, node_ptr> children;
        node_ptr fail;
        char value;

Это метод поиска:

 void
        matchWords(char *text, int startIdx, int endIdx,  cl::sycl::cl_int *matched) {

            node_ptr child = start;
            int item = startIdx;
            for (int i = startIdx; i < endIdx; ++i) {
                child = child->nextNode(text[i]);
                if (child == nullptr) {
                    child = start;
                    continue;
                }
                for (const auto &returns: child->getRetVals()) {
                    matched[item++] = returns.first;
                    if (item == endIdx) item = startIdx;
                }
            }
        }

Буферы:

cl::sycl::buffer<char, 1> fasta_buf(tempFasta.data(), cl::sycl::range<1>(len));
cl::sycl::buffer<cl::sycl::cl_int, 1> vec_buf(vec.data(), cl::sycl::range<1>(len));
cl::sycl::buffer<aho_corasick::AhoCorasick, 1> aho_buf(a, cl::sycl::range<1>(1));

и очередь очереди:

q.submit([&](cl::sycl::handler &cgh) {
        auto string_acc = fasta_buf.get_access<cl::sycl::access::mode::read>(cgh);
        auto vec_acc = vec_buf.get_access<cl::sycl::access::mode::read_write>(cgh);
        auto aho_acc = aho_buf.get_access<cl::sycl::access::mode::read>(cgh);

        cgh.parallel_for<class dummy>(
                cl::sycl::range<1>(10), [=](cl::sycl::item<1> i) {
                    // 10 is the number of workers I want 
                    int startInx = (int) (i.get_linear_id() * (len / 10)); 
                    int endInx = (int) ((i.get_linear_id() + 1) * (len / 10));
                    aho_acc.get_pointer()->matchWords(string_acc.get_pointer(), startInx, endInx, vec_acc.get_pointer());
                });
    });
    q.wait_and_throw();

Я выяснил, что программа не работает после попытки получить доступ к элементам карты детей. Таким образом, я думаю, что проблема в том, что указатели, хранящиеся в карте, являются указателями на память хоста, к которой устройство не имеет доступа.

Answer 1

Если я правильно понял, вы пытаетесь использовать std::unordered_map, std::string а также std::setв коде устройства. Я не эксперт по расширениям oneAPI SYCL для Intel, но в чистом SYCL 1.2.1 это запрещено, и я был бы удивлен, если бы это сработало в DPC++.

Спецификация SYCL 1.2.1 на самом деле не определяет, как SYCL взаимодействует со стандартной библиотекой. Хотя некоторые реализации могут дать некоторые гарантии относительно определенных четко определенных частей стандартной библиотеки, работающих в коде разработчика как расширение (обычно, например,std::математические функции), это не гарантируется во всех реализациях SYCL. Кроме того, поддержка контейнеров STL в коде устройства (что не требуется спецификацией SYCL), на мой взгляд, будет особенно сложной задачей, и я никогда не слышал о реализации SYCL, поддерживающей это. Это связано с тем, что контейнеры обычно используют механизмы, не поддерживаемые в коде устройства SYCL, потому что они требуют поддержки во время выполнения, например, выдачи исключений. Поскольку, скажем, на GPU нет среды выполнения C++, такие механизмы не могут работать в SYCL.

Также важно понимать, что на самом деле это не ограничение, специфичное для SYCL, а общее ограничение среди моделей гетерогенного программирования. Другие модели гетерогенного программирования, такие как CUDA, налагают аналогичные ограничения по тем же причинам.

Другая трудность с контейнерами в ядрах состоит в том, что структуры данных STL обычно не предназначены для массивно-параллельной модели выполнения SIMT на устройстве SYCL, что делает их подверженными условиям гонки.

Последняя проблема, которую вы уже определили: вы копируете указатели в память хоста. Поскольку вы используете oneAPI DPC++, самым простым решением для работы со структурами данных на основе указателей является использование расширения единой разделяемой памяти (USM) Intel SYCL, которое можно использовать для создания указателей, действующих как на хосте, так и на устройстве. Также существует распределитель USM, который можно передать контейнерам, если они поддерживаются в коде устройства.