[ESIMD][NFC][E2E] Fix 570 compilation warnings in ESIMD E2E tests (#12748)

Warnings fixed:
- deprecated scatter_rgba
- deprecated get_cl_code
- deprecated lsc_fence
- deprecated uchar type usage
- deprecated get_access on HOST
- deprecated get_pointer
- usage of isfinite with -ffast-math
- deprecated dpas_argument_type::s1
- deprecated gpu_selector()

Also, the memory alloc/free in historgram*.cpp tests were updated to
simplify the potential memory leak avoidance.

Signed-off-by: Klochkov, Vyacheslav N <[email protected]>

Author: v-klochkov

sycl/test-e2e/ESIMD/accessor_global.cpp

@@ -2,7 +2,7 @@
 // RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
 // RUN: %{run} %t.out
 
-// This test verifies usage of accessor methods operator[] and get_pointer().
+// This test verifies usage of accessor methods operator[] and get_multi_ptr().
 
 #include "esimd_test_utils.hpp"
 
@@ -39,7 +39,8 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
            for (int I = 0; I < VL; I++)
              TmpAcc[GID * VL + I] = GID * 100 + I;
          } else {
-           T *Ptr = TmpAcc.get_pointer();
+           T *Ptr =
+               TmpAcc.template get_multi_ptr<access::decorated::yes>().get();
            simd<int, VL> IntValues(GID * 100, 1);
            simd<T, VL> Values = IntValues;
            block_store(Ptr + GID * VL, Values);
@@ -53,12 +54,13 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
                for (int I = 0; I < VL; I++)
                  Out[(GID + LID) * VL + I] = TmpAcc[(GID + LID) * VL + I];
              } else {
-               T *Ptr = TmpAcc.get_pointer();
+               T *Ptr = TmpAcc.template get_multi_ptr<access::decorated::yes>()
+                            .get();
                simd<T, VL> Values = block_load<T, VL>(Ptr + (GID + LID) * VL);
                Values.template copy_to(Out + (GID + LID) * VL);
              }
            } // end for (int LID = 0; LID < LocalRange; LID++)
-         }   // end if (LID == 0)
+         } // end if (LID == 0)
        });
      }).wait();
   } catch (sycl::exception const &e) {

sycl/test-e2e/ESIMD/accessor_local.cpp

@@ -43,8 +43,10 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
        CGH.parallel_for(NDRange, [=](nd_item<1> Item) SYCL_ESIMD_KERNEL {
          uint32_t GID = Item.get_global_id(0);
          uint32_t LID = Item.get_local_id(0);
-         uint32_t LocalAccOffset = static_cast<uint32_t>(
-             reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get()));
+         uint32_t LocalAccOffset =
+             static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
+                 LocalAcc.template get_multi_ptr<access::decorated::yes>()
+                     .get()));
          if constexpr (TestSubscript) {
            for (int I = 0; I < VL; I++)
              LocalAcc[LID * VL + I] = GID * 100 + I;
@@ -67,7 +69,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
                ValuesFromSLM.copy_to(Out + (GID + LID) * VL);
              }
            } // end for (int LID = 0; LID < LocalRange; LID++)
-         }   // end if (LID == 0)
+         } // end if (LID == 0)
        });
      }).wait();
   } catch (sycl::exception const &e) {

sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp

@@ -24,7 +24,7 @@ using namespace sycl::ext::intel::esimd;
 template <unsigned VL, class T, class F>
 bool test(queue q, std::string str, F funcUnderTest) {
   std::cout << "Testing " << str << ", VL = " << VL << " ...\n";
-  size_t Size = 4 * VL;
+  constexpr size_t Size = 4 * VL;
   T A[Size];
   T B[Size];
   constexpr unsigned HalfVL = VL > 1 ? (VL / 2) : 1;

sycl/test-e2e/ESIMD/dpas/dpas_common.hpp

@@ -51,9 +51,7 @@ std::string toString(dpas_argument_type T) {
     return "bf16";
   case dpas_argument_type::tf32:
     return "tf32";
-  case dpas_argument_type::s1:
-  case dpas_argument_type::u1:
-  case dpas_argument_type::Invalid:
+  default:
     return "UNSUPPORTED";
   }
   return "UNRECOGNIZED";
@@ -127,9 +125,7 @@ template <dpas_argument_type T> constexpr int getBitSize() {
   case dpas_argument_type::tf32:
     return 32;
 
-  case dpas_argument_type::Invalid:
-  case dpas_argument_type::s1:
-  case dpas_argument_type::u1:
+  default:
     break;
   }
   return 0;
@@ -405,7 +401,7 @@ bool test(queue &Q, bool Print) {
                   << ") != expected (" << GoldRes << ")" << std::endl;
       }
     } // end for JJ
-  }   // end for II
+  } // end for II
 
   free(Res, Q);
   free(APacked, Q);

