Я новичок в Sycl. Следующий код компилируется нормально, но дает неправильный результат.
Код вычисляет продукт двух матриц (a, r m x k и b, r k x n ) в матрицу результатов (c, r m x n ).
#include
#include
constexpr int M = 17,
K = 18,
N = 19,
tile_size = 9;
int main( void )
{
// Initialize matrices to be zero and "1" in the diagonal.
std::array< int, M * K > a;
for( int i = 0; i < M * K; ++i ) a[i] = 0;
for( int i = 0; i < (M < K ? M : K); ++i ) a[i * (M < K ? K + 1 : M + 1)] = 1;
std::array< int, K * N > b;
for( int i = 0; i < K * N; ++i ) b[i] = 0;
for( int i = 0; i < (K < N ? K : N); ++i ) b[i * (K < N ? N + 1 : K + 1)] = 1;
// Result matrix all zeros for now.
std::array< int, M * N > c;
for( int i = 0; i < M * N; ++i ) c[i] = 0;
sycl::queue q;
sycl::buffer a_buf{ a },
b_buf{ b },
c_buf{ c };
q.submit( [&]( sycl::handler& h )
{
sycl::accessor matrix_a{ a_buf, h, sycl::read_only };
sycl::accessor matrix_b{ b_buf, h, sycl::read_only };
sycl::accessor matrix_c{ c_buf, h, sycl::write_only, sycl::no_init };
h.parallel_for( sycl::nd_range< 2 >{ { M, N }, { 1, tile_size } },
[=]( sycl::nd_item< 2 > item )
{
auto sg = item.get_sub_group();
const int m = item.get_global_id()[0],
n = item.get_global_id()[1];
int i = item.get_local_id()[1];
int sum = 0;
for( int k = 0; k < K; k += tile_size )
{
int tile_a = matrix_a[m * K + k + i];
for( int t = 0; t < tile_size; ++t )
sum += sycl::group_broadcast( sg, tile_a, t ) *
matrix_b[(k + t) * N + n];
}
matrix_c[m * N + n] = sum;
}
);
}
);
sycl::host_accessor host_c{ c_buf, sycl::read_only };
for( int i = 0; i < N * M; ++i )
c[i] = host_c[i];
std::cout
Подробнее здесь: [url]https://stackoverflow.com/questions/79461715/matrix-multiplication-with-sycl-sub-groups-goes-wrong[/url]
Я новичок в Sycl. Следующий код компилируется нормально, но дает неправильный результат. Код вычисляет продукт двух матриц (a, r m x k и b, r k x n ) в матрицу результатов (c, r m x n ). [code]#include #include
constexpr int M = 17, K = 18, N = 19, tile_size = 9;
int main( void ) { // Initialize matrices to be zero and "1" in the diagonal. std::array< int, M * K > a; for( int i = 0; i < M * K; ++i ) a[i] = 0; for( int i = 0; i < (M < K ? M : K); ++i ) a[i * (M < K ? K + 1 : M + 1)] = 1; std::array< int, K * N > b; for( int i = 0; i < K * N; ++i ) b[i] = 0; for( int i = 0; i < (K < N ? K : N); ++i ) b[i * (K < N ? N + 1 : K + 1)] = 1; // Result matrix all zeros for now. std::array< int, M * N > c; for( int i = 0; i < M * N; ++i ) c[i] = 0;
h.parallel_for( sycl::nd_range< 2 >{ { M, N }, { 1, tile_size } }, [=]( sycl::nd_item< 2 > item ) { auto sg = item.get_sub_group(); const int m = item.get_global_id()[0], n = item.get_global_id()[1]; int i = item.get_local_id()[1]; int sum = 0; for( int k = 0; k < K; k += tile_size ) { int tile_a = matrix_a[m * K + k + i];
for( int t = 0; t < tile_size; ++t ) sum += sycl::group_broadcast( sg, tile_a, t ) * matrix_b[(k + t) * N + n]; } matrix_c[m * N + n] = sum; } ); } );
sycl::host_accessor host_c{ c_buf, sycl::read_only }; for( int i = 0; i < N * M; ++i ) c[i] = host_c[i];