1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
|
# Copyright (C) 2019-2023 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/>.
#
# Support library for testing ROCm (AMD GPU) GDB features.
# 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
}
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.
if {[target_info gdb_protocol] != ""} {
return 0
}
# Ensure that GDB is built with amd-dbgapi support.
set output [remote_exec host $::GDB "$::INTERNAL_GDBFLAGS --configuration"]
if { [string first "--with-amd-dbgapi" $output] == -1 } {
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 <hip/hip_runtime.h>
__global__ void
kern () {}
int
main ()
{
kern<<<1, 1>>> ();
hipDeviceSynchronize ();
return 0;
}
} executable $flags]} {
return 0
}
return 1
}
# The lock file used to ensure that only one GDB has access to the GPU
# at a time.
set gpu_lock_filename $objdir/gpu-parallel.lock
# Acquire lock file LOCKFILE. Tries forever until the lock file is
# successfully created.
proc lock_file_acquire {lockfile} {
verbose -log "acquiring lock file: $::subdir/${::gdb_test_file_name}.exp"
while {true} {
if {![catch {open $lockfile {WRONLY CREAT EXCL}} rc]} {
set msg "locked by $::subdir/${::gdb_test_file_name}.exp"
verbose -log "lock file: $msg"
# For debugging, put info in the lockfile about who owns
# it.
puts $rc $msg
flush $rc
return [list $rc $lockfile]
}
after 10
}
}
# Release a lock file.
proc lock_file_release {info} {
verbose -log "releasing lock file: $::subdir/${::gdb_test_file_name}.exp"
if {![catch {fconfigure [lindex $info 0]}]} {
if {![catch {
close [lindex $info 0]
file delete -force [lindex $info 1]
} rc]} {
return ""
} else {
return -code error "Error releasing lockfile: '$rc'"
}
} else {
error "invalid lock"
}
}
# Run body under the GPU lock. Also calls gdb_exit before releasing
# the GPU lock.
proc with_rocm_gpu_lock { body } {
if {[info exists ::GDB_PARALLEL]} {
set lock_rc [lock_file_acquire $::gpu_lock_filename]
}
set code [catch {uplevel 1 $body} result]
# In case BODY returned early due to some testcase failing, and
# left GDB running, debugging the GPU.
gdb_exit
if {[info exists ::GDB_PARALLEL]} {
lock_file_release $lock_rc
}
if {$code == 1} {
global errorInfo errorCode
return -code $code -errorinfo $errorInfo -errorcode $errorCode $result
} else {
return -code $code $result
}
}
|
