Wrong results of a simple kernel

[gogo2]

I'm running some simple test with recentunified/next and have encountered what's described below. Last time I run those tests was about a year ago and, as far as I remember, results were correct (I believe it was still on opencl flow)

Describe the bug At the end of the program vec_c[].c should contain the sequence 70 74 78 ... 194, but instead it is 78 82 ... 194 0 0 When I comment anything in the kernel code other than c_acc[id].x = ... (for loop, other assignments), the results in vec_c[].c become OK. I didn't check if other values are correct.

To Reproduce

#include <iostream>
#include <array>
#include <CL/sycl.hpp>
#include "device_selectors.hpp"

using namespace cl;

struct TestStruct {
    float y{}, z{};
    uint64_t x{}, a{}, b{};
    std::array<std::pair<double, uint64_t>, 8> data{};
};

bool check(const std::vector<TestStruct> &vec) {
    for (auto i = 0; i < vec.size(); ++i) {
        std::cout << vec[i].x << ' ' << 4 * i + 2 * vec.size() + 2 + 4 << std::endl;
    }
    return true;
}

int main(int argc, char *argv[]) {
    std::string device_search_guide = "xilinx";
    std::size_t size = 32;

    std::cout << "test size: " << size << std::endl;
    std::cout << "device search guide: " << device_search_guide << std::endl;

    sycl::queue queue{faisycl::name_selector{device_search_guide}};

    std::cout << "SYCL device: " << queue.get_device().get_info<sycl::info::device::name>() << std::endl;

    std::vector<TestStruct> vec_a{}, vec_b{}, vec_c{};
    vec_a.resize(size);
    vec_b.resize(size);
    vec_c.resize(size);

    for (auto i = 0; i < vec_a.size(); ++i) {
        vec_a[i].a = i + 1;
        vec_b[i].a = i + 2;
        vec_a[i].b = vec_a[i].a + size;
        vec_b[i].b = vec_b[i].a + size;
    }

    sycl::range<1> rng{N};
    {
        sycl::buffer<TestStruct, 1> a_buff(vec_a.data(), rng);
        sycl::buffer<TestStruct, 1> b_buff(vec_b.data(), rng);
        sycl::buffer<TestStruct, 1> c_buff(vec_c.data(), rng);

        queue.submit([&](sycl::handler &cgh) {
            const auto a_acc = a_buff.get_access<sycl::access::mode::read>(cgh);
            const auto b_acc = b_buff.get_access<sycl::access::mode::read>(cgh);
            auto c_acc = c_buff.get_access<sycl::access::mode::write>(cgh);

            auto kernel = [=](sycl::id<1> id) {
                c_acc[id].x = a_acc[id].a + a_acc[id].b + b_acc[id].a + b_acc[id].b;
                c_acc[id].y = sycl::sqrt(static_cast<float>(a_acc[id].a * b_acc[id].b));
                c_acc[id].z = sycl::fabs(static_cast<float>(a_acc[id].a - b_acc[id].b));

                for (auto i = 0; i < c_acc[id].data.size(); ++i) {
                    c_acc[id].data[i].first = c_acc[id].y * c_acc[id].y * c_acc[id].y;
                    c_acc[id].data[i].second = c_acc[id].x * 2 * i;
                }

                c_acc[id].y = 3.f;
                c_acc[id].z = 5.f;
            };

            cgh.parallel_for<class TestStructorAdd>(rng, kernel);
        });
    }
    check(vec_c);

    return 0;
}

Environment

• Ubuntu 22.04
• Vitis 2022.1.2
• target platform: xilinx_u280_gen3x16_xdma_1_202211_1
• xrt at https://github.com/Xilinx/XRT/tree/836a1f4e39f04d989b7709f3c41d337e8a073197
• sycl at https://github.com/triSYCL/sycl/tree/ec1938ff0593ed06a330d272f03101eae88d2bc0
• hls flow, hw_emu mode

[keryell]

Curious.

[Ralender]

After some investigation I think it is an issue with the HLS backend. I filled an internal issue about it. But the SYCL implementation is not using the HLS backend in a supported way. So I don't have a clear idea of whether they will fix it or when the will fix it.

[gogo2]

After some investigation I think it is an issue with the HLS backend. I filled an internal issue about it. But the SYCL implementation is not using the HLS backend in a supported way. So I don't have a clear idea of whether they will fix it or when the will fix it.

If this was HLS issue, it seems to be fixed in 2022.2, at least for this test case.