Умножение матриц в SYCL с использованием nd_range - PullRequest
Новая
       69

Умножение матриц в SYCL с использованием nd_range

0 голосов
/ 24 октября 2019

Я делаю матричное умножение в sycl, но у меня есть некоторые проблемы. Я использую 2 (4x4) матрицы для умножения и на первой итерации цикла for она работает , но на второй итерации, когда i = 1, она работает нормально до C [11] = A [11] *B [15] , но затем пропускает 1 умножение и движется вперед. Я знаю проблему, почему она пропускает, но, к сожалению, я не смог правильно изменить индекс матрицы B. Пожалуйста, если кто-то может помочь, я буду очень признателен. Спасибо

Вот код Matsize = 4, Blocksize = 4 также я знаю, что цикл будет равен matsize, это 2, просто чтобы получить ясное представление о потоке выполнения 

{
    range<1> dimensions(matSize * matSize);
    const property_list props = { property::buffer::use_host_ptr() };
    buffer<T> A_buf(MA, dimensions, props);
    buffer<T> B_buf(MB, dimensions, props);
    buffer<T> C_buf(MC, dimensions, props);

    myQueue.submit([&](handler& cgh) {
        auto A_ptr = A_buf.template get_access<access::mode::read>(cgh);
        auto B_ptr = B_buf.template get_access<access::mode::read_write>(cgh);
        auto C_ptr = C_buf.template get_access<access::mode::write>(cgh);
        auto localRange = range<1>(blockSize* blockSize);

        accessor<T, 1, access::mode::read_write, access::target::local>
            C(matSize * matSize, cgh);

        cgh.parallel_for<mxm_kernel>(
            nd_range<2>(range<2>{matSize, matSize},   
                range<2>{blockSize, blockSize}),  
            [=](nd_item<2> item) {

            const auto id_x = item.get_global_id(0);
            const auto id_y = item.get_global_id(1);

            const auto width = item.get_group_range(0) * item.get_local_range(0);

            const auto index = id_x * width + id_y;
            const auto index2 = id_y * width + id_x;

            for (int i = 0; i < 2 ; i++) {

                C[index] += A_ptr[index] * B_ptr[index2 + i ];
            }
            out << "C is!" << C[index] << sycl::endl;

            item.barrier(cl::sycl::access::fence_space::local_space);
            C_ptr[index] = C[index];
        });
    });
}
...