sycl/test-e2e/ESIMD/esimd_test_utils.hpp

@@ -27,6 +27,10 @@ using namespace sycl;
 
 namespace esimd_test {
 
+template <typename T>
+using shared_allocator = sycl::usm_allocator<T, sycl::usm::alloc::shared>;
+template <typename T> using shared_vector = std::vector<T, shared_allocator<T>>;
+
 // This is the function provided to SYCL runtime by the application to decide
 // on which device to run, or whether to run at all.
 // When selecting a device, SYCL runtime first takes (1) a selector provided by

sycl/test-e2e/ESIMD/ext_math.cpp

@@ -388,7 +388,10 @@ bool test(queue &Q, const std::string &Name, InitF Init = InitNarrow<T>{},
     if constexpr (sizeof(T) <= 2)
       delta = delta + delta;
 
-    bool BothFinite = std::isfinite(Test) && std::isfinite(Gold);
+    bool BothFinite = true;
+#ifndef TEST_FAST_MATH
+    BothFinite = std::isfinite(Test) && std::isfinite(Gold);
+#endif
     if (BothFinite && std::abs(Test - Gold) > delta) {
       if (++ErrCnt < 10) {
         std::cout << "    failed at index " << I << ", " << Test

sycl/test-e2e/ESIMD/grf.cpp

@@ -71,13 +71,10 @@ int main(void) {
     A[i] = i;
   }
 
+  queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+  esimd_test::printTestLabel(q);
   try {
     buffer<float, 1> bufa(A.data(), range<1>(Size));
-    queue q(gpu_selector{}, esimd_test::createExceptionHandler());
-
-    auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
-
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<class SyclKernel>(Size,
@@ -98,11 +95,6 @@ int main(void) {
 
   try {
     buffer<float, 1> bufa(A.data(), range<1>(Size));
-    queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
-
-    auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
-
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<class EsimdKernel>(Size, [=](id<1> i) SYCL_ESIMD_KERNEL {
@@ -128,7 +120,6 @@ int main(void) {
 
   try {
     buffer<float, 1> bufa(A.data(), range<1>(Size));
-    queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 #ifdef USE_AUTO
     sycl::ext::oneapi::experimental::properties prop{grf_size_automatic};
 #elif defined(USE_NEW_API)
@@ -137,9 +128,6 @@ int main(void) {
     sycl::ext::oneapi::experimental::properties prop{
         register_alloc_mode<register_alloc_mode_enum::large>};
 #endif
-    auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
-
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<class EsimdKernelSpecifiedGRF>(

sycl/test-e2e/ESIMD/histogram.cpp

@@ -85,15 +85,15 @@ int main(int argc, char *argv[]) {
 
   // Allocate Input Buffer
   queue q = esimd_test::createQueue();
+  esimd_test::printTestLabel(q);
 
-  auto dev = q.get_device();
-  unsigned char *srcY = malloc_shared<unsigned char>(width * height, q);
-  if (srcY == NULL) {
-    std::cerr << "Out of memory\n";
-    exit(1);
-  }
-  unsigned int *bins = malloc_shared<unsigned int>(NUM_BINS, q);
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  esimd_test::shared_vector<uint8_t> srcY_vec(
+      width * height, esimd_test::shared_allocator<uint8_t>{q});
+  esimd_test::shared_vector<unsigned int> bins_vec(
+      NUM_BINS, esimd_test::shared_allocator<unsigned int>{q});
+  uint8_t *srcY = srcY_vec.data();
+  ;
+  unsigned int *bins = bins_vec.data();
 
   uint range_width = width / BLOCK_WIDTH;
   uint range_height = height / BLOCK_HEIGHT;
@@ -106,16 +106,12 @@ int main(int argc, char *argv[]) {
     FILE *f = fopen(input_file, "rb");
     if (f == NULL) {
       std::cerr << "Error opening file " << input_file;
-      free(srcY, q);
-      free(bins, q);
       std::exit(1);
     }
 
     unsigned int cnt = fread(srcY, sizeof(unsigned char), input_size, f);
     if (cnt != input_size) {
       std::cerr << "Error reading input from " << input_file;
-      free(srcY, q);
-      free(bins, q);
       std::exit(1);
     }
   } else {
@@ -171,18 +167,17 @@ int main(int argc, char *argv[]) {
               uint h_pos = (tid % range_width) * BLOCK_WIDTH;
               uint v_pos = (tid / range_width) * BLOCK_HEIGHT;
 
-              // Declare a 8x32 uchar matrix to store the input block pixel
+              // Declare a 8x32 uint8_t matrix to store the input block pixel
               // value
-              simd<unsigned char, 8 * 32> in;
+              simd<uint8_t, 8 * 32> in;
 
               // Declare a vector to store the local histogram
               simd<unsigned int, NUM_BINS> histogram(0);
 
               // Each thread handles BLOCK_HEIGHTxBLOCK_WIDTH pixel block
               for (int y = 0; y < BLOCK_HEIGHT / 8; y++) {
                 // Perform 2D media block read to load 8x32 pixel block
-                in = media_block_load<unsigned char, 8, 32>(readAcc, h_pos,
-                                                            v_pos);
+                in = media_block_load<uint8_t, 8, 32>(readAcc, h_pos, v_pos);
 
             // Accumulate local histogram for each pixel value
 #pragma unroll
@@ -236,8 +231,6 @@ int main(int argc, char *argv[]) {
     // make sure data is given back to the host at the end of this scope
   } catch (sycl::exception const &e) {
     std::cerr << "SYCL exception caught: " << e.what() << '\n';
-    free(srcY, q);
-    free(bins, q);
     return 1;
   }
 
@@ -251,8 +244,6 @@ int main(int argc, char *argv[]) {
   writeHist(cpuHistogram);
   // Checking Histogram
   bool Success = checkHistogram(cpuHistogram, bins);
-  free(srcY, q);
-  free(bins, q);
 
   if (!Success) {
     std::cerr << "FAILED\n";

sycl/test-e2e/ESIMD/histogram_256_slm.cpp

@@ -45,7 +45,7 @@ ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
     auto start_addr = ((unsigned int *)input_ptr) + start_off;
     simd<uint, 32> data;
     data.copy_from(start_addr);
-    auto in = data.bit_cast_view<uchar>();
+    auto in = data.bit_cast_view<uint8_t>();
 
 #pragma unroll
     for (int j = 0; j < BLOCK_WIDTH * sizeof(int); j += 16) {
@@ -75,7 +75,7 @@ void HistogramCPU(unsigned int size, unsigned int *src,
                   unsigned int *cpu_histogram) {
   for (int i = 0; i < size; i++) {
     unsigned int x = src[i];
-    cpu_histogram[(x)&0xFFU] += 1;
+    cpu_histogram[(x) & 0xFFU] += 1;
     cpu_histogram[(x >> 8) & 0xFFU] += 1;
     cpu_histogram[(x >> 16) & 0xFFU] += 1;
     cpu_histogram[(x >> 24) & 0xFFU] += 1;
@@ -104,14 +104,18 @@ int CheckHistogram(unsigned int *cpu_histogram, unsigned int *gpu_histogram) {
 
 int main() {
   queue q = esimd_test::createQueue();
+  esimd_test::printTestLabel(q);
 
   const char *input_file = nullptr;
   unsigned int width = 1024;
   unsigned int height = 1024;
 
   // Initializes input.
   unsigned int input_size = width * height;
-  unsigned int *input_ptr = malloc_shared<unsigned int>(input_size, q);
+
+  esimd_test::shared_vector<unsigned int> input_vec(
+      input_size, esimd_test::shared_allocator<unsigned int>{q});
+  unsigned int *input_ptr = input_vec.data();
   printf("Processing %dx%d inputs\n", width, height);
 
   srand(2009);
@@ -124,13 +128,8 @@ int main() {
 
   // Allocates system memory for output buffer.
   int buffer_size = sizeof(unsigned int) * NUM_BINS;
-  unsigned int *hist = new unsigned int[buffer_size];
-  if (hist == nullptr) {
-    free(input_ptr, q);
-    std::cerr << "Out of memory\n";
-    exit(1);
-  }
-  memset(hist, 0, buffer_size);
+  std::vector<unsigned int> hist_vec(buffer_size, 0);
+  unsigned int *hist = hist_vec.data();
 
   // Uses the CPU to calculate the histogram output data.
   unsigned int cpu_histogram[NUM_BINS];
@@ -141,7 +140,9 @@ int main() {
   std::cout << "finish cpu_histogram\n";
 
   // Uses the GPU to calculate the histogram output data.
-  unsigned int *output_surface = malloc_shared<unsigned int>(NUM_BINS, q);
+  esimd_test::shared_vector<unsigned int> output_vec(
+      NUM_BINS, esimd_test::shared_allocator<unsigned int>{q});
+  unsigned int *output_surface = output_vec.data();
 
   unsigned int num_threads;
   num_threads = width * height / (NUM_BLOCKS * BLOCK_WIDTH);
@@ -194,9 +195,6 @@ int main() {
 
   memcpy(hist, output_surface, 4 * NUM_BINS);
 
-  free(output_surface, q);
-  free(input_ptr, q);
-
   // Compares the CPU histogram output data with the
   // GPU histogram output data.
   // If there is no difference, the result is correct.
@@ -207,7 +205,5 @@ int main() {
   else
     std::cout << "FAILED\n";
 
-  delete[] hist;
-
   return res ? 0 : -1;
 }