From patchwork Mon Aug 7 17:20:06 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 798791 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-459971-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="epeFqVFz"; 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 3xR46s5rhHz9s7C for ; Tue, 8 Aug 2017 03:20:35 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to:cc :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=yZgXq9J3L7e+zX7z5+00A0b2dMR9msLqAqPGzzEEuNNZPZAeBX 88ZdZ2hDz4MsVW0CFOvzW+giCg329eYwHKriGRwMtNLOy7VzI1H1hH0ROGVZow9r cVglq8xyWCqeqira7gBo9uTN8/xMK8u7loiRhxs1XchkBUvcGh4PtT8yo= 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:cc :from:subject:message-id:date:mime-version:content-type; s= default; bh=AZNK9jdLPxkxuj2M8xG2Oc8B9EU=; b=epeFqVFz1fkJ7prA/29P exNzRPqEeKiXeOpgweyzOpL82oKNnOUi6mUIIGDNhaOlsr+GbbuCmGzC0Havpo6V 1yRQva3rNMA1LDo2XEgA6apk8NtefoxUfpHii20IwjWeSeW629nDxlcBZOAKeKwN 13rVlpZSncFSVO4CCFI8sj4= 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-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 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, 7 Aug 2017 19:20:06 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.2.1 MIME-Version: 1.0 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) 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 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