From patchwork Fri Dec 13 14:34:56 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1209171 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-515891-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="c2UGzbzv"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 47ZCrK57mNz9sPn for ; Sat, 14 Dec 2019 01:35:28 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=ZvBbDeOwkOOquzAFBfV3h4epgULHKik6COgXSnteddQ/kABYEZ B4YXRY0nuZ2brn3Q1o8f6+oWE9cV5rvOvB8TAdZO0PSF9iIeXHelxYpsDKqLtMhb enUJPad4iFw0FjZcGdPr1NVpmnxpZf9OAVmZ90Tw7kiCjK32VBu0pgiAc= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to :from:subject:message-id:date:mime-version:content-type; s= default; bh=aVr5IbX26H5box74PEkOmUdC0LM=; b=c2UGzbzvpr+nSDYJ2kFx PZaYxjKDgsbsDfO0NoqzYvz1WTAIp2IMcoBmbMRH2a8OCvikU6dJG5zHfifgGBzK iq1RkTFvroS50kRX8mL9NnVyl0WbYYdZ4a/3rEKE00EfsxlHn6AqtNZh6p+j1onc fVlwCvliA85VWXiQGogDC6M= Received: (qmail 106346 invoked by alias); 13 Dec 2019 14:35:20 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 106302 invoked by uid 89); 13 Dec 2019 14:35:19 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-16.8 required=5.0 tests=AWL, BAYES_00, GARBLED_SUBJECT, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=Suggestions, sk:BUILDIN, sk:offload, cherry X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 13 Dec 2019 14:35:17 +0000 IronPort-SDR: um5QUuqgx2FguEI4Tj5iKf8O+kE1147Ulr9t8CpEHTvj+F7tn8FlpOzj7E3he42fJz4a3IYRDI W58XkJSrPwkL+ylIzFmE54LjrwSRypVLQFte5EwqBFM1RNX59hmCPz6v1gLRteMyUbTLYbnGiS 1jVlJuLF+SfPZjT7quLVzUvdcX0FW7DD4+PsnDz6UKdM8yXXsVWkEshDo3eCjkx7QF2MsAx4+2 dMGIPZ4hGsbMTKLrd5ruwwPbm94/0o0+j6cb5HyGaBXqHYeD5nBUW70puaDGPg2YbdsOrciP8r ntw= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 13 Dec 2019 06:35:15 -0800 IronPort-SDR: x4sJ9BjDXBXnBVNjV3gghf+zySJwvwkhbXqRaRCVTp0nrocmqhggNWBq7diVLCRZ4y/SlvhwPj W7PKC2aYeX30FWx/HzvKk08K6zJNMIgp48a+2CBeq+snwxFI9I3iXFD9cukoThYxcfL07ibHjO H8Ao3uy5LwOz3xhEaHLbgXb4zHi3PId2FBUihWKs1fQzJsSsIG5hUs+ry+GiYxQW5sZaljxn5o +p1uyro20Yrx9yRaRWJtBQLCQySj6ee94orldlj5UcaKUbOsStnCkbQZg+aK736sUU4bV1S4qG beE= To: Jakub Jelinek , Richard Biener , gcc-patches From: Tobias Burnus Subject: =?utf-8?q?=5BLTO=5D_PR_86416_=E2=80=93_improve_lto1_diagnostic_i?= =?utf-8?q?f_a_mode_does_not_exist_=28esp=2E_for_offloading_targets?= =?utf-8?q?=29?= Message-ID: Date: Fri, 13 Dec 2019 15:34:56 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.2.2 MIME-Version: 1.0 X-IsSubscribed: yes As long as one compiles for a single target, the message is unlikely to appear. However, when compiling for offloading, the modes supported on the target 'host' and on the target 'device' can be different. In particular, 'long double' (when larger than double) and '__float128' might not be available on the device. This gives currently errors like the following (see PR, comment 0): > lto1: fatal error: unsupported mode TF > > compilation terminated. > mkoffload: fatal error: x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status While the device target is hidden in 'x86_64-pc-linux-gnu-accel-nvptx-none-gcc', it might make more sense to add it more prominently. Additionally, an average user does not understand what 'TF' or 'XF' means. Solutions: (A) Add the target to the output (B) Add a better description for the error I did both in the attached patch, giving: lto1: fatal error: nvptx-none - 80-bit floating-point numbers unsupported (mode 'XF') lto1: fatal error: nvptx-none - 128-bit floating-point numbers unsupported (mode 'TF') * (A) should be fine, I think. * But I am not 100% happy with (B). I think as interim solution, it is acceptable as XF/TF are well defined and probably the most common problem. — Alternatively, one could only commit (A) or solve it more properly (how?). * If, e.g., 'long long' or 'integer(kind=16)' are used, only a generic message is printed. A message such as "'__float128' not supported" or "'real(kind=10)' not supported" is more helpful and supporting all modes and not cherry picking those two would be useful as well. The question is how to pass this information to lto-streamer-in.c; it is available as TYPE_NAME – and, with debugging turned on, this also gets passed on, but is also not be readily available in lto_input_mode_table. – Suggestions? Build on x86-64-gnu-linux and tested without offloading and with nvptx offloading. Tobias PR middle-end/86416 * Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on. * lto-streamer-in.c (lto_input_mode_table): Use it; add special error diagnostic for missing XF and TF modes. PR middle-end/86416 * testsuite/libgomp.c/pr86416-1.c: New. * testsuite/libgomp.c/pr86416-2.c: New. diff --git a/gcc/Makefile.in b/gcc/Makefile.in index 6b857bd75de..657488d416b 100644 --- a/gcc/Makefile.in +++ b/gcc/Makefile.in @@ -2244,6 +2244,8 @@ version.o: $(REVISION) $(DATESTAMP) $(BASEVER) $(DEVPHASE) # lto-compress.o needs $(ZLIBINC) added to the include flags. CFLAGS-lto-compress.o += $(ZLIBINC) +CFLAGS-lto-streamer-in.o += -DTARGET_MACHINE=\"$(target_noncanonical)\" + bversion.h: s-bversion; @true s-bversion: BASE-VER echo "#define BUILDING_GCC_MAJOR `echo $(BASEVER_c) | sed -e 's/^\([0-9]*\).*$$/\1/'`" > bversion.h diff --git a/gcc/lto-streamer-in.c b/gcc/lto-streamer-in.c index 128d7640726..f59e8c5363d 100644 --- a/gcc/lto-streamer-in.c +++ b/gcc/lto-streamer-in.c @@ -1700,7 +1700,19 @@ lto_input_mode_table (struct lto_file_decl_data *file_data) } /* FALLTHRU */ default: - fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname); + /* For offloading-target compilions, this is a user-facing + message. See also target.def and machmode.def. */ + if (strcmp (mname, "XF") == 0) + fatal_error (UNKNOWN_LOCATION, + "%s - 80-bit floating-point numbers unsupported " + "(mode %qs)", TARGET_MACHINE, mname); + else if (strcmp (mname, "TF") == 0) + fatal_error (UNKNOWN_LOCATION, + "%s - 128-bit floating-point numbers unsupported " + "(mode %qs)", TARGET_MACHINE, mname); + else + fatal_error (UNKNOWN_LOCATION, "%s - unsupported mode %qs", + TARGET_MACHINE, mname); break; } } diff --git a/libgomp/testsuite/libgomp.c/pr86416-1.c b/libgomp/testsuite/libgomp.c/pr86416-1.c new file mode 100644 index 00000000000..c7a162d2f41 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr86416-1.c @@ -0,0 +1,22 @@ +/* { dg-do link } */ +/* { dg-require-effective-target large_long_double } */ + +/* PR middle-end/86416 */ +/* { dg-error "80-bit floating-point numbers unsupported .mode 'XF'." "" { target offload_device } 0 } */ +/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */ + +#include /* For abort. */ + +long double foo (long double x) +{ + #pragma omp target map(tofrom:x) + x *= 2.0; + return x; +} + +int main() +{ + long double v = foo (10.0q) - 20.0q; + if (v > 1.0e-5 || v < -1.0e-5) abort(); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr86416-2.c b/libgomp/testsuite/libgomp.c/pr86416-2.c new file mode 100644 index 00000000000..75e9cd773a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr86416-2.c @@ -0,0 +1,22 @@ +/* { dg-do link { target __float128 } } */ +/* { dg-add-options __float128 } */ + +/* PR middle-end/86416 */ +/* { dg-error "128-bit floating-point numbers unsupported .mode 'TF'." "" { target offload_device } 0 } */ +/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } } */ + +#include /* For abort. */ + +__float128 foo(__float128 y) +{ + #pragma omp target map(tofrom: y) + y *= 4.0; + return y; +} + +int main() +{ + __float128 v = foo (5.0L) - 20.0L; + if (v > 1.0e-5 || v < -1.0e-5) abort(); + return 0; +}