git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

amdgcn, libgomp: improve generic device errors 

Switching to use "generic" ISA variants has changed the error modes a bit.

This patch changes the runtime so that it doesn't say to use the
device-specific -march option when the real problem is not the ISA (it'll be a
mismatched xnack setting, probably).

Additionally, the testsuite effective target check needs to see if the xnack
mode is accepted by the runtime, as well as the compiler.

libgomp/ChangeLog:

	* plugin/plugin-gcn.c (generic_isa_code): New function.
	(isa_matches_agent): Use generic ISA details to help select an error
	message on ISA mismatch.
	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_target_amdgcn_with_xnack): Use a
	runtime check.
This commit is contained in:
Andrew Stubbs 2025-12-03 15:21:30 +00:00
parent 56c6612598
commit 1d1d12da6d
2 changed files with 20 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
libgomp/plugin/plugin-gcn.c
View File

@ -1776,6 +1776,22 @@ isa_code(const char *isa) {
return EF_AMDGPU_MACH_UNSUPPORTED;
}
/* Returns the code which is used in the GCN object code to identify the
generic ISA that corresponds to a specific ISA. */
static gcn_isa
generic_isa_code (int isa) {
switch(isa)
{
#define EF_AMDGPU_MACH_AMDGCN_NONE 0
#define GCN_DEVICE(name, NAME, ELF, GCCISA, XNACK, SRAM, WAVE64, CUMODE, \
VGPRS, CO, ARCH, GENERIC_ISA, ...) \
case ELF: return EF_AMDGPU_MACH_AMDGCN_ ## GENERIC_ISA;
#include "../../gcc/config/gcn/gcn-devices.def"
}
return 0;
}
/* CDNA2 devices have twice as many VGPRs compared to older devices. */
static int
@ -2551,7 +2567,9 @@ isa_matches_agent (struct agent_info *agent, Elf64_Ehdr *image,
"Consider using ROCR_VISIBLE_DEVICES to disable incompatible "
"devices or run with LOADER_ENABLE_LOGGING=1 for more details.",
device_isa_s, agent_isa_s, agent->device_id);
else if (strcmp (device_isa_s, agent_isa_s) == 0)
else if (strcmp (device_isa_s, agent_isa_s) == 0
|| (elf_gcn_isa_is_generic (image)
&& generic_isa_code (agent->device_isa) == isa_field))
snprintf (msg, sizeof msg,
"GCN code object features do not match for an unknown reason "
"(device %d).\n"

2
libgomp/testsuite/lib/libgomp.exp
View File

@ -758,7 +758,7 @@ proc check_effective_target_omp_managedmem { } {
proc check_effective_target_offload_target_amdgcn_with_xnack { } {
if { [libgomp_check_effective_target_offload_target "amdgcn"] } {
return [check_no_compiler_messages amd_xnack_ executable {
return [check_runtime amd_xnack_ {
int main () {
#pragma omp target
;