From 46528f91dc23b3fa523c3650facbb1c6b279ef89 Mon Sep 17 00:00:00 2001 From: Wu Yingcong Date: Fri, 12 Jul 2024 18:09:02 +0800 Subject: [PATCH] [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 https://github.com/intel-restricted/applications.compilers.llvm-project/commit/f303cf35eaf173a59ee63e2d47d9201f8c45fed9#diff-bcdb4487c2827e573175267bdceb404b166460192e119ee45199d32d5fc13288 is missing some other necessary changes, causing the OOB access of `buffer1`. --- sycl/test-e2e/Basic/swizzle_op.cpp | 6 ++--- sycl/test-e2e/DeviceLib/cmath_test.cpp | 10 ++++----- .../exclusive_scan_sycl2020.cpp | 22 ++++++++++--------- .../inclusive_scan_sycl2020.cpp | 22 ++++++++++--------- 4 files changed, 32 insertions(+), 28 deletions(-) diff --git a/sycl/test-e2e/Basic/swizzle_op.cpp b/sycl/test-e2e/Basic/swizzle_op.cpp index 659a85749d5fe..9e96c51f9483f 100644 --- a/sycl/test-e2e/Basic/swizzle_op.cpp +++ b/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(); }); }); } diff --git a/sycl/test-e2e/DeviceLib/cmath_test.cpp b/sycl/test-e2e/DeviceLib/cmath_test.cpp index 6d3bc69e8690b..8206cba22bc8b 100644 --- a/sycl/test-e2e/DeviceLib/cmath_test.cpp +++ b/sycl/test-e2e/DeviceLib/cmath_test.cpp @@ -160,15 +160,15 @@ template void device_cmath_test_1(s::queue &deviceQueue) { // support from underlying device. #ifndef _WIN32 template 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 buffer1(result, numOfItems); + s::buffer buffer1(result, NumOfTestItems); s::buffer buffer2(&exponent, s::range<1>{1}); deviceQueue.submit([&](s::handler &cgh) { auto res_access = buffer1.template get_access(cgh); @@ -183,7 +183,7 @@ template 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])); } diff --git a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 2aef518350bc6..e90b70137473a 100644 --- a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/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(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( + 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(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( + 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, diff --git a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp index d83b1b50d32d4..e73215714a7c7 100644 --- a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/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(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( + 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(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( + 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,