[SYCL][ESIMD]Add tests for support for different types for lsc functions

SYCL/ESIMD/lsc/Inputs/common.hpp

@@ -7,12 +7,19 @@
 //===----------------------------------------------------------------------===//
 
 #include <stdlib.h>
+#include <sycl/bit_cast.hpp>
 
 template <int case_num> class KernelID;
 
 template <typename T> T get_rand() {
-  T v = rand();
-  if constexpr (sizeof(T) > 4)
+  using Tuint = std::conditional_t<
+      sizeof(T) == 1, uint8_t,
+      std::conditional_t<
+          sizeof(T) == 2, uint16_t,
+          std::conditional_t<sizeof(T) == 4, uint32_t,
+                             std::conditional_t<sizeof(T) == 8, uint64_t, T>>>>;
+  Tuint v = rand();
+  if constexpr (sizeof(Tuint) > 4)
     v = (v << 32) | rand();
-  return v;
+  return sycl::bit_cast<T>(v);
 }

SYCL/ESIMD/lsc/Inputs/lsc_surf_load.hpp

@@ -38,8 +38,8 @@ bool test(uint32_t pmask = 0xffffffff) {
   }
 
   uint16_t Size = Groups * Threads * VL * VS;
-
-  T vmask = (T)-1;
+  using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
+  Tuint vmask = (Tuint)-1;
   if constexpr (DS == lsc_data_size::u8u32)
     vmask = (T)0xff;
   if constexpr (DS == lsc_data_size::u16u32)
