[PATCH v2 43/47] testsuite, sycl: add test for 1-D and 2-D parallel_for kernels

Tankut Baris Aktemur
2024-12-13T16:30:19Z

Add SYCL test cases for a parallel_for kernel on a 1-dimensional and
2-dimensional index space.
---
gdb/testsuite/gdb.sycl/parallel-for-1D.cpp | 72 +++++++++++++++++++++++++++
gdb/testsuite/gdb.sycl/parallel-for-1D.exp | 55 +++++++++++++++++++++
gdb/testsuite/gdb.sycl/parallel-for-2D.cpp | 73 ++++++++++++++++++++++++++++
gdb/testsuite/gdb.sycl/parallel-for-2D.exp | 55 +++++++++++++++++++++
gdb/testsuite/gdb.sycl/step-parallel-for.exp | 63 ++++++++++++++++++++++++
5 files changed, 318 insertions(+)

diff --git a/gdb/testsuite/gdb.sycl/parallel-for-1D.cpp b/gdb/testsuite/gdb.sycl/parallel-for-1D.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..bbcff7253b61c5ca53c3bb23a27500f46676a43f
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-1D.cpp
@@ -0,0 +1,72 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+ Copyright 2019-2024 Free Software Foundation, Inc.
+
+ This program is free software; you can redistribute it and/or modify
+ it under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3 of the License, or
+ (at your option) any later version.
+
+ This program is distributed in the hope that it will be useful,
+ but WITHOUT ANY WARRANTY; without even the implied warranty of
+ MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ GNU General Public License for more details.
+
+ You should have received a copy of the GNU General Public License
+ along with this program. If not, see <http://www.gnu.org/licenses/>. */
+
+#include <sycl/sycl.hpp>
+#include <iostream>
+#include "../lib/sycl-util.cpp"
+
+static int
+get_dim (sycl::id<1> wi, int index)
+{
+ return wi[index]; /* inside-kernel-callee */
+}
+
+int
+main (int argc, char *argv[])
+{
+ constexpr size_t DIM0 = 1024;
+
+ int in[DIM0];
+ int out[DIM0];
+
+ /* Initialize the input. */
+ for (unsigned int i = 0; i < DIM0; i++)
+ in[i] = i + 123;
+
+ { /* Extra scope enforces waiting on the kernel. */
+ sycl::queue deviceQueue {get_sycl_queue (argc, argv)};
+ sycl::range<1> dataRange {DIM0};
+ sycl::buffer<int, 1> bufferIn {&in[0], dataRange};
+ sycl::buffer<int, 1> bufferOut {&out[0], dataRange};
+
+ deviceQueue.submit ([&] (sycl::handler& cgh) /* line-before-kernel */
+ {
+ auto accessorIn = bufferIn.get_access<sycl::access::mode::read> (cgh);
+ auto accessorOut
+ = bufferOut.get_access<sycl::access::mode::write> (cgh);
+
+ cgh.parallel_for<class kernel> (dataRange, [=] (sycl::id<1> wiID)
+ {
+ int dim0 = get_dim (wiID, 0); /* kernel-first-line */
+ int in_elem = accessorIn[wiID];
+ int in_elem2 = accessorIn[dim0];
+ accessorOut[wiID] = in_elem + 100; /* kernel-last-line */
+ });
+ });
+ }
+
+ /* Verify the output. */
+ for (unsigned int i = 0; i < DIM0; i++)
+ if (out[i] != in[i] + 100)
+ {
+ std::cout << "Element " << i << " is " << out[i] << std::endl;
+ return 1;
+ }
+
+ std::cout << "Correct" << std::endl; /* end-marker */
+ return 0;
+}
diff --git a/gdb/testsuite/gdb.sycl/parallel-for-1D.exp b/gdb/testsuite/gdb.sycl/parallel-for-1D.exp
new file mode 100644
index 0000000000000000000000000000000000000000..53e2d902b0fc8e83a013c899622f217638926ab6
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-1D.exp
@@ -0,0 +1,55 @@
+# Copyright 2019-2024 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program. If not, see <http://www.gnu.org/licenses/>.
+#
+# Tests GDBs support for SYCL, for a parallel_for kernel on a 1D range.
+
+load_lib sycl.exp
+
+standard_testfile .cpp
+
+set sycl_device_list [init_sycl_devices_list]
+if {[llength $sycl_device_list] == 0} {
+ unsupported "target does not support SYCL"
+ return 0
+}
+
+if {[build_executable "failed to compile $srcfile" \
+ "${binfile}" $srcfile {sycl debug}]} {
+ return -1
+}
+
+foreach device $sycl_device_list {
+ sycl_with_intelgt_lock $device {
+ clean_restart "${binfile}"
+
+ if {![sycl_start $device]} {
+ continue
+ }
+
+ set last_line [gdb_get_line_number "kernel-last-line"]
+ gdb_breakpoint $last_line
+
+ # Check that we hit the BP for a number of elements. We do not check
+ # for each element because the number of hits received may depend on
+ # whether the kernel was vectorized, and if so, the width of vectors.
+ # Since the data array in the test program is large, having a small
+ # number of trips here should be safe.
+
+ for {set i 1} {$i <= 5} {incr i} { with_test_prefix "trip $i" {
+ gdb_continue_to_breakpoint "hit the last line" \
+ ".*$srcfile:$last_line.*"
+ }}
+ }
+}
diff --git a/gdb/testsuite/gdb.sycl/parallel-for-2D.cpp b/gdb/testsuite/gdb.sycl/parallel-for-2D.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..0607f8ffd2559c989236dbc0a9dc5220febfd46f
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-2D.cpp
@@ -0,0 +1,73 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+ Copyright 2019-2024 Free Software Foundation, Inc.
+
+ This program is free software; you can redistribute it and/or modify
+ it under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3 of the License, or
+ (at your option) any later version.
+
+ This program is distributed in the hope that it will be useful,
+ but WITHOUT ANY WARRANTY; without even the implied warranty of
+ MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ GNU General Public License for more details.
+
+ You should have received a copy of the GNU General Public License
+ along with this program. If not, see <http://www.gnu.org/licenses/>. */
+
+#include <sycl/sycl.hpp>
+#include <iostream>
+#include "../lib/sycl-util.cpp"
+
+int
+main (int argc, char *argv[])
+{
+ constexpr size_t DIM0 = 128;
+ constexpr size_t DIM1 = 64;
+
+ int in[DIM0][DIM1];
+ int out[DIM1][DIM0]; /* will transpose the input. */
+
+ /* Initialize the input. */
+ int val = 123;
+ for (unsigned int i = 0; i < DIM0; i++)
+ for (unsigned int j = 0; j < DIM1; j++)
+ in[i][j] = val++;
+
+ { /* Extra scope enforces waiting on the kernel. */
+ sycl::queue deviceQueue {get_sycl_queue (argc, argv)};
+ sycl::range<2> dataRangeIn {DIM0, DIM1};
+ sycl::range<2> dataRangeOut {DIM1, DIM0};
+ sycl::buffer<int, 2> bufferIn {&in[0][0], dataRangeIn};
+ sycl::buffer<int, 2> bufferOut {&out[0][0], dataRangeOut};
+
+ deviceQueue.submit ([&] (sycl::handler& cgh)
+ {
+ auto accessorIn = bufferIn.get_access<sycl::access::mode::read> (cgh);
+ auto accessorOut
+ = bufferOut.get_access<sycl::access::mode::write> (cgh);
+
+ cgh.parallel_for<class kernel> (dataRangeIn, [=] (sycl::id<2> wiID)
+ {
+ int dim0 = wiID[0]; /* kernel-first-line */
+ int dim1 = wiID[1];
+ int in_elem = accessorIn[wiID];
+ /* Negate the value, write into the transpositional location. */
+ accessorOut[dim1][dim0] = -1 * in_elem; /* kernel-last-line */
+ });
+ });
+ }
+
+ /* Verify the output. */
+ for (unsigned int i = 0; i < DIM0; i++)
+ for (unsigned int j = 0; j < DIM1; j++)
+ if (in[i][j] != -out[j][i])
+ {
+ std::cout << "Element " << j << "," << i
+ << " is " << out[j][i] << std::endl;
+ return 1;
+ }
+
+ std::cout << "Correct" << std::endl;
+ return 0;
+}
diff --git a/gdb/testsuite/gdb.sycl/parallel-for-2D.exp b/gdb/testsuite/gdb.sycl/parallel-for-2D.exp
new file mode 100644
index 0000000000000000000000000000000000000000..580e65cee15c1722d7760ffac14b969b7219999d
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-2D.exp
@@ -0,0 +1,55 @@
+# Copyright 2019-2024 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program. If not, see <http://www.gnu.org/licenses/>.
+#
+# Tests GDBs support for SYCL, for a parallel_for kernel on a 2D range.
+
+load_lib sycl.exp
+
+standard_testfile .cpp
+
+set sycl_device_list [init_sycl_devices_list]
+if {[llength $sycl_device_list] == 0} {
+ unsupported "target does not support SYCL"
+ return 0
+}
+
+if {[build_executable "failed to compile $srcfile" \
+ "${binfile}" $srcfile {sycl debug}]} {
+ return -1
+}
+
+foreach device $sycl_device_list {
+ sycl_with_intelgt_lock $device {
+ clean_restart "${binfile}"
+
+ if {![sycl_start $device]} {
+ continue
+ }
+
+ set last_line [gdb_get_line_number "kernel-last-line"]
+ gdb_breakpoint $last_line
+
+ # Check that we hit the BP for a number of elements. We do not check
+ # for each element because the number of hits received may depend on
+ # whether the kernel was vectorized, and if so, the width of vectors.
+ # Since the data array in the test program is large, having a small
+ # number of trips here should be safe.
+
+ for {set i 1} {$i <= 5} {incr i} { with_test_prefix "trip $i" {
+ gdb_continue_to_breakpoint "hit the last line" \
+ ".*$srcfile:$last_line.*"
+ }}
+ }
+}
diff --git a/gdb/testsuite/gdb.sycl/step-parallel-for.exp b/gdb/testsuite/gdb.sycl/step-parallel-for.exp
new file mode 100644
index 0000000000000000000000000000000000000000..0504843a36088dc70ec15f526fb1006717da28ea
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/step-parallel-for.exp
@@ -0,0 +1,63 @@
+# Copyright 2019-2024 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program. If not, see <http://www.gnu.org/licenses/>.
+#
+# Tests GDBs support for SYCL, for stepping inside a parallel_for kernel.
+
+load_lib sycl.exp
+
+standard_testfile parallel-for-1D.cpp
+
+set sycl_device_list [init_sycl_devices_list]
+if {[llength $sycl_device_list] == 0} {
+ unsupported "target does not support SYCL"
+ return 0
+}
+
+if {[build_executable "failed to compile $srcfile" \
+ "${binfile}" $srcfile {sycl debug}]} {
+ return -1
+}
+
+foreach device $sycl_device_list {
+ sycl_with_intelgt_lock $device {
+ clean_restart "${binfile}"
+
+ if {![sycl_start $device]} {
+ continue
+ }
+
+ set first_line [gdb_get_line_number "kernel-first-line"]
+ gdb_breakpoint $first_line
+
+ # Check that we can step inside the kernel. We do not check
+ # for each element because the number of hits received may depend on
+ # whether the kernel was vectorized, and if so, the width of vectors.
+ # Since the data array in the test program is large, having a small
+ # number of trips here should be safe.
+ #
+ # Lock the scheduler for stepping to avoid inference.
+
+ gdb_test_no_output "set scheduler-locking step"
+
+ for {set i 1} {$i <= 5} {incr i} { with_test_prefix "trip $i" {
+ gdb_continue_to_breakpoint "hit the first line" \
+ ".*$srcfile:$first_line.*"
+
+ gdb_test "next" "int in_elem = .*" "next 1"
+ gdb_test "next" "int in_elem2 = .*" "next 2"
+ gdb_test "next" "accessorOut.wiID. = in_elem.*" "next 3"
+ }}
+ }
+}