[SYCL][ESIMD] Enable verification for 32 simd vector length (#758)

Author: vasilytric

SYCL/ESIMD/api/functional/ctors/ctor_array_core.cpp

@@ -13,15 +13,6 @@
 // UNSUPPORTED: cuda, hip
 // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
-// TODO Unexpected runtime error "error: unsupported type for load/store" while
-// try to use simd::copy_from(), then simd::copy_to() with fixed-size array that
-// was defined on device side and the SIMD_RUN_TEST_WITH_VECTOR_LEN_1 macros
-// must be enabled when it is resolved.
-//
-// TODO This test disabled due to simd<short, 32> vector filled with unexpected
-// values from 16th element. The issue was created
-// https://github.com/intel/llvm/issues/5245 and and the
-// SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled when it is resolved.
 //
 // Test for simd constructor from an array.
 // This test uses different data types, dimensionality and different simd
@@ -180,16 +171,7 @@ int main(int, char **) {
   bool passed = true;
 
   const auto types = get_tested_types<tested_types::all>();
-#if defined(SIMD_RUN_TEST_WITH_VECTOR_LEN_1) &&                                \
-    defined(SIMD_RUN_TEST_WITH_VECTOR_LEN_32)
   const auto dims = get_all_dimensions();
-#elif SIMD_RUN_TEST_WITH_VECTOR_LEN_32
-  const auto dims = values_pack<8, 16, 32>();
-#elif SIMD_RUN_TEST_WITH_VECTOR_LEN_1
-  const auto dims = values_pack<1, 8, 16>();
-#else
-  const auto dims = values_pack<8, 16>();
-#endif
 
   // Run for specific combinations of types, vector length, and invocation
   // contexts.

SYCL/ESIMD/api/functional/ctors/ctor_fill_accuracy_core.cpp

@@ -15,8 +15,8 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
 // TODO simd<float, 32> fills with unexpected values while base value is denorm
-// and step is ulp. The SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled
-// when it is resolved.
+// and step is ulp. The SIMD_RUN_TEST_WITH_DENORM_INIT_VAL_AND_ULP_STEP macros
+// must be enabled when it is resolved.
 //
 // The test verifies that simd fill constructor has no precision differences.
 // The test do the following actions:

SYCL/ESIMD/api/functional/ctors/ctor_fill_core.cpp

@@ -14,11 +14,6 @@
 // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
-// TODO This test disabled due to simd<short, 32> vector filled with unexpected
-// values from 16th element. The issue was created
-// https://github.com/intel/llvm/issues/5245 and and the
-// SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled when it is resolved.
-//
 // Test for simd fill constructor for core types.
 // This test uses different data types, dimensionality, base and step values and
 // different simd constructor invocation contexts. The test do the following
@@ -89,11 +84,8 @@ int main(int, char **) {
   passed &= ctors::run_verification<ctors::const_ref, ctors::init_val::min_half,
                                     ctors::init_val::positive>(queue, two_dims,
                                                                char_int_types);
-#ifdef SIMD_RUN_TEST_WITH_VECTOR_LEN_32
-  const auto all_dims = values_pack<1, 8, 16, 32>();
-#else
-  const auto all_dims = values_pack<1, 8, 16>();
-#endif
+
+  const auto all_dims = get_all_dimensions();
   const auto all_types = get_tested_types<tested_types::all>();
   passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::min,
                                     ctors::init_val::zero>(queue, all_dims,

SYCL/ESIMD/api/functional/ctors/ctor_move.cpp

@@ -21,12 +21,6 @@
 // source simd instance. It is expected for a new simd instance to store
 // bitwise same data as the one passed as the source simd constructor.
 
-// The following issues for simd<T,32> observed:
-// - freeze with T in {char, unsigned char, singned char};
-// - runtime failure with T in {short, unsigned short}.
-// TODO Remove once the freeze is fixed
-#define SKIP_VECTOR_LEN_32
-
 // The test proxy is used to verify the move constructor was called actually.
 #define __ESIMD_ENABLE_TEST_PROXY
 
@@ -222,11 +216,7 @@ template <typename DataT, int NumElems, typename TestCaseT> class run_test {
 int main(int, char **) {
   bool passed = true;
   const auto types = get_tested_types<tested_types::all>();
-#ifdef SKIP_VECTOR_LEN_32
-  const auto dims = values_pack<1, 8, 16>();
-#else
   const auto dims = get_all_dimensions();
-#endif
 
   sycl::queue queue(esimd_test::ESIMDSelector{},
                     esimd_test::createExceptionHandler());

SYCL/ESIMD/api/functional/ctors/ctor_vector_core.cpp

@@ -14,14 +14,6 @@
 // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
-// TODO The tests freezed during runtime when using simd<char, 32>. The
-// SIMD_RUN_TEST_WITH_CHAR_TYPES macros must be enabled when it is resolved.
-//
-// TODO This test disabled due to simd<short, 32> vector filled with unexpected
-// values from 16th element. The issue was created
-// https://github.com/intel/llvm/issues/5245 and and the
-// SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled when it is resolved.
-//
 // Test for simd constructor from vector.
 // This test uses different data types, dimensionality and different simd
 // constructor invocation contexts.
@@ -172,20 +164,8 @@ int main(int, char **) {
 
   bool passed = true;
 
-#ifdef SIMD_RUN_TEST_WITH_CHAR_TYPES
   const auto types = get_tested_types<tested_types::all>();
-#else
-  const auto types = named_type_pack<short, unsigned short, int, unsigned int,
-                                     long, unsigned long, float, double,
-                                     long long, unsigned long long>(
-      {"short", "unsigned short", "int", "unsigned int", "long",
-       "unsigned long", "float", "double", "long long", "unsigned long long"});
-#endif
-#ifdef SIMD_RUN_TEST_WITH_VECTOR_LEN_32
   const auto dims = get_all_dimensions();
-#else
-  const auto dims = values_pack<1, 8, 16>();
-#endif
 
   // Run for specific combinations of types, vector length and invocation
   // contexts.