# 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 "remote debugging"}
    }

    if {![istarget "*-linux*"]} {
	return {0 "target platform is not Linux"}
    }

    # 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 "amd-dbgapi not supported"}
    }

    # Check we have a working hipcc compiler available.
    set targets [hcc_amdgpu_targets]
    if { [llength $targets] == 0} {
	return {0 "no suitable amdgpu targets found"}
    }

    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>>> ();
		if (hipDeviceSynchronize () != hipSuccess)
		  return -1;
		return 0;
	    }
	} executable $flags]} {
	return {0 "failed to compile hip program"}
    }

    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
    }
}

# Return true if all the devices support debugging multiple processes
# using the GPU.

proc hip_devices_support_debug_multi_process {} {
    set unsupported_targets \
	{gfx900 gfx906 gfx908 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032}

    set targets [hcc_amdgpu_targets]
    if { [llength $targets] == 0 } {
	return 0
    }

    foreach target $targets {
	if { [lsearch -exact $unsupported_targets $target] != -1 } {
	    return 0
	}
    }
    return 1
}

# Return true if all the devices on the host support precise memory.

proc hip_devices_support_precise_memory {} {
    set unsupported_targets \
	{gfx900 gfx906 gfx908 gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032}

    set targets [hcc_amdgpu_targets]
    if { [llength $targets] == 0 } {
	return 0
    }

    foreach target $targets {
	if { [lsearch -exact $unsupported_targets $target] != -1 } {
	    return 0
	}
    }
    return 1
}