From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 6132 invoked by alias); 7 Aug 2017 17:20:27 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 1164 invoked by uid 89); 7 Aug 2017 17:20:24 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.6 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy= X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 07 Aug 2017 17:20:18 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1delhP-000320-FX from Tom_deVries@mentor.com for gcc-patches@gcc.gnu.org; Mon, 07 Aug 2017 10:20:15 -0700 Received: from [127.0.0.1] (137.202.0.87) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1263.5; Mon, 7 Aug 2017 18:20:10 +0100 To: GCC Patches CC: Thomas Schwinge From: Tom de Vries Subject: [openacc, PR78266, committed] Fix diff_type in expand_oacc_for char iter_type Message-ID: <5ae64941-1e8d-a788-8b95-297aacd9bfd2@mentor.com> Date: Mon, 07 Aug 2017 17:20:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.2.1 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------DC7CDE812031592238E1B3F2" X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-SW-Source: 2017-08/txt/msg00534.txt.bz2 --------------DC7CDE812031592238E1B3F2 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 609 Hi, this fixes PR78266, an openacc PR. When compiling a gang loop with an iteration variable of type 'unsigned char' and 256 gangs: ... #pragma acc parallel loop num_gangs (256) for (unsigned char j = 0; j < 5; j++) .. we run into trouble. The 'diff_type' in expand_oacc_for is set to 'signed char', and we generate f.i.: ... _41 = GOACC_DIM_SIZE (0); _29 = (signed char) _41; ... where _41 is 256, so forwprop2 folds _29 to '0'. The patch fixes this by ensuring that diff_type is chosen big enough in expand_oacc_for. Tested libgomp on x86_64 with nvptx accelerator. Committed. Thanks, - Tom --------------DC7CDE812031592238E1B3F2 Content-Type: text/x-patch; name="0001-Fix-diff_type-in-expand_oacc_for-char-iter_type.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename="0001-Fix-diff_type-in-expand_oacc_for-char-iter_type.patch" Content-length: 3405 Fix diff_type in expand_oacc_for char iter_type 2017-08-07 Tom de Vries PR middle-end/78266 * omp-expand.c (expand_oacc_for): Ensure diff_type is large enough. * testsuite/libgomp.oacc-c-c++-common/vprop-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vprop.c: Remove xfail. --- gcc/omp-expand.c | 2 + .../testsuite/libgomp.oacc-c-c++-common/vprop-2.c | 45 ++++++++++++++++++++++ .../testsuite/libgomp.oacc-c-c++-common/vprop.c | 1 - 3 files changed, 47 insertions(+), 1 deletion(-) diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 970e04f..1eef7c0 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -5328,6 +5328,8 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) } if (POINTER_TYPE_P (diff_type) || TYPE_UNSIGNED (diff_type)) diff_type = signed_type_for (diff_type); + if (TYPE_PRECISION (diff_type) < TYPE_PRECISION (integer_type_node)) + diff_type = integer_type_node; basic_block entry_bb = region->entry; /* BB ending in OMP_FOR */ basic_block exit_bb = region->exit; /* BB ending in OMP_RETURN */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c new file mode 100644 index 0000000..046ac68 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c @@ -0,0 +1,45 @@ +/* { dg-do run } */ + +#include + +#define DO_PRAGMA(x) _Pragma (#x) + +#define test(idx,type,ngangs) \ + void \ + test_##idx () \ + { \ + int b[100]; \ + \ + for (unsigned int i = 0; i < 100; i++) \ + b[i] = 0; \ + \ + DO_PRAGMA(acc parallel num_gangs (ngangs) copy (b)) \ + { \ + _Pragma("acc loop gang") \ + for (type j = 0; j < 5; j++) \ + { \ + _Pragma("acc loop vector") \ + for (unsigned int i = 0; i < 20; i++) \ + b[j * 20 + i] = -2; \ + } \ + } \ + \ + for (unsigned int i = 0; i < 100; i++) \ + assert (b[i] == -2); \ + } + +test (0, signed char, 256) +test (1, unsigned char, 256) +test (2, signed short, 65535) +test (3, unsigned short, 65535) + +int +main () +{ + test_0 (); + test_1 (); + test_2 (); + test_3 (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c index 0ac0cf6..e4dd682 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c @@ -1,5 +1,4 @@ /* { dg-do run } */ -/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } } */ #include --------------DC7CDE812031592238E1B3F2--