[SYCL][E2E] fix buffer out-of-bound access for tests (#14524)

- exclusive_scan_sycl2020.cpp
- inclusive_scan_sycl2020.cpp
The `InputContainer` has length of 12 or 128, but the `G` is constant
16, so it would have out-of-bound access for buffer `in_buf` and
`out_buf` when length is 12.

- swizzle_op.cpp
Open-and-shut OOB access.

- cmath_test.cpp
Change in
intel-restricted/applications.compilers.llvm-project@f303cf3#diff-bcdb4487c2827e573175267bdceb404b166460192e119ee45199d32d5fc13288
is missing some other necessary changes, causing the OOB access of
`buffer1`.

sycl/test-e2e/Basic/swizzle_op.cpp

@@ -93,7 +93,7 @@ int main() {
  B[0] = ab.x();
  B[1] = ab.y();
  B[2] = c.x();
- B[4] = c.y();
+ B[3] = c.y();
  });
  });
  }
@@ -117,7 +117,7 @@ int main() {
  B[0] = ab.x();
  B[1] = ab.y();
  B[2] = c.x();
- B[4] = c.y();
+ B[3] = c.y();
  });
  });
  }
@@ -141,7 +141,7 @@ int main() {
  B[0] = ab.x();
  B[1] = ab.y();
  B[2] = c.x();
- B[4] = c.y();
+ B[3] = c.y();
  });
  });
  }

sycl/test-e2e/DeviceLib/cmath_test.cpp

@@ -160,15 +160,15 @@ template <class T> void device_cmath_test_1(s::queue &deviceQueue) {
 // support from underlying device.
 #ifndef _WIN32
 template <class T> void device_cmath_test_2(s::queue &deviceQueue) {
- s::range<1> numOfItems{2};
- T result[2] = {-1};
- T ref[3] = {0, 2, 1};
+ constexpr int NumOfTestItems = 3;
+ T result[NumOfTestItems] = {-1};
+ T ref[NumOfTestItems] = {0, 2, 1};
  // Variable exponent is an integer value to store the exponent in frexp
  // function
  int exponent = -1;
 
  {
- s::buffer<T, 1> buffer1(result, numOfItems);
+ s::buffer<T, 1> buffer1(result, NumOfTestItems);
  s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
  deviceQueue.submit([&](s::handler &cgh) {
  auto res_access = buffer1.template get_access<sycl_write>(cgh);
@@ -183,7 +183,7 @@ template <class T> void device_cmath_test_2(s::queue &deviceQueue) {
  }
 
  // Compare result with reference
- for (int i = 0; i < 2; ++i) {
+ for (int i = 0; i < NumOfTestItems; ++i) {
  assert(approx_equal_fp(result[i], ref[i]));
  }
 

sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp

@@ -41,11 +41,12 @@ void test(const InputContainer &input, BinaryOperation binary_op,
  q.submit([&](handler &cgh) {
  accessor in{in_buf, cgh, sycl::read_only};
  accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
- cgh.parallel_for<kernel_name0>(nd_range<1>(G, G), [=](nd_item<1> it) {
- group<1> g = it.get_group();
- int lid = it.get_local_id(0);
- out[lid] = exclusive_scan_over_group(g, in[lid], binary_op);
- });
+ cgh.parallel_for<kernel_name0>(
+ nd_range<1>(confirmRange, confirmRange), [=](nd_item<1> it) {
+ group<1> g = it.get_group();
+ int lid = it.get_local_id(0);
+ out[lid] = exclusive_scan_over_group(g, in[lid], binary_op);
+ });
  });
  }
  emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
@@ -66,11 +67,12 @@ void test(const InputContainer &input, BinaryOperation binary_op,
  q.submit([&](handler &cgh) {
  accessor in{in_buf, cgh, sycl::read_only};
  accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
- cgh.parallel_for<kernel_name1>(nd_range<1>(G, G), [=](nd_item<1> it) {
- group<1> g = it.get_group();
- int lid = it.get_local_id(0);
- out[lid] = exclusive_scan_over_group(g, in[lid], init, binary_op);
- });
+ cgh.parallel_for<kernel_name1>(
+ nd_range<1>(confirmRange, confirmRange), [=](nd_item<1> it) {
+ group<1> g = it.get_group();
+ int lid = it.get_local_id(0);
+ out[lid] = exclusive_scan_over_group(g, in[lid], init, binary_op);
+ });
  });
  }
  emu::exclusive_scan(input.begin(), input.begin() + confirmRange,

sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp

@@ -41,11 +41,12 @@ void test(const InputContainer &input, BinaryOperation binary_op,
  q.submit([&](handler &cgh) {
  accessor in{in_buf, cgh, sycl::read_only};
  accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
- cgh.parallel_for<kernel_name0>(nd_range<1>(G, G), [=](nd_item<1> it) {
- group<1> g = it.get_group();
- int lid = it.get_local_id(0);
- out[lid] = inclusive_scan_over_group(g, in[lid], binary_op);
- });
+ cgh.parallel_for<kernel_name0>(
+ nd_range<1>(confirmRange, confirmRange), [=](nd_item<1> it) {
+ group<1> g = it.get_group();
+ int lid = it.get_local_id(0);
+ out[lid] = inclusive_scan_over_group(g, in[lid], binary_op);
+ });
  });
  }
  emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
@@ -66,11 +67,12 @@ void test(const InputContainer &input, BinaryOperation binary_op,
  q.submit([&](handler &cgh) {
  accessor in{in_buf, cgh, sycl::read_only};
  accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
- cgh.parallel_for<kernel_name1>(nd_range<1>(G, G), [=](nd_item<1> it) {
- group<1> g = it.get_group();
- int lid = it.get_local_id(0);
- out[lid] = inclusive_scan_over_group(g, in[lid], binary_op, init);
- });
+ cgh.parallel_for<kernel_name1>(
+ nd_range<1>(confirmRange, confirmRange), [=](nd_item<1> it) {
+ group<1> g = it.get_group();
+ int lid = it.get_local_id(0);
+ out[lid] = inclusive_scan_over_group(g, in[lid], binary_op, init);
+ });
  });
  }
  emu::inclusive_scan(input.begin(), input.begin() + confirmRange,