c++ - 可能的 ComputeCPP SYCL 错误读取嵌套缓冲区

Question

我正在尝试在 GPU 上访问具有多个间接级别的数据结构。我现在拥有的示例层次结构是 A 包含 B 包含 C。每个包含数据。 A 包含指向 B 的指针，B 包含指向 C 的指针。当释放包含访问器的堆分配数据结构时，SYCL 实现会在访问器的析构函数中出现段错误。当 BView 被销毁时，就会发生段错误。

我正在为 Ubuntu 使用 ComputeCPP 实现。这似乎是运行时中的错误，因为与 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;
}

如上所述，在BView 中释放m_cAcc 时会出现段错误。这是 stacktrace从它的外观来看，正在访问的缓冲区(m_cBuff)的访问器中的shared_ptr的整个内存都是无效的(不是指向的内存，shared_ptr中的实际数据包括计数)。怎么会这样？ BView 不会被多次释放、复制、移动等。

Answer 1

很久以前我就这个领域的一些实验做了一个介绍，解释了主机和设备之间不同内存 View 的问题https://github.com/keryell/ronan/blob/gh-pages/Talks/2016/2016-03-13-PPoPP-SYCL-triSYCL/2016-03-13-PPoPP-SYCL-triSYCL-expose.pdf

更有趣的是，英特尔最近提出了解决此问题的建议，您可以查看/参与:https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/USM/USM.adoc