[PATCH 42/46] testsuite, sycl: add test for 1-D and 2-D parallel_for kernels

[Tankut Baris Aktemur <tankut.baris.aktemur@intel.com>]

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