Возможная ошибка ComputeCPP SYCL при чтении вложенных буферов

Question

Я пытаюсь получить доступ к структуре данных с несколькими уровнями косвенности на GPU. Примерная иерархия, которая у меня сейчас есть: A содержит B, содержит C. Каждая содержит данные. A содержит указатель на B, B содержит указатель на C. Когда выделенная в куче структура данных, содержащая методы доступа, освобождается, в реализации деструкторов SYCL возникают ошибки в деструкторах методов доступа. Когда BView уничтожен, происходит сбой.

Я использую реализацию ComputeCPP для Ubuntu. Это похоже на ошибку во время выполнения, потому что буферы, связанные с аксессорами в BView все еще действительны на момент уничтожения BView. Других ошибок нет. Я тоже пробовал течь BView обойти ошибку. Однако, поскольку методы доступа BView содержат ссылку на буферы для B и C, возникает тупиковая ситуация. Это также означает, что буферы, на которые ссылаются методы доступа BView, действительны. Является ли нарушением спецификации SYCL выделение динамической памяти аксессорам или буферам? Может быть, это может быть причиной проблем, так как AView освобождает без проблем.

#include "SYCL/sycl.hpp"
#include <vector>
#include <utility>
#include <iostream>
#include <memory>

struct C {
    int m_cData;

    C() : m_cData(0) {}

    ~C() {
        std::cout << "C deallocating" << std::endl;
    }
};

struct B {
    int m_bData;
    std::shared_ptr<C> m_c;

    B() : m_bData(0), m_c(std::make_shared<C>()) {}

    ~B() {
        std::cout << "B deallocating" << std::endl;
    }
};

struct BBuff {
    cl::sycl::buffer<B> m_bBuff;
    cl::sycl::buffer<C> m_cBuff;

    BBuff(const std::shared_ptr<B>& b) : m_bBuff(b, cl::sycl::range<1>(1)),
                                  m_cBuff(b->m_c, cl::sycl::range<1>(1)) {}

    ~BBuff() {
        std::cout << "BBuff deallocating" << std::endl;
    }
};

template<cl::sycl::access::target target>
struct BView
{
    cl::sycl::accessor<B, 1, cl::sycl::access::mode::read_write, target,
            cl::sycl::access::placeholder::true_t> m_bDataAcc;
    cl::sycl::accessor<C, 1, cl::sycl::access::mode::read_write, target,
            cl::sycl::access::placeholder::true_t> m_cAcc;

    BView(const std::shared_ptr<BBuff>& bBuff) : m_bDataAcc(bBuff->m_bBuff), m_cAcc(bBuff->m_cBuff)
    {
    }

    void RequireForHandler(cl::sycl::handler& cgh) {
        cgh.require(m_bDataAcc);
        cgh.require(m_cAcc);
    }

    ~BView()
    {
        std::cout << "BView deallocating" << std::endl;
    }
};

struct A {
    int m_aData;
    std::shared_ptr<B> m_b;

    A() : m_aData(0), m_b(std::make_shared<B>()) {}

    ~A()
    {
        std::cout << "A deallocating" << std::endl;
    }
};

template<cl::sycl::access::target target>
struct ABuff {
    cl::sycl::buffer<A> m_aBuff;
    std::shared_ptr<BBuff> m_bBuff;

    std::shared_ptr<BView<target>> m_bViewBuffData;
    std::shared_ptr<cl::sycl::buffer<BView<target>>> m_bViewBuff;

    ABuff(const std::shared_ptr<A>& a): m_aBuff(a, cl::sycl::range<1>(1)),
        m_bBuff(std::make_shared<BBuff>(a->m_b)) {
        m_bViewBuffData = std::make_shared<BView<target>>(m_bBuff);
        m_bViewBuff = std::make_shared<cl::sycl::buffer<BView<target>>>(m_bViewBuffData, cl::sycl::range<1>(1));
    }

    ~ABuff()
    {
        std::cout << "ABuff deallocating" << std::endl;
    }
};

template<cl::sycl::access::target target>
struct AView {
    cl::sycl::accessor<BView<target>, 1, cl::sycl::access::mode::read_write, target,
            cl::sycl::access::placeholder::true_t> m_bAcc;
    cl::sycl::accessor<A, 1, cl::sycl::access::mode::read_write, target,
            cl::sycl::access::placeholder::true_t> m_aDataAcc;

    ABuff<target>* m_aBuff;
    AView(ABuff<target>* aBuff): m_bAcc(*aBuff->m_bViewBuff), m_aDataAcc(aBuff->m_aBuff),
                                                 m_aBuff(aBuff) {}

    void RequireForHandler(cl::sycl::handler& cgh) {
        m_aBuff->m_bViewBuffData->RequireForHandler(cgh);
        cgh.require(m_bAcc);
        cgh.require(m_aDataAcc);
    }
};

class init_first_block;

int main(int argc, char** argv)
{
    std::shared_ptr<A> a = std::make_shared<A>();
    try
    {
        cl::sycl::queue workQueue;
        ABuff<cl::sycl::access::target::global_buffer> aGlobalBuff(a);
        AView<cl::sycl::access::target::global_buffer> aAccDevice(&aGlobalBuff);
        workQueue.submit([&aAccDevice](cl::sycl::handler &cgh) {
            aAccDevice.RequireForHandler(cgh);

            cgh.single_task<class init_first_block>([aAccDevice]() {
                aAccDevice.m_aDataAcc[0].m_aData = 1;
                aAccDevice.m_bAcc[0].m_bDataAcc[0].m_bData = 2;
                aAccDevice.m_bAcc[0].m_cAcc[0].m_cData = 3;
            });
        });

        workQueue.wait();
    }
    catch (...)
    {
        std::cout << "Failure running nested accessor test" << std::endl;
    }

    std::cout << "A data: " << a->m_aData << std::endl;
    std::cout << "B data: " << a->m_b->m_bData << std::endl;
    std::cout << "C data: " << a->m_b->m_c->m_cData << std::endl;

    return 0;
}

Как уже упоминалось выше, при освобождении происходит ошибка m_cAcc в BView, Вот трассировка стека. Судя по всему, вся память shared_ptr в методе доступа к буферу, к которому осуществляется доступ (m_cBuff), недопустима (не указанная память, фактические данные в shared_ptr, включая счетчик). Как это может быть? BView не освобождается несколько раз, копируется, перемещается и т. д.

Answer 1

Я давным-давно выступил с презентацией о некоторых экспериментах в этой области, в которой объясняется проблема различных представлений памяти между хостом и устройством https://github.com/keryell/ronan/blob/gh-pages/Talks/2016/2016-03-13-PPoPP-SYCL-triSYCL/2016-03-13-PPoPP-SYCL-triSYCL-expose.pdf

Что еще интереснее, у Intel есть недавнее предложение по решению этой проблемы, на которое вы можете посмотреть / внести свой вклад: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/USM/USM.adoc

Answer 2

Причиной сбоя было то, что средства доступа в BView не действительны. Запись данных, на которые они указывают, повредила внутреннюю часть средства доступа, что привело к сбою при разрушении. Тот же код работает, если BView не находится в буфере SYCL, а вместо этого расположен в стеке перед группой команд. Удаление записи в m_cData предотвращает сбой, но показывает, что запись в m_bData не проходит успешно. Кажется, что размещение аксессоров внутри буферов SYCL, к которым затем осуществляется доступ к устройству, на данный момент не поддерживается.