[PATCH v4 40/44] testsuite, sycl: add test for 1-D and 2-D parallel_for kernels
Markus Metzger <[email protected]>
| Newsgroups | gmane.comp.gdb.patches |
|---|---|
| Message-ID | <[email protected]> |
From: Tankut Baris Aktemur <[email protected]> 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 | 57 +++++++++++++++++ gdb/testsuite/gdb.sycl/parallel-for-2D.cpp | 73 ++++++++++++++++++++++ gdb/testsuite/gdb.sycl/parallel-for-2D.exp | 57 +++++++++++++++++ 4 files changed, 259 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 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..fa9f01c04aa --- /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-2026 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 "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> inBuffer {&in[0], dataRange}; + sycl::buffer<int, 1> outBuffer {&out[0], dataRange}; + + deviceQueue.submit ([&] (sycl::handler& cgh) /* line-before-kernel */ + { + auto accessorIn = inBuffer.get_access<sycl::access::mode::read> (cgh); + auto accessorOut + = outBuffer.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..2fd99df3ca0 --- /dev/null +++ b/gdb/testsuite/gdb.sycl/parallel-for-1D.exp @@ -0,0 +1,57 @@ +# Copyright 2019-2026 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 +} + +if {[build_executable "failed to compile $srcfile" "$binfile" "$srcfile" \ + {sycl debug}]} { + return +} + +foreach device $sycl_device_list { + sycl_with_intelgt_lock $device { + clean_restart $testfile + + 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..7811e57faf1 --- /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-2026 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 "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> inBuffer {&in[0][0], dataRangeIn}; + sycl::buffer<int, 2> outBuffer {&out[0][0], dataRangeOut}; + + deviceQueue.submit ([&] (sycl::handler& cgh) + { + auto accessorIn = inBuffer.get_access<sycl::access::mode::read> (cgh); + auto accessorOut + = outBuffer.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..0e1a68bda67 --- /dev/null +++ b/gdb/testsuite/gdb.sycl/parallel-for-2D.exp @@ -0,0 +1,57 @@ +# Copyright 2019-2026 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 +} + +if {[build_executable "failed to compile $srcfile" "$binfile" "$srcfile" \ + {sycl debug}]} { + return +} + +foreach device $sycl_device_list { + sycl_with_intelgt_lock $device { + clean_restart $testfile + + 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.*" + } + } + } +} -- 2.43.0 ________________________________________ Intel Deutschland GmbH Registered Address: Dornacher Strasse 1, 85622 Feldkirchen, Germany Tel: +49 (89) 99143-0 www.intel.de Managing Directors: Candice Moore, Jeffrey Schneiderman, Ramachandran Sitaraman Chairperson of the Supervisory Board: Sonja Pierer Registered Seat: Munich Commercial Register B: Amtsgericht Munich HRB 186928 This e-mail and any attachments may contain confidential material for the sole use of the intended recipient(s). Any review or distribution by others is strictly prohibited. If you are not the intended recipient, please contact the sender and delete all copies.