Open
Description
I am not sure if here is the right place but I have a problem with intel oneapi. Here is the simple matrix multiplication code I am using:
/*
Intel oneAPI DPC++
dpcpp -Qstd=c++17 /EHsc hellocl.cpp -Qtbb opencl.lib -o d.exe
Microsoft C++ Compiler
cl /EHsc /std:c++17 hellocl.cpp opencl.lib /Fe: m.exe
clang++ -std=c++17 hellocl.cpp -ltbb -lopencl -o c.exe
g++ -std=c++17 hellocl.cpp -ltbb -lopencl -o c.exe
*/
/*
1. How to use Random Number Generator
2. How to use std::vector as 2-dimensional array
3. How to suppress warning in clang compiler
4. How to use Tpf_FormatWidth, Tpf_FormatPrecision macros
dpcpp naive.cpp tbbmalloc.lib -o d.exe
*/
#if defined(__clang__)
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wpass-failed"
#endif
#define Tpf_FormatWidth 6
#define Tpf_FormatPrecision 4
#include "tpf_linear_algebra.hpp"
#include <cl/sycl.hpp>
namespace chr = tpf::chrono_random;
namespace mtx = tpf::matrix;
tpf::sstream stream;
auto& nl = tpf::nl; // single carriage-return
auto& nL = tpf::nL; // two carriage-return
auto& endl = tpf::endl; // sing carriage-return and flush out to console
auto& endL = tpf::endL; // two carriage-returns and flush out to console
void test_random_number_generator()
{
using element_t = double;
using matrix_t = mtx::scalable_fast_matrix_t<element_t>;
size_t N = 10; // number of rows
size_t M = N; // number of columns
matrix_t A{ N, M }; // N x M matrix
matrix_t B{ N, M };
// we created a random number generator
// <int> means we generator integer
// (-10, 10) means from -10 to 10, inclusive
auto generator = chr::random_generator<int>(-10, 10);
chr::random_parallel_fill(A.array(), generator);
chr::random_parallel_fill(B.array(), generator);
auto C = A * B; // matrix multiplication
stream << "A = " << nl << A << endl;
stream << "B = " << nl << B << endl;
stream << "A x B = " << nl << C << endL;
}
void test_naive_matrix_multiplication()
{
size_t N = 10;
size_t M = N;
using element_t = double;
using vectrix_t = std::vector<element_t>;
vectrix_t A(N * M);
vectrix_t B(N * M);
vectrix_t C(N * M);
vectrix_t D(N * M);
auto generator = chr::random_generator<int>(-10, 10);
chr::random_parallel_fill(A, generator);
chr::random_parallel_fill(B, generator);
auto out_A = mtx::create_formatter(A, N, M);
auto out_B = mtx::create_formatter(B, N, M);
auto out_C = mtx::create_formatter(C, N, M);
auto out_D = mtx::create_formatter(D, N, M);
stream << "A = " << nl << out_A() << endl;
stream << "B = " << nl << out_B() << endl;
auto idx_A = mtx::create_indexer(A, N, M);
auto idx_B = mtx::create_indexer(B, N, M);
auto idx_C = mtx::create_indexer(C, N, M);
for (int i = 0; i < (int)N; ++i)
{
for (int j = 0; j < (int)M; ++j)
{
for (int k = 0; k < (int)M; ++k)
idx_C(i, j) += idx_A(i, k) * idx_B(k, j); // matrix multiplication
}
}
stream << "CPU: A x B = " << nl << out_C() << endl;
sycl::queue queue{ sycl::gpu_selector{} };
sycl::buffer buf_A{ &A[0], sycl::range{N, M} };
sycl::buffer buf_B{ &B[0], sycl::range{N, M} };
sycl::buffer buf_D{ &D[0], sycl::range{N, M} };
queue.submit([&](sycl::handler& cgh)
{
auto a = buf_A.get_access<sycl::access::mode::read>(cgh);
auto b = buf_B.get_access<sycl::access::mode::read>(cgh);
auto d = buf_D.get_access<sycl::access::mode::read_write>(cgh);
constexpr int tile_size = 16;
sycl::local_accessor<int> tileA{tile_size, cgh};
cgh.parallel_for(
sycl::nd_range<2>{{N, N}, {1, tile_size}}, [=](sycl::nd_item<2> it) {
// Indices in the global index space:
int m = it.get_global_id()[0];
int n = it.get_global_id()[1];
// Index in the local index space:
int i = it.get_local_id()[1];
size_t sum = 0;
for (int kk = 0; kk < 496; kk += tile_size) {
// Load the matrix tile from matrix A, and synchronize
// to ensure all work-items have a consistent view
// of the matrix tile in local memory.
tileA[i] = a[m][kk + i];
it.barrier();
// Perform computation using the local memory tile, and
// matrix B in global memory.
for (int k = 0; k < tile_size; k++)
sum += tileA[k] * b[kk + k][n];
// After computation, synchronize again, to ensure all
// reads from the local memory tile are complete.
it.barrier();
}
// Write the final result to global memory.
d[m][n] = sum;
});
});
// when this block goes off,
// the destructor of buf_D waits until it is released by the queue
// and copies to its host memory D
stream << "GPU: A x B = " << nl << out_D() << endl;
}
#if defined(__clang__)
#pragma clang diagnostic pop
#endif
int main()
{
// test_random_number_generator();
test_naive_matrix_multiplication();
}
Here is the error I am getting: error: no template named 'local_accessor' in namespace 'sycl'; did you mean 'host_accessor'?
sycl::local_accessor tileA{tile_size, cgh};
I am using intel oneapi dpc++ compiler.
Is this a wrong way to implement local_accessor or intel oneapi SYCL implementation? or it does not support it?
dpcpp -Qstd=c++17 /EHsc hellocl.cpp -Qtbb opencl.lib -o d.exe this is the comment that I am using to produce the executable.
OS: windows
target device is intel GPU
dpc++ version : Intel(R) oneAPI DPC++/C++ Compiler 2021.3.0 (2021.3.0.20210619)
Target: x86_64-pc-windows-msvc
Thread model: posix
InstalledDir: C:\Program Files (x86)\Intel\oneAPI\compiler\latest\windows\bin