[PATCH 42/46] testsuite, sycl: add test for 1-D and 2-D parallel_for kernels
Tankut Baris Aktemur
tankut.baris.aktemur@intel.com
Tue Jul 2 11:43:13 GMT 2024
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 | 54 +++++++++++++++
gdb/testsuite/gdb.sycl/parallel-for-2D.cpp | 73 ++++++++++++++++++++
gdb/testsuite/gdb.sycl/parallel-for-2D.exp | 54 +++++++++++++++
gdb/testsuite/gdb.sycl/step-parallel-for.exp | 62 +++++++++++++++++
5 files changed, 315 insertions(+)
create mode 100644 gdb/testsuite/gdb.sycl/parallel-for-1D.cpp
create mode 100644 gdb/testsuite/gdb.sycl/parallel-for-1D.exp
create mode 100644 gdb/testsuite/gdb.sycl/parallel-for-2D.cpp
create mode 100644 gdb/testsuite/gdb.sycl/parallel-for-2D.exp
create mode 100644 gdb/testsuite/gdb.sycl/step-parallel-for.exp
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 00000000000..bbcff7253b6
--- /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 00000000000..56ea6b421f0
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-1D.exp
@@ -0,0 +1,54 @@
+# 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 {
+ clean_restart "${binfile}"
+ with_test_prefix [sycl_get_device_prefix $device] {
+ 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 00000000000..0607f8ffd25
--- /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 00000000000..3457fe7d657
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/parallel-for-2D.exp
@@ -0,0 +1,54 @@
+# 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 {
+ clean_restart "${binfile}"
+ with_test_prefix [sycl_get_device_prefix $device] {
+ 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 00000000000..4b62d869991
--- /dev/null
+++ b/gdb/testsuite/gdb.sycl/step-parallel-for.exp
@@ -0,0 +1,62 @@
+# 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 {
+ clean_restart "${binfile}"
+ with_test_prefix [sycl_get_device_prefix $device] {
+ 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"
+ }}
+ }
+}
--
2.34.1
Intel Deutschland GmbH
Registered Address: Am Campeon 10, 85579 Neubiberg, Germany
Tel: +49 89 99 8853-0, www.intel.de
Managing Directors: Sean Fennelly, Jeffrey Schneiderman, Tiffany Doon Silva
Chairperson of the Supervisory Board: Nicole Lau
Registered Office: Munich
Commercial Register: Amtsgericht Muenchen HRB 186928
More information about the Gdb-patches
mailing list