Update Intel SYCL compiler driver. Update device selectors and access…

…ors to SYCL2020. (ParRes#629) Signed-off-by: James Brodman <[email protected]>
s-sajid-ali · Jul 19, 2023 · 9272001 · 9272001
1 parent 252bbb5
commit 9272001
Show file tree

Hide file tree

Showing 23 changed files with 80 additions and 294 deletions.
diff --git a/Cxx11/dgemm-onemkl.cc b/Cxx11/dgemm-onemkl.cc
@@ -126,7 +126,7 @@ int main(int argc, char * argv[])
   }
   std::cout << "Input copy           = " << (input_copy ? "yes" : "no") << std::endl;
 
-  sycl::queue q(sycl::default_selector{});
+  sycl::queue q(sycl::default_selector_v);
   prk::SYCL::print_device_platform(q);
 
   //////////////////////////////////////////////////////////////////////

diff --git a/Cxx11/dgemm-sycl.cc b/Cxx11/dgemm-sycl.cc
@@ -73,9 +73,9 @@ void prk_dgemm(sycl::queue & q,
 {
     q.submit([&](sycl::handler& h) {
 
-      auto A = d_A.get_access<sycl::access::mode::read>(h);
-      auto B = d_B.get_access<sycl::access::mode::read>(h);
-      auto C = d_C.get_access<sycl::access::mode::read_write>(h);
+      sycl::accessor A(d_A, h, sycl::read_only);
+      sycl::accessor B(d_B, h, sycl::read_only);
+      sycl::accessor C(d_C, h);
 
       h.parallel_for<class dgemm>( sycl::range<2>{order,order}, [=] (sycl::id<2> it) {
 
@@ -130,7 +130,7 @@ int main(int argc, char * argv[])
     return 1;
   }
 
-  sycl::queue q(sycl::default_selector{});
+  sycl::queue q(sycl::default_selector_v);
   prk::SYCL::print_device_platform(q);
 
   if (tile_size < order) {

diff --git a/Cxx11/generate-sycl-stencil.py b/Cxx11/generate-sycl-stencil.py
@@ -26,8 +26,8 @@ def codegen(src,pattern,stencil_size,radius,model,dim,usm):
     src.write('{\n')
     src.write('  q.submit([&](sycl::handler& h) {\n')
     if (not usm):
-        src.write('    auto in  = d_in.template get_access<sycl::access::mode::read>(h);\n')
-        src.write('    auto out = d_out.template get_access<sycl::access::mode::read_write>(h);\n')
+        src.write('    sycl::accessor in(d_in, h, sycl::read_only);\n')
+        src.write('    sycl::accessor out(d_out, h);\n')
     if (dim==2):
         for r in range(1,radius+1):
             src.write('    sycl::id<2> dx'+str(r)+'(sycl::range<2> {'+str(r)+',0});\n')

diff --git a/Cxx11/nstream-dpcpp.cc b/Cxx11/nstream-dpcpp.cc
@@ -106,7 +106,7 @@ int main(int argc, char * argv[])
   std::cout << "Vector length        = " << length << std::endl;
   std::cout << "Block size           = " << block_size << std::endl;
 
-  sycl::queue q(sycl::default_selector{});
+  sycl::queue q(sycl::default_selector_v);
   prk::SYCL::print_device_platform(q);
 
   size_t padded_length = block_size * prk::divceil(length,block_size);

diff --git a/Cxx11/nstream-onedpl.cc b/Cxx11/nstream-onedpl.cc
@@ -101,7 +101,7 @@ int main(int argc, char *argv[]) {
   std::cout << "Number of iterations = " << iterations << std::endl;
   std::cout << "Vector length        = " << length << std::endl;
 
-  sycl::queue q(sycl::default_selector{});
+  sycl::queue q(sycl::default_selector_v);
   prk::SYCL::print_device_platform(q);
 
   //////////////////////////////////////////////////////////////////////

diff --git a/Cxx11/nstream-onemkl.cc b/Cxx11/nstream-onemkl.cc
@@ -106,7 +106,7 @@ int main(int argc, char * argv[])
   std::cout << "Number of iterations = " << iterations << std::endl;
   std::cout << "Vector length        = " << length << std::endl;
 
-  sycl::queue q(sycl::default_selector{}, sycl::property::queue::in_order{});
+  sycl::queue q(sycl::default_selector_v, sycl::property::queue::in_order{});
 
   //////////////////////////////////////////////////////////////////////
   // Allocate space and perform the computation

diff --git a/Cxx11/nstream-sycl-explicit-usm.cc b/Cxx11/nstream-sycl-explicit-usm.cc
@@ -275,7 +275,7 @@ int main(int argc, char * argv[])
   //////////////////////////////////////////////////////////////////////
 
   try {
-    sycl::queue q{sycl::host_selector{}};
+    sycl::queue q{sycl::cpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE
@@ -294,26 +294,7 @@ int main(int argc, char * argv[])
   }
 
   try {
-    sycl::queue q{sycl::cpu_selector{}};
-    prk::SYCL::print_device_platform(q);
-    run<float>(q, iterations, length, block_size);
-#ifndef DPCPP_NO_DOUBLE
-    run<double>(q, iterations, length, block_size);
-#endif
-  }
-  catch (sycl::exception & e) {
-    std::cout << e.what() << std::endl;
-    prk::SYCL::print_exception_details(e);
-  }
-  catch (std::exception & e) {
-    std::cout << e.what() << std::endl;
-  }
-  catch (const char * e) {
-    std::cout << e << std::endl;
-  }
-
-  try {
-    sycl::queue q{sycl::gpu_selector{}};
+    sycl::queue q{sycl::gpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE

diff --git a/Cxx11/nstream-sycl-explicit.cc b/Cxx11/nstream-sycl-explicit.cc
@@ -100,15 +100,15 @@ void run(sycl::queue & q, int iterations, size_t length, size_t block_size)
     sycl::buffer<T> d_C { sycl::range<1>{length} };
 
     q.submit([&](sycl::handler& h) {
-        sycl::accessor<T, 1, sycl::access::mode::discard_write, sycl::access::target::global_buffer> A(d_A, h, sycl::range<1>(length), sycl::id<1>(0));
+        sycl::accessor A(d_A, h, sycl::no_init);
         h.fill(A,(T)0);
     });
     q.submit([&](sycl::handler& h) {
-        sycl::accessor<T, 1, sycl::access::mode::discard_write, sycl::access::target::global_buffer> B(d_B, h, sycl::range<1>(length), sycl::id<1>(0));
+        sycl::accessor B(d_B, h, sycl::no_init);
         h.fill(B,(T)2);
     });
     q.submit([&](sycl::handler& h) {
-        sycl::accessor<T, 1, sycl::access::mode::discard_write, sycl::access::target::global_buffer> C(d_C, h, sycl::range<1>(length), sycl::id<1>(0));
+        sycl::accessor C(d_C, h, sycl::no_init);
         h.fill(C,(T)2);
     });
     q.wait();
@@ -118,10 +118,9 @@ void run(sycl::queue & q, int iterations, size_t length, size_t block_size)
       if (iter==1) nstream_time = prk::wtime();
 
       q.submit([&](sycl::handler& h) {
-
-        auto A = d_A.template get_access<sycl::access::mode::read_write>(h);
-        auto B = d_B.template get_access<sycl::access::mode::read>(h);
-        auto C = d_C.template get_access<sycl::access::mode::read>(h);
+        sycl::accessor A(d_A, h);
+        sycl::accessor B(d_B, h, sycl::read_only);
+        sycl::accessor C(d_C, h, sycl::read_only);
 
         if (block_size == 0) {
             // hipSYCL prefers range to nd_range because no barriers
@@ -164,7 +163,7 @@ void run(sycl::queue & q, int iterations, size_t length, size_t block_size)
     nstream_time = prk::wtime() - nstream_time;
 
     q.submit([&](sycl::handler& h) {
-        sycl::accessor<T, 1, sycl::access::mode::read, sycl::access::target::global_buffer> A(d_A, h, sycl::range<1>(length), sycl::id<1>(0));
+        sycl::accessor A(d_A, h, sycl::read_only);
         h.copy(A,h_A.data());
     });
     q.wait();
@@ -268,26 +267,7 @@ int main(int argc, char * argv[])
   //////////////////////////////////////////////////////////////////////
 
   try {
-    sycl::queue q{sycl::host_selector{}};
-    prk::SYCL::print_device_platform(q);
-    run<float>(q, iterations, length, block_size);
-#ifndef DPCPP_NO_DOUBLE
-    run<double>(q, iterations, length, block_size);
-#endif
-  }
-  catch (sycl::exception & e) {
-    std::cout << e.what() << std::endl;
-    prk::SYCL::print_exception_details(e);
-  }
-  catch (std::exception & e) {
-    std::cout << e.what() << std::endl;
-  }
-  catch (const char * e) {
-    std::cout << e << std::endl;
-  }
-
-  try {
-    sycl::queue q{sycl::cpu_selector{}};
+    sycl::queue q{sycl::cpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE
@@ -306,7 +286,7 @@ int main(int argc, char * argv[])
   }
 
   try {
-    sycl::queue q{sycl::gpu_selector{}};
+    sycl::queue q{sycl::gpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE

diff --git a/Cxx11/nstream-sycl-usm.cc b/Cxx11/nstream-sycl-usm.cc
@@ -253,7 +253,7 @@ int main(int argc, char * argv[])
   //////////////////////////////////////////////////////////////////////
 
   try {
-    sycl::queue q{sycl::host_selector{}};
+    sycl::queue q{sycl::cpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE
@@ -272,26 +272,7 @@ int main(int argc, char * argv[])
   }
 
   try {
-    sycl::queue q{sycl::cpu_selector{}};
-    prk::SYCL::print_device_platform(q);
-    run<float>(q, iterations, length, block_size);
-#ifndef DPCPP_NO_DOUBLE
-    run<double>(q, iterations, length, block_size);
-#endif
-  }
-  catch (sycl::exception & e) {
-    std::cout << e.what() << std::endl;
-    prk::SYCL::print_exception_details(e);
-  }
-  catch (std::exception & e) {
-    std::cout << e.what() << std::endl;
-  }
-  catch (const char * e) {
-    std::cout << e << std::endl;
-  }
-
-  try {
-    sycl::queue q{sycl::gpu_selector{}};
+    sycl::queue q{sycl::gpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE

diff --git a/Cxx11/nstream-sycl.cc b/Cxx11/nstream-sycl.cc
@@ -106,10 +106,9 @@ void run(sycl::queue & q, int iterations, size_t length, size_t block_size)
       if (iter==1) nstream_time = prk::wtime();
 
       q.submit([&](sycl::handler& h) {
-
-        auto A = d_A.template get_access<sycl::access::mode::read_write>(h);
-        auto B = d_B.template get_access<sycl::access::mode::read>(h);
-        auto C = d_C.template get_access<sycl::access::mode::read>(h);
+        sycl::accessor A(d_A, h);
+        sycl::accessor B(d_B, h, sycl::read_only);
+        sycl::accessor C(d_C, h, sycl::read_only);
 
         if (block_size == 0) {
             // hipSYCL prefers range to nd_range because no barriers
@@ -250,26 +249,7 @@ int main(int argc, char * argv[])
   //////////////////////////////////////////////////////////////////////
 
   try {
-    sycl::queue q{sycl::host_selector{}};
-    prk::SYCL::print_device_platform(q);
-    run<float>(q, iterations, length, block_size);
-#ifndef DPCPP_NO_DOUBLE
-    run<double>(q, iterations, length, block_size);
-#endif
-  }
-  catch (sycl::exception & e) {
-    std::cout << e.what() << std::endl;
-    prk::SYCL::print_exception_details(e);
-  }
-  catch (std::exception & e) {
-    std::cout << e.what() << std::endl;
-  }
-  catch (const char * e) {
-    std::cout << e << std::endl;
-  }
-
-  try {
-    sycl::queue q{sycl::cpu_selector{}};
+    sycl::queue q{sycl::cpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE
@@ -288,7 +268,7 @@ int main(int argc, char * argv[])
   }
 
   try {
-    sycl::queue q{sycl::gpu_selector{}};
+    sycl::queue q{sycl::gpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, length, block_size);
 #ifndef DPCPP_NO_DOUBLE

diff --git a/Cxx11/p2p-hyperplane-sycl.cc b/Cxx11/p2p-hyperplane-sycl.cc
@@ -148,7 +148,7 @@ int main(int argc, char* argv[])
 
         q.submit([&](sycl::handler& h) {
 
-          auto grid = d_grid.get_access<sycl::access::mode::read_write>(h);
+          sycl::accessor grid(d_grid, h);
 
           unsigned begin = std::max(2,i-n+2);
           unsigned end   = std::min(i,n)+1;
@@ -172,7 +172,7 @@ int main(int argc, char* argv[])
       }
       q.submit([&](sycl::handler& h) {
 
-        auto grid = d_grid.get_access<sycl::access::mode::read_write>(h);
+        sycl::accessor grid(d_grid, h);
 
         h.single_task<class corner>([=] {
             grid[0*n+0] = -grid[(n-1)*n+(n-1)];

diff --git a/Cxx11/pic-sycl.cc b/Cxx11/pic-sycl.cc
@@ -523,14 +523,12 @@ int main(int argc, char ** argv) {
   std::string devname = (devchar==NULL ? "None" : devchar);
   sycl::device d;
   if (devname == "CPU") {
-      d = sycl::cpu_selector{}.select_device();
+      d = sycl::device{sycl::cpu_selector_v};
   } else if (devname == "GPU") {
-      d = sycl::gpu_selector{}.select_device();
-  } else if (devname == "HOST") {
-      d = sycl::host_selector{}.select_device();
+      d = sycl::device{sycl::gpu_selector_v};
   } else {
-      std::cout << "PRK_DEVICE should be CPU, GPU or HOST" << std::endl;
-      d = sycl::default_selector{}.select_device();
+      std::cout << "PRK_DEVICE should be CPU or GPU" << std::endl;
+      d = sycl::device{sycl::default_selector_v};
   }
   sycl::queue q(d);
   prk::SYCL::print_device_platform(q);
@@ -603,9 +601,8 @@ int main(int argc, char ** argv) {
 
           /* Calculate forces on particles and update positions */
           q.submit([&](sycl::handler& cgh) {
-
-              auto p = d_particles.get_access<sycl::access::mode::read_write>(cgh);
-              auto q = d_Qgrid.get_access<sycl::access::mode::read>(cgh);
+              sycl::accessor p(d_particles, cgh);
+              sycl::accessor q(d_Qgrid, cgh, sycl::read_only);
 
               cgh.parallel_for<class pic>(sycl::nd_range<1>(sycl::range<1>(global_work_size), sycl::range<1>(local_work_size)), [=] (sycl::nd_item<1> item) {
                   auto i = item.get_global_id(0);

diff --git a/Cxx11/prk_sycl.h b/Cxx11/prk_sycl.h
@@ -6,8 +6,6 @@
 
 #include "CL/sycl.hpp"
 
-namespace sycl = cl::sycl;
-
 #if defined(__LIBSYCL_MAJOR_VERSION) && defined(__LIBSYCL_MINOR_VERSION) && defined(__LIBSYCL_PATCH_VERSION)
 #    define __LIBSYCL_VERSION                                                                                          \
         (__LIBSYCL_MAJOR_VERSION * 10000 + __LIBSYCL_MINOR_VERSION * 100 + __LIBSYCL_PATCH_VERSION)

diff --git a/Cxx11/stencil-2d-sycl.cc b/Cxx11/stencil-2d-sycl.cc
@@ -123,7 +123,7 @@ void run(sycl::queue & q, int iterations, size_t n, size_t block_size, bool star
     q.submit([&](sycl::handler& h) {
 
       // accessor methods
-      auto in  = d_in.template get_access<sycl::access::mode::read_write>(h);
+      sycl::accessor in(d_in, h);
 
       h.parallel_for<class init<T>>(sycl::range<2> {n, n}, [=] (sycl::item<2> it) {
           sycl::id<2> xy = it.get_id();
@@ -142,7 +142,7 @@ void run(sycl::queue & q, int iterations, size_t n, size_t block_size, bool star
       q.wait();
 
       q.submit([&](sycl::handler& h) {
-        auto in  = d_in.template get_access<sycl::access::mode::read_write>(h);
+        sycl::accessor in(d_in, h);
         // Add constant to solution to force refresh of neighbor data, if any
         h.parallel_for<class add<T>>(sycl::range<2> {n, n}, [=] (sycl::item<2> it) {
             sycl::id<2> xy = it.get_id();
@@ -278,7 +278,7 @@ int main(int argc, char * argv[])
   //////////////////////////////////////////////////////////////////////
 
   try {
-    sycl::queue q{sycl::host_selector{}};
+    sycl::queue q{sycl::cpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, n, block_size, star, radius);
 #ifndef DPCPP_NO_DOUBLE
@@ -297,26 +297,7 @@ int main(int argc, char * argv[])
   }
 
   try {
-    sycl::queue q{sycl::cpu_selector{}};
-    prk::SYCL::print_device_platform(q);
-    run<float>(q, iterations, n, block_size, star, radius);
-#ifndef DPCPP_NO_DOUBLE
-    run<double>(q, iterations, n, block_size, star, radius);
-#endif
-  }
-  catch (sycl::exception & e) {
-    std::cout << e.what() << std::endl;
-    prk::SYCL::print_exception_details(e);
-  }
-  catch (std::exception & e) {
-    std::cout << e.what() << std::endl;
-  }
-  catch (const char * e) {
-    std::cout << e << std::endl;
-  }
-
-  try {
-    sycl::queue q{sycl::gpu_selector{}};
+    sycl::queue q{sycl::gpu_selector_v};
     prk::SYCL::print_device_platform(q);
     run<float>(q, iterations, n, block_size, star, radius);
 #ifndef DPCPP_NO_DOUBLE