From patchwork Thu Dec 14 20:22:38 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Pedro Alves X-Patchwork-Id: 82174 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 8CD0A3861880 for ; Thu, 14 Dec 2023 20:23:57 +0000 (GMT) X-Original-To: gdb-patches@sourceware.org Delivered-To: gdb-patches@sourceware.org Received: from mail-wm1-f51.google.com (mail-wm1-f51.google.com [209.85.128.51]) by sourceware.org (Postfix) with ESMTPS id CD7FC38618AB for ; Thu, 14 Dec 2023 20:23:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CD7FC38618AB Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=palves.net Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=gmail.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org CD7FC38618AB Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=209.85.128.51 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702585387; cv=none; b=gaCtF/YLM1DYTfqc3gNwMICga3Wi1HXnjgxugWc6zVWtllXwkBoPSgIT5Iffa92+ymy7ZAilESYHeCm+2+8foSCbdVGafyxy6SdzpJy0f1eYDDTYHYqCNdL4Jks7RdWVzZz8z6AX1o9XhdN4G/i3QdZaH4F+SpLUL2Nv1mA9oF0= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1702585387; c=relaxed/simple; bh=LZnyR8ApVvwsa2D2h+0tT8iukG/p2mF9a8o3o3zdCm4=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=qb3YUErkELI+FxIRHF3/dJHTwUrCzxix+JxPsIH5QmbY7X70ti3ri6ZwhSbcV86eVUQzt05EWrw6QVS61q/hq2uq/vp+VH6wl46VYPU6Y5P4BFvjIm3GXwfRirawxDcIoDb9LQVGFma8OfYJPrEARSEqjgs6D08rWghnJ049k3E= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-f51.google.com with SMTP id 5b1f17b1804b1-40c2db2ee28so91858385e9.2 for ; Thu, 14 Dec 2023 12:23:05 -0800 (PST) X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1702585384; x=1703190184; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=eROZZW26mn+2YuM6OcsIO1ppUF8zBUFQTgYW7snKLnw=; b=GdJgsW7gP2f3Wbq3Na+fgMQEvpSpibPRl5CeMTL3iL04Ed0PWF+m7GQi2jaFSMD6wR WwU5Plek3IC48UpDKSkMAkYRRk7jKw4mMRfX943GAqO+0hGbq3osnvMY9AqFcupo0V9d sZpz/q8il8vHQHjx/5gKdY9hZ+VYigIuDV0oAbpffxSNRDzf+hqAlf+tXtS8pblY0o3r sofZi/kwoNNbi2szETbsTZu3FfL48xqUWlDdmxr4gS1W3hB6ehOXSkhhnHfElWewGQgy VFewzDBxey+HluIAzb7F+VAMsJVCabJP9G9fwsNNL0f54gQG/WIs2XLXIauR36yA/+sK fdyA== X-Gm-Message-State: AOJu0Yy5a5Jfes8X+/FdzxxypWbmjjBkvZuOWo4jDbJfhO23G0Yh6h+y zlx1mXNPjXy31CRJv5900lEC2eOrbgq5Tw== X-Google-Smtp-Source: AGHT+IFFKXNfBGxSJvCPodOVvn/pinM3ezUNwZYoDfT72k6TwIIVel9PXkc+hIhXdGyOrnM080KbsA== X-Received: by 2002:a05:600c:21c7:b0:40c:2699:b625 with SMTP id x7-20020a05600c21c700b0040c2699b625mr6076150wmj.138.1702585384210; Thu, 14 Dec 2023 12:23:04 -0800 (PST) Received: from localhost ([2001:8a0:f923:4f00:2646:535c:5a04:e380]) by smtp.gmail.com with UTF8SMTPSA id gw18-20020a05600c851200b004063c9f68f2sm25507262wmb.26.2023.12.14.12.23.03 for (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Thu, 14 Dec 2023 12:23:03 -0800 (PST) From: Pedro Alves To: gdb-patches@sourceware.org Subject: [PATCH 8/8] Add tests for handling of vanishing threads that were stepping/stopping Date: Thu, 14 Dec 2023 20:22:38 +0000 Message-ID: <20231214202238.1065676-9-pedro@palves.net> X-Mailer: git-send-email 2.43.0 In-Reply-To: <20231214202238.1065676-1-pedro@palves.net> References: <20231214202238.1065676-1-pedro@palves.net> MIME-Version: 1.0 X-Spam-Status: No, score=-10.3 required=5.0 tests=BAYES_00, FREEMAIL_FORGED_FROMDOMAIN, FREEMAIL_FROM, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, RCVD_IN_DNSWL_NONE, RCVD_IN_MSPIKE_H2, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org 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+patchwork=sourceware.org@sourceware.org Not for commit. This won't work with current upstream, unfortunately. Change-Id: I43a66f060c35aad1fe0d9ff022ce2afd0537f028 --- .../gdb.rocm/continue-over-kernel-exit.cpp | 66 +++++++ .../gdb.rocm/continue-over-kernel-exit.exp | 165 ++++++++++++++++++ 2 files changed, 231 insertions(+) create mode 100644 gdb/testsuite/gdb.rocm/continue-over-kernel-exit.cpp create mode 100644 gdb/testsuite/gdb.rocm/continue-over-kernel-exit.exp diff --git a/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.cpp b/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.cpp new file mode 100644 index 00000000000..bad7064f30b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.cpp @@ -0,0 +1,66 @@ +/* Copyright (C) 2023 Free Software Foundation, Inc. + Copyright (C) 2023 Advanced Micro Devices, Inc. All rights reserved. + + This file is part of GDB. + + 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 + +#define CHECK(cmd) \ + do \ + { \ + hipError_t error = cmd; \ + if (error != hipSuccess) \ + { \ + fprintf (stderr, "error: '%s'(%d) at %s:%d\n", \ + hipGetErrorString (error), error, \ + __FILE__, __LINE__); \ + exit (EXIT_FAILURE); \ + } \ + } while (0) + +__global__ void +kern () +{ + asm ("s_endpgm_insn: s_endpgm"); +} + +__global__ void +second_kernel () +{ +} + +int +main () +{ + /* Use 1-thread blocks to easily control number of waves. */ + size_t blocksize = 1; + size_t gridsize = 10; + + kern<<>> (); + + /* Stopping at this second kernel after the first kernel completely + finishes makes GDB refresh its thread list while the + amd-dbgapi-target is still active, which triggers different code + paths in GDB that lead to deleting exited threads. We test both + stopping here, and not stopping here. */ + second_kernel<<<1, 1>>> (); + + CHECK (hipDeviceSynchronize ()); + + return 0; +} diff --git a/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.exp b/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.exp new file mode 100644 index 00000000000..ed98ab4697b --- /dev/null +++ b/gdb/testsuite/gdb.rocm/continue-over-kernel-exit.exp @@ -0,0 +1,165 @@ +# Copyright (C) 2023 Free Software Foundation, Inc. +# Copyright (C) 2023 Advanced Micro Devices, Inc. All rights reserved. + +# This file is part of GDB. + +# 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 . + +# Test stopping at and continuing from a s_endpgm instruction, with +# and without stepping over a breakpoint on top of it. + +load_lib rocm.exp + +standard_testfile .cpp + +require allow_hipcc_tests + +if { [build_executable "failed to prepare" \ + $testfile $srcfile {debug hip}] == -1 } { + return -1 +} + +# Test continuing from an s_endpgm instruction with multiple waves. +# +# STEP_OVER_MODE can be one of: +# +# - none: don't put a breakpoint on the s_endpgm instruction. +# +# - inline: put a breakpoint on the s_endpgm instruction, and use +# in-line stepping to step over it (disable displaced-stepping). +# +# - displaced: same, but use displaced stepping. +# +# FINISH_HOW can be one of: +# +# - second_kernel: stop at a breakpoint at "second_kernel" after the +# first kernel finishes. +# +# - normal_exit: let the inferior exit normally after the first +# kernel finishes, without stopping at second_kernel. +# +proc do_test { step_over_mode finish_how } { + with_rocm_gpu_lock { + clean_restart $::binfile + + if { $step_over_mode == "none" } { + # Nothing to do. + } elseif { $step_over_mode == "inline" } { + gdb_test_no_output "set displaced-stepping off" + } elseif { $step_over_mode == "displaced" } { + gdb_test_no_output "set displaced-stepping on" + } else { + error "Invalid step_over_mode value: $step_over_mode" + } + + if ![runto_main] { + fail "can't run to main" + return -1 + } + + # Put a breakpoint on the s_endpgm instruction, and continue + # to it. If testing a step-over is requested, leave the + # breakpoint inserted after the initial hit to force a + # step-over. + if { $step_over_mode != "none" } { + gdb_breakpoint "s_endpgm_insn" allow-pending + set bp_hits 10 + } else { + gdb_breakpoint "s_endpgm_insn" allow-pending temporary + set bp_hits 0 + } + gdb_continue_to_breakpoint "s_endpgm_insn" + + gdb_test "x/i \$pc" \ + "$::hex <\[^\r\n\]*>:\[ \t\]+s_endpgm.*" \ + "stopped at s_endpgm" + + if {$finish_how == "second_kernel"} { + gdb_breakpoint "second_kernel" + } + + for {set i 1} {$i < $bp_hits} {incr i} { + with_test_prefix "iter $i" { + gdb_test_multiple "continue" "continue to s_endpgm" { + -re -wrap "Continuing\\.\r\n.*hit Breakpoint $::decimal, .* kern .*\"s_endpgm_insn: .*" { + pass $gdb_test_name + } + } + } + } + + # GDB used to mishandle wave exits resulting in + # WAVE_COMMAND_TERMINATED events being left in the + # amd-dbgapi-target's event queue _after_ the wave had already + # been removed from the GDB thread list. That in turn would + # result in seeing already-dead waves re-added to GDB's thread + # list, and then immediately deleted, with these user + # notifications: + # + # [New AMDGPU Wave ?:?:?:1 (?,?,?)/?] + # [AMDGPU Wave ?:?:?:1 (?,?,?)/? exited] + # + # The test below fails if we see any "?" in the wave's target + # id. It is written this way instead of a tighter match to + # increase the chances of the problem being caught if the GDB + # output ever changes. + # + # The bad thread additions mentioned above were done without + # adding the amd-dbgapi-target-specific wave info to the + # amd-dbgapi-target data structures, resulting in GDB + # assertion failures in the amd-dbgapi-target if GDB handled + # any other stop event after the first kernel finishes. We + # exercise that with FINISH_HOW=second_kernel. + + set bad_coords 0 + set exited_normally 0 + set second_kernel_breakpoint 0 + gdb_test_multiple "continue" "last continue" -lbl { + -re "AMDGPU Wave (\[^\r\n\]*)(?=\r\n)" { + set wave_coords $expect_out(1,string) + if {[string first "?" $wave_coords] != -1} { + incr bad_coords + } + exp_continue + } + -re "Inferior 1 \\(process $::decimal\\) exited normally" { + incr exited_normally + exp_continue + } + -re "hit Breakpoint $::decimal, with lane 0, second_kernel " { + incr second_kernel_breakpoint + exp_continue + } + -re -wrap "" { + if {$bad_coords > 0} { + fail "$gdb_test_name (bad coords)" + } elseif {[expr ($second_kernel_breakpoint + $exited_normally) > 1]} { + fail "$gdb_test_name (bad finish)" + } elseif {$finish_how == "second_kernel" && $second_kernel_breakpoint != 1} { + fail "$gdb_test_name (no second_kernel breakpoint)" + } elseif {$finish_how == "normal_exit" && $exited_normally != 1} { + fail "$gdb_test_name (not normal exit)" + } else { + pass $gdb_test_name + } + } + } + } +} + +foreach_with_prefix step_over_mode {none inline displaced} { + foreach_with_prefix finish_how {second_kernel normal_exit} { + do_test $step_over_mode $finish_how + } +}