@@ -124,20 +124,24 @@ bool test(uint32_t pmask = 0xffffffff) {
 
   if constexpr (transpose) {
     for (int i = 0; i < Size; i++) {
-      T e = in[i];
-      if (out[i] != e) {
+      Tuint e = sycl::bit_cast<Tuint>(in[i]);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   } else {
     for (int i = 0; i < Size; i++) {
-      T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
-      if (out[i] != e) {
+      Tuint in_val = sycl::bit_cast<Tuint>(in[i]);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      Tuint e = (pmask >> ((i / VS) % VL)) & 1 ? in_val & vmask
+                                               : sycl::bit_cast<Tuint>(old_val);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   }

SYCL/ESIMD/lsc/Inputs/lsc_surf_store.hpp

@@ -37,8 +37,8 @@ bool test(uint32_t pmask = 0xffffffff) {
   }
 
   uint16_t Size = Groups * Threads * VL * VS;
-
-  T vmask = (T)-1;
+  using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
+  Tuint vmask = (Tuint)-1;
   if constexpr (DS == lsc_data_size::u8u32)
     vmask = (T)0xff;
   if constexpr (DS == lsc_data_size::u16u32)
@@ -104,22 +104,28 @@ bool test(uint32_t pmask = 0xffffffff) {
 
   if constexpr (transpose) {
     for (int i = 0; i < Size; i++) {
-      T e = new_val + i;
-      if (out[i] != e) {
+      T expected_value = new_val + i;
+      Tuint e = sycl::bit_cast<Tuint>(expected_value);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   } else {
     for (int i = 0; i < Size; i++) {
-      T e = (pmask >> ((i / VS) % VL)) & 1
-                ? ((new_val + i) & vmask) | (old_val & ~vmask)
-                : old_val;
-      if (out[i] != e) {
+      T expected_value = new_val + i;
+      Tuint in_val = sycl::bit_cast<Tuint>(expected_value);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      Tuint e =
+          (pmask >> ((i / VS) % VL)) & 1
+              ? (in_val & vmask) | (sycl::bit_cast<Tuint>(old_val) & ~vmask)
+              : sycl::bit_cast<Tuint>(old_val);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   }

SYCL/ESIMD/lsc/Inputs/lsc_usm_load.hpp

@@ -38,8 +38,8 @@ bool test(uint32_t pmask = 0xffffffff) {
   }
 
   uint16_t Size = Groups * Threads * VL * VS;
-
-  T vmask = (T)-1;
+  using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
+  Tuint vmask = (Tuint)-1;
   if constexpr (DS == lsc_data_size::u8u32)
     vmask = (T)0xff;
   if constexpr (DS == lsc_data_size::u16u32)
@@ -125,20 +125,24 @@ bool test(uint32_t pmask = 0xffffffff) {
 
   if constexpr (transpose) {
     for (int i = 0; i < Size; i++) {
-      T e = in[i];
-      if (out[i] != e) {
+      Tuint e = sycl::bit_cast<Tuint>(in[i]);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   } else {
     for (int i = 0; i < Size; i++) {
-      T e = (pmask >> ((i / VS) % VL)) & 1 ? in[i] & vmask : old_val;
-      if (out[i] != e) {
+      Tuint in_val = sycl::bit_cast<Tuint>(in[i]);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      Tuint e = (pmask >> ((i / VS) % VL)) & 1 ? in_val & vmask
+                                               : sycl::bit_cast<Tuint>(old_val);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   }

SYCL/ESIMD/lsc/Inputs/lsc_usm_store.hpp

@@ -37,8 +37,9 @@ bool test(uint32_t pmask = 0xffffffff) {
   }
 
   uint16_t Size = Groups * Threads * VL * VS;
+  using Tuint = sycl::_V1::ext::intel::esimd::detail::uint_type_t<sizeof(T)>;
 
-  T vmask = (T)-1;
+  Tuint vmask = (Tuint)-1;
   if constexpr (DS == lsc_data_size::u8u32)
     vmask = (T)0xff;
   if constexpr (DS == lsc_data_size::u16u32)
@@ -104,22 +105,29 @@ bool test(uint32_t pmask = 0xffffffff) {
 
   if constexpr (transpose) {
     for (int i = 0; i < Size; i++) {
-      T e = new_val + i;
-      if (out[i] != e) {
+      T expected_value = new_val + i;
+      Tuint e = sycl::bit_cast<Tuint>(expected_value);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   } else {
     for (int i = 0; i < Size; i++) {
-      T e = (pmask >> ((i / VS) % VL)) & 1
-                ? ((new_val + i) & vmask) | (old_val & ~vmask)
-                : old_val;
-      if (out[i] != e) {
+      T expected_value = new_val + i;
+      Tuint in_val = sycl::bit_cast<Tuint>(expected_value);
+      Tuint out_val = sycl::bit_cast<Tuint>(out[i]);
+
+      Tuint e =
+          (pmask >> ((i / VS) % VL)) & 1
+              ? (in_val & vmask) | (sycl::bit_cast<Tuint>(old_val) & ~vmask)
+              : sycl::bit_cast<Tuint>(old_val);
+      if (out_val != e) {
         passed = false;
-        std::cout << "out[" << i << "] = 0x" << std::hex << (uint64_t)out[i]
-                  << " vs etalon = 0x" << (uint64_t)e << std::dec << std::endl;
+        std::cout << "out[" << i << "] = 0x" << std::hex << out_val
+                  << " vs etalon = 0x" << e << std::dec << std::endl;
       }
     }
   }

SYCL/ESIMD/lsc/lsc_surf_load_u32.cpp

@@ -14,25 +14,34 @@
 
 constexpr uint32_t seed = 199;
 
-int main(void) {
-  srand(seed);
+template <int TestCastNum, typename T> bool tests() {
   bool passed = true;
 
   // non transpose
-  passed &= test<0, uint32_t, 1, 4, 32, 1, false>(rand());
-  passed &= test<1, uint32_t, 1, 4, 32, 2, false>(rand());
-  passed &= test<2, uint32_t, 1, 4, 16, 2, false>(rand());
-  passed &= test<3, uint32_t, 1, 4, 4, 1, false>(rand());
-  passed &= test<4, uint32_t, 1, 1, 1, 1, false>(1);
-  passed &= test<5, uint32_t, 2, 1, 1, 1, false>(1);
+  passed &= test<TestCastNum, T, 1, 4, 32, 1, false>(rand());
+  passed &= test<TestCastNum + 1, T, 1, 4, 32, 2, false>(rand());
+  passed &= test<TestCastNum + 2, T, 1, 4, 16, 2, false>(rand());
+  passed &= test<TestCastNum + 3, T, 1, 4, 4, 1, false>(rand());
+  passed &= test<TestCastNum + 4, T, 1, 1, 1, 1, false>(1);
+  passed &= test<TestCastNum + 5, T, 2, 1, 1, 1, false>(1);
 
-  // passed &= test<6, uint32_t, 1, 4, 8, 2, false>(rand());
-  // passed &= test<7, uint32_t, 1, 4, 8, 3, false>(rand());
+  // passed &= test<TestCastNum+6, T, 1, 4, 8, 2, false>(rand());
+  // passed &= test<TestCastNum+7, T, 1, 4, 8, 3, false>(rand());
 
   // transpose
-  passed &= test<8, uint32_t, 1, 4, 1, 32, true>();
-  passed &= test<9, uint32_t, 2, 2, 1, 16, true>();
-  passed &= test<10, uint32_t, 4, 4, 1, 4, true>();
+  passed &= test<TestCastNum + 8, T, 1, 4, 1, 32, true>();
+  passed &= test<TestCastNum + 9, T, 2, 2, 1, 16, true>();
+  passed &= test<TestCastNum + 10, T, 4, 4, 1, 4, true>();
+  return passed;
+}
+
+int main(void) {
+  srand(seed);
+  bool passed = true;
+
+  passed &= tests<0, uint32_t>();
+  passed &= tests<11, float>();
+  passed &= tests<22, sycl::ext::intel::experimental::esimd::tfloat32>();
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/lsc/lsc_surf_load_u32_stateless.cpp

@@ -0,0 +1,13 @@
+//==----- lsc_surf_load_u32_stateless.cpp - DPC++ ESIMD on-device test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu-intel-pvc || esimd_emulator
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include "lsc_surf_load_u32.cpp"

SYCL/ESIMD/lsc/lsc_surf_load_u64.cpp

@@ -13,26 +13,32 @@
 #include "Inputs/lsc_surf_load.hpp"
 
 constexpr uint32_t seed = 198;
-
-int main(void) {
-  srand(seed);
+template <int TestCastNum, typename T> bool tests() {
   bool passed = true;
 
   // non transpose
-  passed &= test<0, uint64_t, 1, 4, 32, 1, false>(rand());
-  passed &= test<1, uint64_t, 1, 4, 32, 2, false>(rand());
-  passed &= test<2, uint64_t, 1, 4, 16, 2, false>(rand());
-  passed &= test<3, uint64_t, 1, 4, 4, 1, false>(rand());
-  passed &= test<4, uint64_t, 1, 1, 1, 1, false>(1);
-  passed &= test<5, uint64_t, 2, 1, 1, 1, false>(1);
+  passed &= test<TestCastNum, T, 1, 4, 32, 1, false>(rand());
+  passed &= test<TestCastNum + 1, T, 1, 4, 32, 2, false>(rand());
+  passed &= test<TestCastNum + 2, T, 1, 4, 16, 2, false>(rand());
+  passed &= test<TestCastNum + 3, T, 1, 4, 4, 1, false>(rand());
+  passed &= test<TestCastNum + 4, T, 1, 1, 1, 1, false>(1);
+  passed &= test<TestCastNum + 5, T, 2, 1, 1, 1, false>(1);
 
-  // passed &= test<6, uint64_t, 1, 4, 8, 2, false>(rand());
-  // passed &= test<7, uint64_t, 1, 4, 8, 3, false>(rand());
+  // passed &= test<TestCastNum+6, T, 1, 4, 8, 2, false>(rand());
+  // passed &= test<TestCastNum+7, T, 1, 4, 8, 3, false>(rand());
 
   // transpose
-  passed &= test<8, uint64_t, 1, 4, 1, 32, true>();
-  passed &= test<9, uint64_t, 2, 2, 1, 16, true>();
-  passed &= test<10, uint64_t, 4, 4, 1, 4, true>();
+  passed &= test<TestCastNum + 8, T, 1, 4, 1, 32, true>();
+  passed &= test<TestCastNum + 9, T, 2, 2, 1, 16, true>();
+  passed &= test<TestCastNum + 10, T, 4, 4, 1, 4, true>();
+  return passed;
+}
+int main(void) {
+  srand(seed);
+  bool passed = true;
+
+  passed &= tests<0, uint64_t>();
+  passed &= tests<11, double>();
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/lsc/lsc_surf_load_u64_stateless.cpp

@@ -0,0 +1,13 @@
+//==----- lsc_surf_load_u64_stateless.cpp - DPC++ ESIMD on-device test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu-intel-pvc || esimd_emulator
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include "lsc_surf_load_u64.cpp"

SYCL/ESIMD/lsc/lsc_surf_load_u8_u16.cpp

@@ -29,6 +29,8 @@ int main(void) {
 
   passed &= tests<0, uint8_t>();
   passed &= tests<3, uint16_t>();
+  passed &= tests<6, sycl::ext::oneapi::experimental::bfloat16>();
+  passed &= tests<9, half>();
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/lsc/lsc_surf_load_u8_u16_stateless.cpp

@@ -0,0 +1,13 @@
+//==----- lsc_surf_load_u8_u16_stateless.cpp - DPC++ ESIMD on-device test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu-intel-pvc || esimd_emulator
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -fsycl-esimd-force-stateless-mem %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include "lsc_surf_load_u8_u16.cpp"