From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from simark.ca (simark.ca [158.69.221.121]) by sourceware.org (Postfix) with ESMTPS id 223BD3858D33 for ; Tue, 7 Feb 2023 14:12:09 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 223BD3858D33 Authentication-Results: sourceware.org; dmarc=pass (p=none dis=none) header.from=simark.ca Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=simark.ca Received: from [10.0.0.11] (unknown [217.28.27.60]) (using TLSv1.3 with cipher TLS_AES_128_GCM_SHA256 (128/128 bits) key-exchange X25519 server-signature RSA-PSS (2048 bits)) (No client certificate requested) by simark.ca (Postfix) with ESMTPSA id 8B5E41E0D3; Tue, 7 Feb 2023 09:12:08 -0500 (EST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=simark.ca; s=mail; t=1675779128; bh=yVWmFwpi1oBRTqDZqsxiArLpGPkLXRSax5k3fEDZmFw=; h=Date:Subject:To:Cc:References:From:In-Reply-To:From; b=CKo/dmb/lq+qpEqob7hsgX49dQZcjOxrr8w6r08zzDRP1JmotYXyQkUQ7CSQK6x/a EZUbUCyRK6Z+MNY+gMCkTqZNzIiqoCXmgPKo0vrieDTFt8RyAWDE2uSbikK0BD4nZG IVZ5Zq1TJ7Uprv2B1ORBiFcX7kcGVACWzjOe9VSA= Message-ID: <026f2c76-47fe-9115-b1f2-ec4e0de6eefd@simark.ca> Date: Tue, 7 Feb 2023 09:12:08 -0500 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.7.1 Subject: Re: [PATCH 4/4] gdb/testsuite: allow_hipcc_tests tests the hipcc compiler Content-Language: en-US To: Lancelot SIX , gdb-patches@sourceware.org Cc: lsix@lancelotsix.com References: <20230207132802.223510-1-lancelot.six@amd.com> <20230207132802.223510-5-lancelot.six@amd.com> From: Simon Marchi In-Reply-To: <20230207132802.223510-5-lancelot.six@amd.com> Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 7bit X-Spam-Status: No, score=-11.4 required=5.0 tests=BAYES_00,DKIM_SIGNED,DKIM_VALID,DKIM_VALID_AU,DKIM_VALID_EF,GIT_PATCH_0,NICE_REPLY_A,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: On 2/7/23 08:28, Lancelot SIX via Gdb-patches wrote: > Update allow_hipcc_tests so all gdb.rocm tests are skipped if we do not > have a working hipcc compiler available. > > To achieve this, adjust gdb_simple_compile to ensure that the hip > program is saved in a ".cpp" file before calling hipcc otherwise > compilation will fail. > > One thing to note is that it is possible to have a hipcc installed with > a CUDA backend. Compiling with this back-end will successfully result > in an application, but GDB cannot debug it (at least for the offload > part). In the context of the gdb.rocm tests, we want to detect such > situation where gdb_simple_compile would give a false positive. > > To achieve this, this patch checks that there is at least one AMDGPU > device available and that hipcc can compile for this or those targets. > Detecting the device is done using the rocm_agent_enumerator tool which > is installed with the all ROCm installations (it is used by hipcc to > detect identify targets if this is not specified on the comand line). > > This patch also makes the allow_hipcc_tests proc a cached proc. > --- > gdb/testsuite/lib/gdb.exp | 4 +++ > gdb/testsuite/lib/rocm.exp | 69 +++++++++++++++++++++++++++++++++++++- > 2 files changed, 72 insertions(+), 1 deletion(-) > > diff --git a/gdb/testsuite/lib/gdb.exp b/gdb/testsuite/lib/gdb.exp > index faa0ac05a9a..6333728f71e 100644 > --- a/gdb/testsuite/lib/gdb.exp > +++ b/gdb/testsuite/lib/gdb.exp > @@ -4581,6 +4581,10 @@ proc gdb_simple_compile {name code {type object} {compile_flags {}} {object obj} > set ext "go" > break > } > + if { "$flag" eq "hip" } { > + set ext "cpp" > + break > + } > } > set src [standard_temp_file $name-[pid].$ext] > set obj [standard_temp_file $name-[pid].$postfix] > diff --git a/gdb/testsuite/lib/rocm.exp b/gdb/testsuite/lib/rocm.exp > index b5b59748c27..06d5b3988b8 100644 > --- a/gdb/testsuite/lib/rocm.exp > +++ b/gdb/testsuite/lib/rocm.exp > @@ -15,7 +15,51 @@ > # > # Support library for testing ROCm (AMD GPU) GDB features. > > -proc allow_hipcc_tests { } { > +# Get the list of gpu targets to compile for. > +# > +# If HCC_AMDGPU_TARGET is set in the environment, use it. Otherwise, > +# try reading it from the system using the rocm_agent_enumerator > +# utility. > + > +proc hcc_amdgpu_targets {} { > + # Look for HCC_AMDGPU_TARGET (same env var hipcc uses). If > + # that fails, try using rocm_agent_enumerator (again, same as > + # hipcc does). > + if {[info exists ::env(HCC_AMDGPU_TARGET)]} { > + return [split $::env(HCC_AMDGPU_TARGET) ","] > + } > + > + set rocm_agent_enumerator "rocm_agent_enumerator" > + > + # If available, use ROCM_PATH to locate rocm_agent_enumerator. > + if { [info exists ::env(ROCM_PATH)] } { > + set rocm_agent_enumerator \ > + "$::env(ROCM_PATH)/bin/rocm_agent_enumerator" > + } > + > + # If we fail to locate the rocm_agent_enumerator, just return an empty > + # list of targets and let the caller decide if this should be an error. > + if { [which $rocm_agent_enumerator] == 0 } { > + return [list] > + } > + > + set result [remote_exec host $rocm_agent_enumerator] > + if { [lindex $result 0] != 0 } { > + error "rocm_agent_enumerator failed" > + } > + > + set targets [list] > + foreach target [lindex $result 1] { > + # Ignore gfx000 which is the host CPU. > + if { $target ne "gfx000" } { > + lappend targets $target > + } > + } > + > + return $targets I typically don't have ROCM_PATH set, and I don't add /opt/rocm/bin to my PATH, so rocm_agent_enumerator will not be found by default here. However, gdb_find_hipcc does fall back on /opt/rocm/bin, so it does find my hipcc. I think we should use similar strategies for both cases. I don't mind having to set ROCM_PATH or adding /opt/rocm/bin to my PATH if needed, as long as we're consistent about it. > +} > + > +gdb_caching_proc allow_hipcc_tests { > # Only the native target supports ROCm debugging. E.g., when > # testing against GDBserver, there's no point in running the ROCm > # tests. > @@ -29,6 +73,29 @@ proc allow_hipcc_tests { } { > return 0 > } > > + # Check we have a working hipcc compiler available. > + set targets [hcc_amdgpu_targets] > + if { [llength $targets] == 0} { > + return 0 > + } > + > + set flags [list hip additional_flags=--offload-arch=[join $targets ","]] > + if {![gdb_simple_compile hipprobe { > + #include > + __global__ void > + kern () {} > + > + int > + main () > + { > + kern<<<1, 1>>> (); > + hipDeviceSynchronize (); > + return 0; > + } > + } executable $flags]} { > + return 0 > + } So, this last part ensures we don't have a "CUDA hipcc false positive"? It would be good to put a comment above it to indicate the precise intent. Simon