可能的 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 看起来，访问器中shared_ptr整个内存是无效的（m_cBuff）（不是指向的内存，实际shared_ptr 中的数据，包括计数）。怎么会这样？ BView 未被多次释放、复制、移动等

Answer 1

崩溃的原因是 BView 中的访问器无效。写入它们指向的数据会破坏访问器的内部结构，导致销毁时崩溃。如果 BView 不位于 SYCL 缓冲区中，而是分配在命令组之前的堆栈上，则相同的代码有效。删除对 m_cData 的写入可防止崩溃，但说明对 m_bData 的写入未成功完成。目前似乎不支持将访问器放在 SYCL 缓冲区内，然后在设备上访问这些缓冲区。

Answer 2

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

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

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

Possible ComputeCPP SYCL bug reading nested buffers

c++

sycl