From patchwork Wed Jan 24 12:43:04 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Andrew Stubbs X-Patchwork-Id: 191566 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a05:7300:2553:b0:103:945f:af90 with SMTP id p19csp961215dyi; Wed, 24 Jan 2024 04:44:19 -0800 (PST) X-Google-Smtp-Source: AGHT+IHU/XdbKq7BmS1zFs4jC+5YG616Mwdc1KNuO/xzZO6yynEIXwXZh/6Re03CaPWTHetZ0xOi X-Received: by 2002:a05:620a:6229:b0:783:8d13:998e with SMTP id ou41-20020a05620a622900b007838d13998emr7440030qkn.36.1706100259497; Wed, 24 Jan 2024 04:44:19 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1706100259; cv=pass; d=google.com; s=arc-20160816; b=jUs90jqSf/9sI9dVFDC0mKrk+dTh9p2+8Up5MemAXpytyzA1rLuQv41/WlSdihsBpb iZb8UH5cil1o5GPzgPAHM6l98i3uVsaaCFhtcnVBJ7HksZFM67ojF3/nCzKc/ZFOUxmM yWFJRSaDCcYJ2CZJi1V6VL4zKLCdvxD1V0BB1tNu76rBt2ORRliTje5CyzS+/2IXmlJ5 ycU3VjuqbhkRFxOwTmergsNW+ww9v2DBlXbnMXn44Dov5CEMcQ8McND1pmK6kiiK3aqM xd7rFAj1AduBobnRcYFwLYXA6JUqT0Ea4klDRNOTijf+xzdlC4hLLoJT1IIt3KSs1x2E WZwA== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:message-id:date:subject:cc:to:from:dkim-signature :arc-filter:dmarc-filter:delivered-to; bh=mVJtGKKFbsqlb2QqxK8bHN7fVA/OtWNnmC2QVXZPZtg=; fh=i5P1UXIEs2wYl1vak5AL6fwt9PbcC+sB98HlI7I75ZU=; b=DjpwyuNGHQCSB9NJ8lQtnYHi6xgWJqPllD1lY5+Fy2fRbHk9/9m2vwROJfwR+yRLwz Ln59ZHqjLJFLCYb7QWGUwAdgUIAtrKEd360wRiImXvyj0rJyK3R5By4rZdasV68ON7sE dyNon/y+cgNn8NxiQHj65C0PvoodXWICmJb/i02JnOyuj7tOPvtxPIwvC81OOQJeydNu Nh+dK3UfaUUp1U3YcwPrySuOmSELBYBgM2o+pr+mKGVlPGd0fbEkB3CQrrNPLfl+smrn 3ZqgRMBUwxU4LZEuMnpL1RsromWnRBfKjlT/U/tvt/65maAQzq+hMyYhlRD5aK8ysKjR wBgg== ARC-Authentication-Results: i=2; mx.google.com; dkim=pass header.i=@baylibre-com.20230601.gappssmtp.com header.s=20230601 header.b=ssnlA8RA; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id m1-20020a05620a220100b00783100ada16si9876860qkh.259.2024.01.24.04.44.19 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 24 Jan 2024 04:44:19 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; dkim=pass header.i=@baylibre-com.20230601.gappssmtp.com header.s=20230601 header.b=ssnlA8RA; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 305BB3858425 for ; Wed, 24 Jan 2024 12:44:19 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32a.google.com (mail-wm1-x32a.google.com [IPv6:2a00:1450:4864:20::32a]) by sourceware.org (Postfix) with ESMTPS id 55CCB3858C62 for ; Wed, 24 Jan 2024 12:43:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 55CCB3858C62 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 55CCB3858C62 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32a ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1706100211; cv=none; b=IKSWAlZw9qVffefiNJ7P225oRQCxmfoE8yHsmNu3IGUit6EUCDdmqaYcLtfk+r/db19mLkAynRs8ZwENbeAtawDnwXk7LRki/4e89+IW5uNElPOWkgeAlI9zx4G0S3zAz3LMgZqw2AxTeicTcAWQaxcyhcTM96qGw3RB3BMqIok= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1706100211; c=relaxed/simple; bh=KBqSyGEZHhDHlCrZpPdUYvMqFBZb6L2XzkY0Q27+pv8=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=uvMy96bfNsw+7Rjhy3KVTdOApKbdjJpMPk1KBlM12kguBE5ieqKK+7a7A5d7QtdW3m2c1N5X8rHSFHFVK5qsmnjLIuN3wv2QTwOWsmUhyAToQABn7BdMI0a3zlxrwAx0tvLlY5MhL5zu5IckN9YEkGiMI+rU6fKY0QdVoJ3sAkw= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32a.google.com with SMTP id 5b1f17b1804b1-40ec715f1efso4175085e9.2 for ; Wed, 24 Jan 2024 04:43:23 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1706100201; x=1706705001; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:from:to:cc:subject:date:message-id:reply-to; bh=mVJtGKKFbsqlb2QqxK8bHN7fVA/OtWNnmC2QVXZPZtg=; b=ssnlA8RADr/yDqg2C5p5rA7k8zrO6nL/wXri0XdHyemfmlUn/TgG2k7EMmHlNdaNhy rIIP4z7+Sx2bbINYy/E+navOP72AfBfUtabxhlvqK+tr1uOzj9wS+xWed9XbQAt6l3om 4wfaMSYkqWtLFP3zQsSbSYlYINakuMIoZaEGGLbS8W0Kt1De2y2X0B6FSph6xWtl9Nez c2O80g+QbdxS4Vby58ByadSSxnXwq3kYljWvw4wBG3QJBjKo5VRhMBgyjZnKzflBb6ia Hfsc234fljvvNO4u3Ix5WdEOkaBwdiLhBWY+45QM8YMpvWHl5UP4OwMH1JzOFhOnC7nF 1k0Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1706100201; x=1706705001; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:x-gm-message-state:from:to:cc:subject:date:message-id :reply-to; bh=mVJtGKKFbsqlb2QqxK8bHN7fVA/OtWNnmC2QVXZPZtg=; b=v+8CFKLEYQM+5sGvI8z8o+5waLqkUHwSu/82vTDjR5wlRH+a2qqpXAyjJPqmOYdf7P f44KfusCz1CH3HU63gjwjGiygMbtVMN52Ru5+xNjnu5c/sscwoZlDvSslfypDHhYfnEs 9i+UgmBL0Ehf7Emi+Ua70Lng1Zwn/1eWAHG3TA7tBvu7q6aEId0h4N1OLIOgYxsR1hGe dQYhGJ2CvBnBM17tPPHhEopRgYvEnMED6JBxWLw0pYIpLeMKDTncWwOAot9WRSjSeLhh v1pOUbklASeF7LF4Lnw3KXgUHhYPJ7Ap6czDa4huCGW0WZrlJhxhH5yHjtbmUpbY375v X/0g== X-Gm-Message-State: AOJu0Yw7PtcM/5MEGg/KjF2V7k3JtgoVldcu8gdp9RBMzk/HCvUyuVtz vQTKezF8jvGeZYT5/T9XUA8y0ETYTHEJOSSjF6Wx5I4+K0UpWIrzavCZwmGLl8wkv7fr0dYRk01 GzOs= X-Received: by 2002:a7b:c7c6:0:b0:40e:5e29:cca2 with SMTP id z6-20020a7bc7c6000000b0040e5e29cca2mr1483977wmk.44.1706100201071; Wed, 24 Jan 2024 04:43:21 -0800 (PST) Received: from carlos.baylibre (laubervilliers-658-1-213-31.w90-63.abo.wanadoo.fr. [90.63.244.31]) by smtp.googlemail.com with ESMTPSA id v16-20020a5d4b10000000b00337d4eed87asm15685140wrq.115.2024.01.24.04.43.20 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 24 Jan 2024 04:43:20 -0800 (PST) From: Andrew Stubbs To: gcc-patches@gcc.gnu.org Cc: pa@codesourcery.com, rguenther@suse.de Subject: [PATCH] amdgcn: additional gfx1100 support Date: Wed, 24 Jan 2024 12:43:04 +0000 Message-ID: <20240124124304.1780645-1-ams@baylibre.com> X-Mailer: git-send-email 2.41.0 MIME-Version: 1.0 X-Spam-Status: No, score=-11.0 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, 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: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1788975785689741129 X-GMAIL-MSGID: 1788975785689741129 This is enough to get gfx1100 working for most purposes, on top of the patch that Tobias committed a week or so ago; there are still some test failures to investigate, and probably some tuning to do. It might also get gfx1030 working too. @Richi, could you test it, please? I can't test the other multilibs right now. @PA, can you test it please? I can self-approve the patch, but I'll hold off the commit until the test results come back. Andrew gcc/ChangeLog: * config/gcn/gcn-opts.h (TARGET_PACKED_WORK_ITEMS): Add TARGET_RDNA3. * config/gcn/gcn-valu.md (all_convert): New iterator. (2): New define_expand, and rename the old one to ... (*_sdwa): ... this. (extend2): Likewise, to ... (extend_sdwa): .. this. (*_shift): New. * config/gcn/gcn.cc (gcn_global_address_p): Use "offsetbits" correctly. (gcn_hsa_declare_function_name): Update the vgpr counting for gfx1100. * config/gcn/gcn.md (mulhisi3): Disable on RDNA3. (mulqihi3_scalar): Likewise. libgcc/ChangeLog: * config/gcn/amdgcn_veclib.h (CDNA3_PLUS): Handle RDNA3. libgomp/ChangeLog: * config/gcn/time.c (RTC_TICKS): Configure RDNA3. (omp_get_wtime): Add RDNA3-compatible variant. * plugin/plugin-gcn.c (max_isa_vgprs): Tune for gfx1030 and gfx1100. Signed-off-by: Andrew Stubbs --- gcc/config/gcn/gcn-opts.h | 2 +- gcc/config/gcn/gcn-valu.md | 41 ++++++++++++++++++++++++++++--- gcc/config/gcn/gcn.cc | 31 ++++++++++++++++------- gcc/config/gcn/gcn.md | 4 +-- libgcc/config/gcn/amdgcn_veclib.h | 2 +- libgomp/config/gcn/time.c | 10 ++++++++ libgomp/plugin/plugin-gcn.c | 6 +++-- 7 files changed, 77 insertions(+), 19 deletions(-) diff --git a/gcc/config/gcn/gcn-opts.h b/gcc/config/gcn/gcn-opts.h index 79fbda3ab25..6be2c9204fa 100644 --- a/gcc/config/gcn/gcn-opts.h +++ b/gcc/config/gcn/gcn-opts.h @@ -62,7 +62,7 @@ extern enum gcn_isa { #define TARGET_M0_LDS_LIMIT (TARGET_GCN3) -#define TARGET_PACKED_WORK_ITEMS (TARGET_CDNA2_PLUS) +#define TARGET_PACKED_WORK_ITEMS (TARGET_CDNA2_PLUS || TARGET_RDNA3) #define TARGET_XNACK (flag_xnack != HSACO_ATTR_OFF) diff --git a/gcc/config/gcn/gcn-valu.md b/gcc/config/gcn/gcn-valu.md index 3d5b6271ee6..cd027f8b369 100644 --- a/gcc/config/gcn/gcn-valu.md +++ b/gcc/config/gcn/gcn-valu.md @@ -3555,30 +3555,63 @@ ;; }}} ;; {{{ Int/int conversions +(define_code_iterator all_convert [truncate zero_extend sign_extend]) (define_code_iterator zero_convert [truncate zero_extend]) (define_code_attr convop [ (sign_extend "extend") (zero_extend "zero_extend") (truncate "trunc")]) -(define_insn "2" +(define_expand "2" + [(set (match_operand:V_INT_1REG 0 "register_operand" "=v") + (all_convert:V_INT_1REG + (match_operand:V_INT_1REG_ALT 1 "gcn_alu_operand" " v")))] + "") + +(define_insn "*_sdwa" [(set (match_operand:V_INT_1REG 0 "register_operand" "=v") (zero_convert:V_INT_1REG (match_operand:V_INT_1REG_ALT 1 "gcn_alu_operand" " v")))] - "" + "!TARGET_RDNA3" "v_mov_b32_sdwa\t%0, %1 dst_sel: dst_unused:UNUSED_PAD src0_sel:" [(set_attr "type" "vop_sdwa") (set_attr "length" "8")]) -(define_insn "extend2" +(define_insn "extend_sdwa" [(set (match_operand:V_INT_1REG 0 "register_operand" "=v") (sign_extend:V_INT_1REG (match_operand:V_INT_1REG_ALT 1 "gcn_alu_operand" " v")))] - "" + "!TARGET_RDNA3" "v_mov_b32_sdwa\t%0, sext(%1) src0_sel:" [(set_attr "type" "vop_sdwa") (set_attr "length" "8")]) +(define_insn "*_shift" + [(set (match_operand:V_INT_1REG 0 "register_operand" "=v") + (all_convert:V_INT_1REG + (match_operand:V_INT_1REG_ALT 1 "gcn_alu_operand" " v")))] + "TARGET_RDNA3" + { + enum {extend, zero_extend, trunc}; + rtx shiftwidth = (mode == QImode + || mode == QImode + ? GEN_INT (24) + : mode == HImode + || mode == HImode + ? GEN_INT (16) + : NULL); + operands[2] = shiftwidth; + + if (!shiftwidth) + return "v_mov_b32 %0, %1"; + else if ( == extend || == trunc) + return "v_lshlrev_b32\t%0, %2, %1\;v_ashrrev_i32\t%0, %2, %0"; + else + return "v_lshlrev_b32\t%0, %2, %1\;v_lshrrev_b32\t%0, %2, %0"; + } + [(set_attr "type" "mult") + (set_attr "length" "8")]) + ;; GCC can already do these for scalar types, but not for vector types. ;; Unfortunately you can't just do SUBREG on a vector to select the low part, ;; so there must be a few tricks here. diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc index e668ce7c69e..e80de2ce056 100644 --- a/gcc/config/gcn/gcn.cc +++ b/gcc/config/gcn/gcn.cc @@ -1597,8 +1597,8 @@ gcn_global_address_p (rtx addr) rtx offset = XEXP (addr, 1); int offsetbits = (TARGET_RDNA2_PLUS ? 11 : 12); bool immediate_p = (CONST_INT_P (offset) - && INTVAL (offset) >= -(1 << 12) - && INTVAL (offset) < (1 << 12)); + && INTVAL (offset) >= -(1 << offsetbits) + && INTVAL (offset) < (1 << offsetbits)); if ((gcn_address_register_p (base, DImode, false) || gcn_vec_address_register_p (base, DImode, false)) @@ -6597,8 +6597,10 @@ gcn_hsa_declare_function_name (FILE *file, const char *name, if (df_regs_ever_live_p (FIRST_AVGPR_REG + avgpr)) break; avgpr++; - vgpr = (vgpr + 3) & ~3; - avgpr = (avgpr + 3) & ~3; + + /* The main function epilogue uses v8, but df doesn't see that. */ + if (vgpr < 9) + vgpr = 9; if (!leaf_function_p ()) { @@ -6611,9 +6613,18 @@ gcn_hsa_declare_function_name (FILE *file, const char *name, avgpr = MAX_NORMAL_AVGPR_COUNT; } - /* The gfx90a accum_offset field can't represent 0 registers. */ - if (gcn_arch == PROCESSOR_GFX90a && vgpr < 4) - vgpr = 4; + /* SIMD32 devices count double in wavefront64 mode. */ + if (TARGET_RDNA2_PLUS) + vgpr *= 2; + + /* Round up to the allocation block size. */ + int vgpr_block_size = (TARGET_RDNA3 ? 12 + : TARGET_RDNA2_PLUS || TARGET_CDNA2_PLUS ? 8 + : 4); + if (vgpr % vgpr_block_size) + vgpr += vgpr_block_size - (vgpr % vgpr_block_size); + if (avgpr % vgpr_block_size) + avgpr += vgpr_block_size - (avgpr % vgpr_block_size); fputs ("\t.rodata\n" "\t.p2align\t6\n" @@ -6714,12 +6725,14 @@ gcn_hsa_declare_function_name (FILE *file, const char *name, " .private_segment_fixed_size: 0\n" " .wavefront_size: 64\n" " .sgpr_count: %i\n" - " .vgpr_count: %i\n" + " .vgpr_count: %i%s\n" " .max_flat_workgroup_size: 1024\n", cfun->machine->kernarg_segment_byte_size, cfun->machine->kernarg_segment_alignment, LDS_SIZE, - sgpr, next_free_vgpr); + sgpr, next_free_vgpr, + (TARGET_RDNA2_PLUS ? " ; wavefrontsize64 counts double on SIMD32" + : "")); if (gcn_arch == PROCESSOR_GFX90a || gcn_arch == PROCESSOR_GFX908) fprintf (file, " .agpr_count: %i\n", avgpr); fputs (" .end_amdgpu_metadata\n", file); diff --git a/gcc/config/gcn/gcn.md b/gcc/config/gcn/gcn.md index 492b833e255..1f3c692b7a6 100644 --- a/gcc/config/gcn/gcn.md +++ b/gcc/config/gcn/gcn.md @@ -1618,7 +1618,7 @@ (mult:SI (any_extend:SI (match_operand:HI 1 "register_operand" "%v")) (any_extend:SI (match_operand:HI 2 "register_operand" " v"))))] - "" + "!TARGET_RDNA3" "v_mul_32_24_sdwa\t%0, %1, %2 src0_sel:WORD_0 src1_sel:WORD_0" [(set_attr "type" "vop_sdwa") (set_attr "length" "8")]) @@ -1628,7 +1628,7 @@ (mult:HI (any_extend:HI (match_operand:QI 1 "register_operand" "%v")) (any_extend:HI (match_operand:QI 2 "register_operand" " v"))))] - "" + "!TARGET_RDNA3" "v_mul_32_24_sdwa\t%0, %1, %2 src0_sel:BYTE_0 src1_sel:BYTE_0" [(set_attr "type" "vop_sdwa") (set_attr "length" "8")]) diff --git a/libgcc/config/gcn/amdgcn_veclib.h b/libgcc/config/gcn/amdgcn_veclib.h index 821f6386dd6..d268c6cac16 100644 --- a/libgcc/config/gcn/amdgcn_veclib.h +++ b/libgcc/config/gcn/amdgcn_veclib.h @@ -230,7 +230,7 @@ do { \ #if defined (__GCN3__) || defined (__GCN5__) \ || defined (__CDNA1__) || defined (__CDNA2__) \ - || defined (__RDNA2__) + || defined (__RDNA2__) || defined (__RDNA3__) #define CDNA3_PLUS 0 #else #define CDNA3_PLUS 1 diff --git a/libgomp/config/gcn/time.c b/libgomp/config/gcn/time.c index 30a0d0188e4..efcd04f5f43 100644 --- a/libgomp/config/gcn/time.c +++ b/libgomp/config/gcn/time.c @@ -30,15 +30,25 @@ /* According to AMD: dGPU RTC is 27MHz AGPU RTC is 100MHz + RDNA3 ISA manual states "typically 100MHz" FIXME: DTRT on an APU. */ +#ifdef __RDNA3__ +#define RTC_TICKS (1.0 / 100000000.0) /* 100MHz */ +#else #define RTC_TICKS (1.0 / 27000000.0) /* 27MHz */ +#endif double omp_get_wtime (void) { uint64_t clock; +#ifdef __RDNA3__ + asm ("s_sendmsg_rtn_b64 %0 0x83 ;Get REALTIME\n\t" + "s_waitcnt 0" : "=r" (clock)); +#else asm ("s_memrealtime %0\n\t" "s_waitcnt 0" : "=r" (clock)); +#endif return clock * RTC_TICKS; } diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index 0339848451e..db28781dedb 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -1741,11 +1741,13 @@ max_isa_vgprs (int isa) case EF_AMDGPU_MACH_AMDGCN_GFX900: case EF_AMDGPU_MACH_AMDGCN_GFX906: case EF_AMDGPU_MACH_AMDGCN_GFX908: - case EF_AMDGPU_MACH_AMDGCN_GFX1030: - case EF_AMDGPU_MACH_AMDGCN_GFX1100: return 256; case EF_AMDGPU_MACH_AMDGCN_GFX90a: return 512; + case EF_AMDGPU_MACH_AMDGCN_GFX1030: + return 512; /* 512 SIMD32 = 256 wavefrontsize64. */ + case EF_AMDGPU_MACH_AMDGCN_GFX1100: + return 1536; /* 1536 SIMD32 = 768 wavefrontsize64. */ } GOMP_PLUGIN_fatal ("unhandled ISA in max_isa_vgprs"); }