From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca by simark.ca with LMTP id IUjpA2l2fGoxkyAAWB0awg (envelope-from ) for ; Wed, 12 Aug 2026 09:34:33 -0400 Authentication-Results: simark.ca; dkim=pass (2048-bit key; unprotected) header.d=intel.com header.i=@intel.com header.a=rsa-sha256 header.s=Intel header.b=KRVxY9Ei; dkim-atps=neutral Received: by simark.ca (Postfix, from userid 112) id 0AC861E166; Wed, 12 Aug 2026 09:34:33 -0400 (EDT) X-Spam-Checker-Version: SpamAssassin 4.0.1 (2024-03-25) on simark.ca X-Spam-Level: X-Spam-Status: No, score=-6.4 required=5.0 tests=ARC_SIGNED,ARC_VALID,BAYES_00, DKIMWL_WL_HIGH,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,MAILING_LIST_MULTI, RCVD_IN_DNSWL_MED autolearn=ham autolearn_force=no version=4.0.1 Received: from vm01.sourceware.org (vm01.sourceware.org [38.145.34.32]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange x25519 server-signature ECDSA (prime256v1) server-digest SHA256) (No client certificate requested) by simark.ca (Postfix) with ESMTPS id 0C27A1E033 for ; Wed, 12 Aug 2026 09:34:32 -0400 (EDT) Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 69AE94BB3B87 for ; Wed, 12 Aug 2026 13:34:31 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 69AE94BB3B87 Authentication-Results: sourceware.org; dkim=pass (2048-bit key, unprotected) header.d=intel.com header.i=@intel.com header.a=rsa-sha256 header.s=Intel header.b=KRVxY9Ei Received: from mgamail.intel.com (mgamail.intel.com [192.198.163.12]) by sourceware.org (Postfix) with ESMTPS id D396A4BB3BA2 for ; Wed, 12 Aug 2026 13:30:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org D396A4BB3BA2 Authentication-Results: sourceware.org; dmarc=pass (p=none dis=none) header.from=intel.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=intel.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org D396A4BB3BA2 Authentication-Results: sourceware.org; arc=none smtp.remote-ip=192.198.163.12 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1786541414; cv=none; b=HAdCIFAFxkfbuJsWpFSZxx9Y33xMnun4cyxhtf6vseqytI8HlsNuMOsNXemPlFOthnZI+ez8MIZxWg7a/JhetS4imkDqfVnV4IooLVHlqkv/0yz+WoLd0FCXm2PWfrjl6EAXctJ1R/x0SPfHoLC3RM9CdfVTnm2j1sGbmqfvUTc= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1786541414; c=relaxed/simple; bh=zTbpK5m4XxNt2TFr5KGAovE4YRZUAwduTpC/t2yA5pw=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=O6lDYOwyfJchz+M8n7DGFLV2A60fpitzpTjRFKacwommzSAPnTlv3QQyrbrB1nTPmsGygRBMv6rpb0GriB42ODfz2Ubzj7qay9RsQbazAPFQ6N8KTFWTJI6dy2UIZtwGdQHleGdi5rHQsCjg5YX9AxYaSv4dHl7Gwk8Qc00SnkI= ARC-Authentication-Results: i=1; sourceware.org; dkim=pass (2048-bit key, unprotected) header.d=intel.com header.i=@intel.com header.a=rsa-sha256 header.s=Intel header.b=KRVxY9Ei DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org D396A4BB3BA2 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=intel.com; i=@intel.com; q=dns/txt; s=Intel; t=1786541414; x=1818077414; h=from:to:cc:subject:date:message-id:in-reply-to: references:mime-version:content-transfer-encoding; bh=zTbpK5m4XxNt2TFr5KGAovE4YRZUAwduTpC/t2yA5pw=; b=KRVxY9Eimyi7A5bGz6JW6W8OHvyJ90xIMN+zykn2PXOEY8GQ1QNjShOu UtV7mKtmvQLQ2gxXIrOsw2pYHnyZxJ2vyzKVMK4ngRJz1SSbVpuQz5PxS sDWyOZNaSVR/i/iYsdN3nt6hWnRiKwBwCn9NoeMfP91TyZMcVS5MeC1ru FoDFLWyWjTcTTme+kQODHk88MLrKTHFye0+1R64e2INx0JFXPQt8LziKY psd3sh/sc+y5ueLny8bVq1QI4x/5AQQ0ZYaeXsXYjXI8LvlNARwsXZgHk jlag6pvpp+Lf1VWtZOLzc1VEGwIcLrv567vM1eWMzdpNHO/k5AA7Yw/W/ Q==; X-CSE-ConnectionGUID: AbgPmTT8RJacxnpfXq/z9Q== X-CSE-MsgGUID: WU5K2MR/T9m/eQsdC0seLQ== X-IronPort-AV: E=McAfee;i="6800,10657,11872"; a="90910791" X-IronPort-AV: E=Sophos;i="6.25,219,1779174000"; d="scan'208";a="90910791" Received: from fmviesa005.fm.intel.com ([10.60.135.145]) by fmvoesa106.fm.intel.com with ESMTP/TLS/ECDHE-RSA-AES256-GCM-SHA384; 12 Aug 2026 06:30:13 -0700 X-CSE-ConnectionGUID: cBXL2vOPQsWIiak0i0DRlw== X-CSE-MsgGUID: VuDF7t5uQ0ikC5PxzmMOng== X-ExtLoop1: 1 X-IronPort-AV: E=Sophos;i="6.25,219,1779174000"; d="scan'208";a="268804819" Received: from gkldtt-dev-004.igk.intel.com (HELO localhost) ([10.123.221.202]) by fmviesa005-auth.fm.intel.com with ESMTP/TLS/ECDHE-RSA-AES256-GCM-SHA384; 12 Aug 2026 06:30:12 -0700 From: Markus Metzger To: gdb-patches@sourceware.org Cc: Tankut Baris Aktemur Subject: [PATCH v4 40/44] testsuite, sycl: add test for 1-D and 2-D parallel_for kernels Date: Wed, 12 Aug 2026 15:28:00 +0200 Message-ID: <20260812132805.380163-41-markus.t.metzger@intel.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260812132805.380163-1-markus.t.metzger@intel.com> References: <20260812132805.380163-1-markus.t.metzger@intel.com> MIME-Version: 1.0 Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: 7bit X-BeenThere: gdb-patches@sourceware.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gdb-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gdb-patches-bounces~public-inbox=simark.ca@sourceware.org From: Tankut Baris Aktemur 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 . */ + +#include +#include +#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 inBuffer {&in[0], dataRange}; + sycl::buffer outBuffer {&out[0], dataRange}; + + deviceQueue.submit ([&] (sycl::handler& cgh) /* line-before-kernel */ + { + auto accessorIn = inBuffer.get_access (cgh); + auto accessorOut + = outBuffer.get_access (cgh); + + cgh.parallel_for (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 . +# +# 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 . */ + +#include +#include +#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 inBuffer {&in[0][0], dataRangeIn}; + sycl::buffer outBuffer {&out[0][0], dataRangeOut}; + + deviceQueue.submit ([&] (sycl::handler& cgh) + { + auto accessorIn = inBuffer.get_access (cgh); + auto accessorOut + = outBuffer.get_access (cgh); + + cgh.parallel_for (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 . +# +# 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.